Merge pull request #1544 from vsytch/master

QoL changes to the hipMemset family

[ROCm/hip commit: e1aac060da]

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -1742,10 +1742,10 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes);
  *
  *  @param[out] dst Data ptr to be filled
  *  @param[in]  constant value to be set
- *  @param[in]  sizeBytes Data size in bytes
+ *  @param[in]  number of values to be set
  *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
  */
-hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes);
+hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t count);
 
 /**
  *  @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant
@@ -1758,11 +1758,11 @@ hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeByte
  *
  *  @param[out] dst Data ptr to be filled
  *  @param[in]  constant value to be set
- *  @param[in]  sizeBytes Data size in bytes
+ *  @param[in]  number of values to be set
  *  @param[in]  stream - Stream identifier
  *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
  */
-hipError_t hipMemsetD8Async(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes, hipStream_t stream __dparm(0));
+hipError_t hipMemsetD8Async(hipDeviceptr_t dest, unsigned char value, size_t count, hipStream_t stream __dparm(0));
 
 /**
  *  @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant
@@ -1770,10 +1770,10 @@ hipError_t hipMemsetD8Async(hipDeviceptr_t dest, unsigned char value, size_t siz
  *
  *  @param[out] dst Data ptr to be filled
  *  @param[in]  constant value to be set
- *  @param[in]  sizeBytes Data size in bytes
+ *  @param[in]  number of values to be set
  *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
  */
-hipError_t hipMemsetD16(hipDeviceptr_t dest, unsigned short value, size_t sizeBytes);
+hipError_t hipMemsetD16(hipDeviceptr_t dest, unsigned short value, size_t count);
 
 /**
  *  @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant
@@ -1786,11 +1786,11 @@ hipError_t hipMemsetD16(hipDeviceptr_t dest, unsigned short value, size_t sizeBy
  *
  *  @param[out] dst Data ptr to be filled
  *  @param[in]  constant value to be set
- *  @param[in]  sizeBytes Data size in bytes
+ *  @param[in]  number of values to be set
  *  @param[in]  stream - Stream identifier
  *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
  */
-hipError_t hipMemsetD16Async(hipDeviceptr_t dest, unsigned short value, size_t sizeBytes, hipStream_t stream __dparm(0));
+hipError_t hipMemsetD16Async(hipDeviceptr_t dest, unsigned short value, size_t count, hipStream_t stream __dparm(0));
 
 /**
  *  @brief Fills the memory area pointed to by dest with the constant integer

projects/hip/src/hip_memory.cpp

@@ -1935,15 +1935,15 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, si
     return ihipLogStatus(e);
 };
 
-hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
+hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) {
-    HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, sizeBytes);
+    HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, count);
 
     hipError_t e = hipSuccess;
 
     hipStream_t stream = hipStreamNull;
     stream = ihipSyncAndResolveStream(stream);
     if (stream) {
-        e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
+        e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeChar);
         stream->locked_wait();
     } else {
         e = hipErrorInvalidValue;
@@ -1951,23 +1951,23 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
     return ihipLogStatus(e);
 }
 
-hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes , hipStream_t stream ) {
+hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t count , hipStream_t stream ) {
-    HIP_INIT_SPECIAL_API(hipMemsetD8Async, (TRACE_MCMD), dst, value, sizeBytes, stream);
+    HIP_INIT_SPECIAL_API(hipMemsetD8Async, (TRACE_MCMD), dst, value, count, stream);
 
     stream = ihipSyncAndResolveStream(stream);
     if (stream) {
-        return ihipLogStatus(ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar));
+        return ihipLogStatus(ihipMemset(dst, value, count, stream, ihipMemsetDataTypeChar));
     } else {
         return ihipLogStatus(hipErrorInvalidValue);
     }
 }
 
-hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t sizeBytes){
+hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t count){
-    HIP_INIT_SPECIAL_API(hipMemsetD16, (TRACE_MCMD), dst, value, sizeBytes);
+    HIP_INIT_SPECIAL_API(hipMemsetD16, (TRACE_MCMD), dst, value, count);
     hipError_t e = hipSuccess;
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
     if (stream) {
-        e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeShort);
+        e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeShort);
         if(hipSuccess == e)
             stream->locked_wait();
     } else {
@@ -1976,12 +1976,12 @@ hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t sizeByt
     return ihipLogStatus(e);
 }
 
-hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t sizeBytes, hipStream_t stream ){
+hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t count, hipStream_t stream ){
-    HIP_INIT_SPECIAL_API(hipMemsetD16Async, (TRACE_MCMD), dst, value, sizeBytes, stream);
+    HIP_INIT_SPECIAL_API(hipMemsetD16Async, (TRACE_MCMD), dst, value, count, stream);
 
     stream = ihipSyncAndResolveStream(stream);
     if (stream) {
-        return ihipLogStatus(ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeShort));
+        return ihipLogStatus(ihipMemset(dst, value, count, stream, ihipMemsetDataTypeShort));
     } else {
         return ihipLogStatus(hipErrorInvalidValue);
     }

projects/hip/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(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
-    HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
+    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;
+    bool testResult = true;
-    testResult = testhipMemset2D(memsetval, p_gpuDevice);
+    testResult &= testhipMemset2D(memsetval, p_gpuDevice);
-    testResult = testhipMemset2DAsync(memsetval, p_gpuDevice);
+    testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice);
     hipCtxDestroy(context);
     if(testResult){
        passed();

projects/hip/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);
+    A_h = (char*)malloc(sizeElements);
-	HIPASSERT(A_h != NULL);
+    HIPASSERT(A_h != NULL);
-	for (size_t i=0; i<elements; i++) {
+    for (size_t i=0; i<elements; i++) {
         A_h[i] = 1;
     }
-	HIPCHECK ( hipMemset3D( devPitchedPtr, memsetval, extent) );
+    HIPCHECK(hipMemset3D( devPitchedPtr, memsetval, extent));
-	hipMemcpy3DParms myparms = {0};
+    hipMemcpy3DParms myparms = {0};
-	myparms.srcPos = make_hipPos(0,0,0);
+    myparms.srcPos = make_hipPos(0,0,0);
-	myparms.dstPos = make_hipPos(0,0,0);
+    myparms.dstPos = make_hipPos(0,0,0);
-	myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
+    myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
-	myparms.srcPtr = devPitchedPtr;
+    myparms.srcPtr = devPitchedPtr;
-	myparms.extent = extent;
+    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 {