Fix hipMemcpy3D (#1798)

Fixes #1790 and #1791. hipMemcpy3D still requires further refactoring for different input and output combinations.

[ROCm/hip commit: 8c5e5e435b]

projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp

@@ -0,0 +1,110 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+template <typename T>
+void runTest(int width,int height,int depth, hipChannelFormatKind formatKind)
+{
+    unsigned int size = width * height * depth * sizeof(T);
+    T* hData = (T*) malloc(size);
+    memset(hData, 0, size);
+
+    for (int i = 0; i < depth; i++) {
+        for (int j = 0; j < height; j++) {
+            for (int k = 0; k < width; k++) {
+                hData[i*width*height + j*width +k] = i*width*height + j*width + k;
+            }
+        }
+    }
+    printf("test- sizeof(T) =%d\n", sizeof(T));
+    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, formatKind);
+    hipArray *arr,*arr1;
+
+    HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
+    HIPCHECK(hipMalloc3DArray(&arr1, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
+    hipMemcpy3DParms myparms = {0};
+    myparms.srcPos = make_hipPos(0,0,0);
+    myparms.dstPos = make_hipPos(0,0,0);
+    myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
+    myparms.dstArray = arr;
+    myparms.extent = make_hipExtent(width , height, depth);
+#ifdef __HIP_PLATFORM_NVCC__
+    myparms.kind = cudaMemcpyHostToDevice;
+#else
+    myparms.kind = hipMemcpyHostToDevice;
+#endif
+    HIPCHECK(hipMemcpy3D(&myparms));
+    HIPCHECK(hipDeviceSynchronize());
+    //Array to Array
+    memset(&myparms,0x0, sizeof(hipMemcpy3DParms));
+    myparms.srcPos = make_hipPos(0,0,0);
+    myparms.dstPos = make_hipPos(0,0,0);
+    myparms.srcArray = arr;
+    myparms.dstArray = arr1;
+    myparms.extent = make_hipExtent(width, height, depth);
+#ifdef __HIP_PLATFORM_NVCC__
+    myparms.kind = cudaMemcpyDeviceToDevice;
+#else
+    myparms.kind = hipMemcpyDeviceToDevice;
+#endif
+    HIPCHECK(hipMemcpy3D(&myparms));
+    HIPCHECK(hipDeviceSynchronize());
+
+    T *hOutputData = (T*) malloc(size);
+    memset(hOutputData, 0,  size);
+    //Device to host
+    memset(&myparms,0x0, sizeof(hipMemcpy3DParms));
+    myparms.srcPos = make_hipPos(0,0,0);
+    myparms.dstPos = make_hipPos(0,0,0);
+    myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height);
+    myparms.srcArray = arr1;
+    myparms.extent = make_hipExtent(width, height, depth);
+#ifdef __HIP_PLATFORM_NVCC__
+    myparms.kind = cudaMemcpyDeviceToHost;
+#else
+    myparms.kind = hipMemcpyDeviceToHost;
+#endif
+    HIPCHECK(hipMemcpy3D(&myparms));
+    HIPCHECK(hipDeviceSynchronize());
+
+    // Check result
+    HipTest::checkArray(hData,hOutputData,width,height,depth);
+    hipFreeArray(arr);
+    hipFreeArray(arr1);
+    free(hData);
+    free(hOutputData);
+}
+
+int main(int argc, char **argv)
+{
+    for(int i=1;i<25;i++)
+    {
+        runTest<float>(i,i,i, hipChannelFormatKindFloat);
+        runTest<int>(i+1,i,i, hipChannelFormatKindSigned);
+        runTest<char>(i,i+1,i, hipChannelFormatKindSigned);
+    }
+    passed();
+}

projects/hip/tests/src/texture/simpleTexture2DLayered.cpp

@@ -65,7 +65,7 @@ void runTest(int width,int height,int num_layers,texture<T, hipTextureType2DLaye
     myparms.dstPos = make_hipPos(0,0,0);
     myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
     myparms.dstArray = arr;
-    myparms.extent = make_hipExtent(width, height, num_layers);
+    myparms.extent = make_hipExtent(width , height, num_layers);
     //myparms.kind = hipMemcpyHostToDevice;
     HIPCHECK(hipMemcpy3D(&myparms));
 

projects/hip/tests/src/texture/simpleTexture3D.cpp

@@ -21,15 +21,12 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
  * TEST: %t
  * HIT_END
  */
 #include "test_common.h"
 
-//typedef char T;
-const char *sampleName = "simpleTexture3D";
-
 // Texture reference for 3D texture
 #if __HIP__
 __hip_pinned_shadow__
@@ -47,7 +44,7 @@ __hip_pinned_shadow__
 texture<char, hipTextureType3D, hipReadModeElementType> texc;
 
 template <typename T>
-__global__ void simpleKernel3DArray(T* outputData, 
+__global__ void simpleKernel3DArray(T* outputData,
                                     int width,
                                     int height,int depth)
 {
@@ -55,21 +52,18 @@ __global__ void simpleKernel3DArray(T* outputData,
        for (int j = 0; j < height; j++) {
            for (int k = 0; k < width; k++) {
                if(std::is_same<T, float>::value)
-                   outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i);
+                   outputData[i*width*height + j*width + k] = tex3D(texf, k, j, i);
                else if(std::is_same<T, int>::value)
-                   outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i);
+                   outputData[i*width*height + j*width + k] = tex3D(texi, k, j, i);
                else if(std::is_same<T, char>::value)
-                   outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i);
+                   outputData[i*width*height + j*width + k] = tex3D(texc, k, j, i);
            }
        }
     }
 }
 
-////////////////////////////////////////////////////////////////////////////////
-//! Run a simple test for tex3D
-////////////////////////////////////////////////////////////////////////////////
 template <typename T>
-void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipReadModeElementType> *tex)
+void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipReadModeElementType> *tex, hipChannelFormatKind formatKind)
 {
     unsigned int size = width * height * depth * sizeof(T);
     T* hData = (T*) malloc(size);
@@ -84,17 +78,21 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
     }
 
     // Allocate array and copy image data
-    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindSigned);
+    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, formatKind);
     hipArray *arr;
 
-    HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayCubemap));
+    HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
     hipMemcpy3DParms myparms = {0};
     myparms.srcPos = make_hipPos(0,0,0);
     myparms.dstPos = make_hipPos(0,0,0);
     myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
     myparms.dstArray = arr;
     myparms.extent = make_hipExtent(width, height, depth);
+#ifdef __HIP_PLATFORM_NVCC__
+    myparms.kind = cudaMemcpyHostToDevice;
+#else
     myparms.kind = hipMemcpyHostToDevice;
+#endif
     HIPCHECK(hipMemcpy3D(&myparms));
 
     // set texture parameters
@@ -119,7 +117,7 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
 
     // copy result from device to host
     HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
-    HipTest::checkArray(hData,hOutputData,width,height,depth); 
+    HipTest::checkArray(hData,hOutputData,width,height,depth);
 
     hipFree(dData);
     hipFreeArray(arr);
@@ -127,18 +125,13 @@ void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipRead
     free(hOutputData);
 }
 
-////////////////////////////////////////////////////////////////////////////////
-// Program main
-////////////////////////////////////////////////////////////////////////////////
 int main(int argc, char **argv)
 {
-    printf("%s starting...\n", sampleName);
     for(int i=1;i<25;i++)
     {
-        runTest<float>(i,i,i,&texf);
-        runTest<int>(i+1,i,i,&texi);
-        runTest<char>(i,i+1,i,&texc);
+        runTest<float>(i,i,i,&texf, hipChannelFormatKindFloat);
+        runTest<int>(i+1,i,i,&texi, hipChannelFormatKindSigned);
+        runTest<char>(i,i+1,i,&texc, hipChannelFormatKindSigned);
     }
     passed();
 }
-