From 370d039d241c0e3c193df9e360ac43f134a91702 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Tue, 15 Oct 2019 14:18:42 -0400 Subject: [PATCH 1/6] Update the declarations of hipMemsetD8, hipMemsetD8Async, hipMemsetD16, hipMemsetD16Async. These functions are type aware and take in as their third argument the number of elements in the buffer, not the buffer size. Change the name of this argument from sizeBytes to count to align with the above description. [ROCm/hip commit: 0b52c1d9d839864a36ee019f7ce3345c405e1fe6] --- .../include/hip/hcc_detail/hip_runtime_api.h | 16 ++++++------- projects/hip/src/hip_memory.cpp | 24 +++++++++---------- 2 files changed, 20 insertions(+), 20 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index c07d2ad9f1..b4402fd67a 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/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 diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index ed1422fcda..48b83287f3 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/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) { - HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, sizeBytes); +hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) { + 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 ) { - HIP_INIT_SPECIAL_API(hipMemsetD8Async, (TRACE_MCMD), dst, value, sizeBytes, 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, 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){ - HIP_INIT_SPECIAL_API(hipMemsetD16, (TRACE_MCMD), dst, value, sizeBytes); +hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t count){ + 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 ){ - HIP_INIT_SPECIAL_API(hipMemsetD16Async, (TRACE_MCMD), dst, value, sizeBytes, 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, 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); } From 28089f61b6bcbb5c4f61b42a602c28502f4ff9ac Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Tue, 15 Oct 2019 14:20:14 -0400 Subject: [PATCH 2/6] hipMemset2D test should pass only if both async and sync subtests pass. [ROCm/hip commit: 3db2ecc52b379caaf8f80842089701b37098a323] --- projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp index b3bcf42222..73f3f5d415 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp @@ -113,8 +113,8 @@ int main(int argc, char *argv[]) hipCtxCreate(&context, 0, p_gpuDevice); bool testResult = false; - testResult = testhipMemset2D(memsetval, p_gpuDevice); - testResult = testhipMemset2DAsync(memsetval, p_gpuDevice); + testResult &= testhipMemset2D(memsetval, p_gpuDevice); + testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice); hipCtxDestroy(context); if(testResult){ passed(); From 5a29d10dbbc57d998edf2a6264777b06aa69c771 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Tue, 15 Oct 2019 14:24:04 -0400 Subject: [PATCH 3/6] Add async subtest to hipMemSet3D [ROCm/hip commit: 00425bdf3dd3a26fe502458d0cf77349603be1af] --- .../src/runtimeApi/memory/hipMemset3D.cpp | 55 ++++++++++++++++++- 1 file changed, 53 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp index 11bd656761..ac26280314 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp @@ -79,12 +79,63 @@ bool testhipMemset3D(int memsetval,int p_gpuDevice) 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 Date: Tue, 15 Oct 2019 15:29:14 -0400 Subject: [PATCH 4/6] Update indentation in the hipMemset3D test. Replace all tabs with four spaces. [ROCm/hip commit: 346bfa90d6783c4743f7df2b39750b079678f6dd] --- .../src/runtimeApi/memory/hipMemset3D.cpp | 52 +++++++++---------- 1 file changed, 26 insertions(+), 26 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp index ac26280314..ce2459a438 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp +++ b/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); - HIPASSERT(A_h != NULL); - for (size_t i=0; i Date: Tue, 15 Oct 2019 17:15:49 -0400 Subject: [PATCH 5/6] In the hipMemset2D and hipMemset3D tests synchronize with the default stream after performing an async memset. [ROCm/hip commit: f5af263abae4d9d7f58f1bf529f04fa8caa28423] --- projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp | 7 ++++--- projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp | 5 +++-- 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp index 73f3f5d415..449f0b6f78 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp +++ b/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 Date: Wed, 16 Oct 2019 11:02:38 -0400 Subject: [PATCH 6/6] hipMemset2D and hipMemset3D tests should be passing by default. [ROCm/hip commit: c747b77ac1d8e495d2ad72f795adb46bf1615d69] --- projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp | 2 +- projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp index 449f0b6f78..2eb62a859f 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp @@ -113,7 +113,7 @@ 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 &= testhipMemset2DAsync(memsetval, p_gpuDevice); hipCtxDestroy(context); diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp index a47b609c73..1917559f2a 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp @@ -134,7 +134,7 @@ int main(int argc, char *argv[]) { HipTest::parseStandardArguments(argc, argv, true); HIPCHECK(hipSetDevice(p_gpuDevice)); - bool testResult = false; + bool testResult = true; testResult &= testhipMemset3D(memsetval, p_gpuDevice); testResult &= testhipMemset3DAsync(memsetval, p_gpuDevice); if (testResult) {