Merge pull request #1544 from vsytch/master

QoL changes to the hipMemset family

tests/src/runtimeApi/memory/hipMemset2D.cpp

@@ -45,7 +45,7 @@ bool testhipMemset2D(int memsetval,int p_gpuDevice)
     char *A_d;
     char *A_h;
     bool testResult = true;
-    HIPCHECK ( hipMemAllocPitch((hipDeviceptr_t*)&A_d, &pitch_A, width , numH,16) );
+    HIPCHECK(hipMemAllocPitch((hipDeviceptr_t*)&A_d, &pitch_A, width , numH,16));
     A_h = (char*)malloc(sizeElements);
     HIPASSERT(A_h != NULL);
     for (size_t i=0; i<elements; i++) {
@@ -89,8 +89,9 @@ bool testhipMemset2DAsync(int memsetval,int p_gpuDevice)
     }
     hipStream_t stream;
     HIPCHECK(hipStreamCreate(&stream));
-    HIPCHECK ( hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
-    HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
+    HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
+    HIPCHECK(hipStreamSynchronize(stream));
+    HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
 
     for (int i=0; i<elements; i++) {
         if (A_h[i] != memsetval) {
@@ -112,9 +113,9 @@ int main(int argc, char *argv[])
     hipCtx_t context;
     hipCtxCreate(&context, 0, p_gpuDevice);
 
-    bool testResult = false;
-    testResult = testhipMemset2D(memsetval, p_gpuDevice);
-    testResult = testhipMemset2DAsync(memsetval, p_gpuDevice);
+    bool testResult = true;
+    testResult &= testhipMemset2D(memsetval, p_gpuDevice);
+    testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice);
     hipCtxDestroy(context);
     if(testResult){
        passed();

tests/src/runtimeApi/memory/hipMemset3D.cpp

@@ -41,27 +41,27 @@ bool testhipMemset3D(int memsetval,int p_gpuDevice)
     size_t elements = numW* numH* depth;
 
 
-	printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
+    printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
     char *A_h;
     bool testResult = true;
     hipExtent extent = make_hipExtent(width, numH, depth);
     hipPitchedPtr devPitchedPtr;
 
     HIPCHECK(hipMalloc3D(&devPitchedPtr, extent));
-	A_h = (char*)malloc(sizeElements);
-	HIPASSERT(A_h != NULL);
-	for (size_t i=0; i<elements; i++) {
+    A_h = (char*)malloc(sizeElements);
+    HIPASSERT(A_h != NULL);
+    for (size_t i=0; i<elements; i++) {
         A_h[i] = 1;
     }
-	HIPCHECK ( hipMemset3D( devPitchedPtr, memsetval, extent) );
-	hipMemcpy3DParms myparms = {0};
-	myparms.srcPos = make_hipPos(0,0,0);
-	myparms.dstPos = make_hipPos(0,0,0);
-	myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
-	myparms.srcPtr = devPitchedPtr;
-	myparms.extent = extent;
+    HIPCHECK(hipMemset3D( devPitchedPtr, memsetval, extent));
+    hipMemcpy3DParms myparms = {0};
+    myparms.srcPos = make_hipPos(0,0,0);
+    myparms.dstPos = make_hipPos(0,0,0);
+    myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
+    myparms.srcPtr = devPitchedPtr;
+    myparms.extent = extent;
 #ifdef __HIP_PLATFORM_NVCC__
-	myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
+    myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
 #else
     myparms.kind = hipMemcpyDeviceToHost;
 #endif
@@ -69,7 +69,58 @@ bool testhipMemset3D(int memsetval,int p_gpuDevice)
 
     for (int i=0; i<elements; i++) {
         if (A_h[i] != memsetval) {
-			testResult = false;
+            testResult = false;
+            printf("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
+            break;
+        }
+    }
+    HIPCHECK(hipFree(devPitchedPtr.ptr));
+    free(A_h);
+    return testResult;
+}
+
+bool testhipMemset3DAsync(int memsetval,int p_gpuDevice)
+{
+    size_t numH = 256;
+    size_t numW = 256;
+    size_t depth = 10;
+    size_t width = numW * sizeof(char);
+    size_t sizeElements = width * numH * depth;
+    size_t elements = numW* numH* depth;
+
+
+    printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
+    char *A_h;
+    bool testResult = true;
+    hipExtent extent = make_hipExtent(width, numH, depth);
+    hipPitchedPtr devPitchedPtr;
+
+    HIPCHECK(hipMalloc3D(&devPitchedPtr, extent));
+    A_h = (char*)malloc(sizeElements);
+    HIPASSERT(A_h != NULL);
+    for (size_t i=0; i<elements; i++) {
+        A_h[i] = 1;
+    }
+    hipStream_t stream;
+    HIPCHECK(hipStreamCreate(&stream));
+    HIPCHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
+    HIPCHECK(hipStreamSynchronize(stream));
+    hipMemcpy3DParms myparms = {0};
+    myparms.srcPos = make_hipPos(0,0,0);
+    myparms.dstPos = make_hipPos(0,0,0);
+    myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
+    myparms.srcPtr = devPitchedPtr;
+    myparms.extent = extent;
+#ifdef __HIP_PLATFORM_NVCC__
+    myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
+#else
+    myparms.kind = hipMemcpyDeviceToHost;
+#endif
+    HIPCHECK(hipMemcpy3D(&myparms));
+
+    for (int i=0; i<elements; i++) {
+        if (A_h[i] != memsetval) {
+            testResult = false;
             printf("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
             break;
         }
@@ -82,9 +133,10 @@ bool testhipMemset3D(int memsetval,int p_gpuDevice)
 int main(int argc, char *argv[])
 {
     HipTest::parseStandardArguments(argc, argv, true);
-    bool testResult = false;
     HIPCHECK(hipSetDevice(p_gpuDevice));
-    testResult = testhipMemset3D(memsetval, p_gpuDevice);
+    bool testResult = true;
+    testResult &= testhipMemset3D(memsetval, p_gpuDevice);
+    testResult &= testhipMemset3DAsync(memsetval, p_gpuDevice);
     if (testResult) {
         passed();
     } else {