From ec494ba62851962ff08674626d61c94bd6419c8e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 27 Feb 2019 14:10:21 +0530 Subject: [PATCH 01/15] [hipcofig] Update HIP_PLATFORM detection logic HIP_PLATFORM detection logic relied on finding a working KFD. If it was found, the platform was set as hcc else as nvcc. However this logic is flawed since it is possible for the development system to only have the user mode bits to build HIP application code. Hence the better logic is to rely on finding a suitable compiler. The new logic is as follows: - look for a working HCC. If found, platform is set as hcc. - else look for a working NVCC. If found, platform is set as nvcc. - else the platform defaults to hcc for now. Change-Id: Ifcc42c29a19f722153d5c23c55f1a8765dceaf6b [ROCm/clr commit: 03adc12474cf322e48c6ec80ef707d01f98fc8a2] --- projects/clr/hipamd/bin/hipconfig | 23 ++++++++++++++++++----- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/bin/hipconfig b/projects/clr/hipamd/bin/hipconfig index e70e6ae970..520a58cc2c 100755 --- a/projects/clr/hipamd/bin/hipconfig +++ b/projects/clr/hipamd/bin/hipconfig @@ -59,6 +59,18 @@ sub parse_config_file { } } +#--- +# Function to check if executable can be run +sub can_run { + my ($exe) = @_; + `$exe --version 2>&1`; + if ($? == 0) { + return 1; + } else { + return 0; + } +} + $CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda'; $HCC_HOME=$ENV{'HCC_HOME'} // '/opt/rocm/hcc'; $HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa'; @@ -67,12 +79,13 @@ $HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa'; #HIP_PLATFORM controls whether to use NVCC or HCC for compilation: $HIP_PLATFORM=$ENV{'HIP_PLATFORM'}; if (not defined $HIP_PLATFORM) { - $NAMDGPUNODES=`cat /sys/class/kfd/kfd/topology/nodes/*/properties 2>/dev/null | grep -c 'simd_count [1-9]'`; - - if ($NAMDGPUNODES > 0) { - $HIP_PLATFORM = "hcc" - } else { + if (can_run("$HCC_HOME/bin/hcc") or can_run("hcc")) { + $HIP_PLATFORM = "hcc"; + } elsif (can_run("$CUDA_PATH/bin/nvcc") or can_run("nvcc")) { $HIP_PLATFORM = "nvcc"; + } else { + # Default to hcc for now + $HIP_PLATFORM = "hcc"; } } From 1a97c0fafc8669e613eb79a9dad51f07573f070b Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 28 Feb 2019 06:54:49 +0530 Subject: [PATCH 02/15] Update hipMemset test [ROCm/clr commit: 0156388a6b57864e9034522622087539bd120763] --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 78 +++++++++++++------ 1 file changed, 56 insertions(+), 22 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 2d63133039..1617c0fd9e 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp @@ -37,32 +37,66 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" +bool testhipMemset(int memsetval,int p_gpuDevice) +{ + size_t Nbytes = N*sizeof(char); + printf ("testhipMemset N=%zu memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice); + char *A_d; + char *A_h; + bool testResult = true; -int main(int argc, char* argv[]) { - HipTest::parseStandardArguments(argc, argv, true); - - HIPCHECK(hipSetDevice(p_gpuDevice)); - - size_t Nbytes = N * sizeof(char); - - printf("N=%zu memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice); - - char* A_d; - char* A_h; - - HIPCHECK(hipMalloc(&A_d, Nbytes)); + HIPCHECK ( hipMalloc(&A_d, Nbytes) ); A_h = (char*)malloc(Nbytes); + HIPCHECK ( hipMemset(A_d, memsetval, Nbytes) ); + HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); - HIPCHECK(hipMemset(A_d, memsetval, Nbytes)); - - HIPCHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); - - for (int i = 0; i < N; i++) { + for (int i=0; i Date: Thu, 28 Feb 2019 11:19:35 -0500 Subject: [PATCH 03/15] Revert "Consume the code obj args to prevent duplicates" This reverts commit 0e1fc751eac55be9822ea85e19ea00d1cd1837e4. [ROCm/clr commit: 510590ac1d3551d3a9657fc11216f796879e2cea] --- projects/clr/hipamd/bin/hipcc | 6 ------ 1 file changed, 6 deletions(-) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 4c6a5ded07..b649bb6e7f 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -424,14 +424,8 @@ foreach $arg (@ARGV) $swallowArg = 1; } - if (($trimarg eq '-mno-code-object-v3') and ($HIP_PLATFORM eq 'clang') ) { - $useCodeObjectV3 = 0; - $swallowArg = 1; - } - if (($trimarg eq '-mcode-object-v3') and ($HIP_PLATFORM eq 'clang') ) { $useCodeObjectV3 = 1; - $swallowArg = 1; } if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) { From 89721c8ce88f19c54a2b9532d6e2543c1a4203ab Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 28 Feb 2019 11:20:04 -0500 Subject: [PATCH 04/15] Revert "Change code-object flag to only HIP-Clang" This reverts commit 1e483af21ba9a2388b3823d289ec6069b456408e. [ROCm/clr commit: 9002c7d09d2deecf0ad0caf4dd19274750cd6245] --- projects/clr/hipamd/bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index b649bb6e7f..d6eea8457b 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -763,7 +763,7 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') { if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') { $HIPCXXFLAGS .= " -mcode-object-v3"; -} elsif ($HIP_PLATFORM eq 'clang') { +} else { $HIPCXXFLAGS .= " -mno-code-object-v3"; } From 169bfb6b75b5cdabb8091f0a7a7a93f41a5ce66c Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 28 Feb 2019 11:21:47 -0500 Subject: [PATCH 05/15] Revert "hipcc should consume -mcode-object-v3 flag" This reverts commit ac4b2b03ac76ff0466faa234fc3c137b55d99c33. [ROCm/clr commit: b40d9c7849cee4f21d9d7249629e216bbc384300] --- projects/clr/hipamd/bin/hipcc | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index d6eea8457b..d227eea267 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -320,7 +320,6 @@ my $runCmd = 1; my $buildDeps = 0; my $linkType = 1; my $setLinkType = 0; -my $useCodeObjectV3 = 0; my @options = (); my @inputs = (); @@ -424,10 +423,6 @@ foreach $arg (@ARGV) $swallowArg = 1; } - if (($trimarg eq '-mcode-object-v3') and ($HIP_PLATFORM eq 'clang') ) { - $useCodeObjectV3 = 1; - } - if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) { $arg = "--cuda-device-only"; } @@ -761,12 +756,6 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') { $HIPCXXFLAGS .= " --cuda-host-only"; } -if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') { - $HIPCXXFLAGS .= " -mcode-object-v3"; -} else { - $HIPCXXFLAGS .= " -mno-code-object-v3"; -} - # Add --hip-link only if there are no source files. if (!$needCXXFLAGS and $HIP_PLATFORM eq 'clang') { $HIPLDFLAGS .= " --hip-link"; From e6f5850ebb652e9b65e45f2168b492ba9dad0507 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 1 Mar 2019 03:46:57 +0530 Subject: [PATCH 06/15] Fix hipMemset test for HIP/NVCC [ROCm/clr commit: 41afe4d9478a0934f69a68fec0c7c28d57491643] --- projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 1617c0fd9e..5c8b1a18c1 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp @@ -66,7 +66,7 @@ bool testhipMemsetAsync(int memsetval,int p_gpuDevice) { size_t Nbytes = N*sizeof(int); printf ("testhipMemsetAsync N=%zu memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice); - hipDeviceptr_t A_d; + char *A_d; char *A_h; bool testResult = true; From 19b361281adddac6998bb1488164d842713f5a41 Mon Sep 17 00:00:00 2001 From: Wilkin Chau Date: Thu, 28 Feb 2019 22:42:46 +0000 Subject: [PATCH 07/15] Fix hipMemset3D test Calculate the allocated size based on the width, height and depth. [ROCm/clr commit: 99540373cf1459fcb5b48a055a8b7d549594b864] --- projects/clr/hipamd/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index ef5f80aee4..8e29bee68e 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -350,7 +350,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh } // hardcoded 128 bytes *pitch = ((((int)width - 1) / 128) + 1) * 128; - const size_t sizeBytes = (*pitch) * height; + const size_t sizeBytes = (*pitch) * height * ((depth==0) ? 1 : depth); auto ctx = ihipGetTlsDefaultCtx(); From b3ba23ba812020caf9ca29458498095b4ff13138 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sat, 2 Mar 2019 23:38:37 +0530 Subject: [PATCH 08/15] Fix review comments [ROCm/clr commit: 5900416629c1642396315b6242137e45b72aae24] --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 5c8b1a18c1..bcb0adbeef 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp @@ -57,7 +57,7 @@ bool testhipMemset(int memsetval,int p_gpuDevice) break; } } - hipFree(A_d); + HIPCHECK(hipFree(A_d)); free(A_h); return testResult; } @@ -75,6 +75,7 @@ bool testhipMemsetAsync(int memsetval,int p_gpuDevice) hipStream_t stream; HIPCHECK(hipStreamCreate(&stream)); HIPCHECK ( hipMemsetAsync(A_d, memsetval, Nbytes, stream )); + HIPCHECK ( hipStreamSynchronize(stream)); HIPCHECK ( hipMemcpy(A_h, (void*)A_d, Nbytes, hipMemcpyDeviceToHost)); for (int i=0; i Date: Wed, 27 Feb 2019 15:42:54 +0000 Subject: [PATCH 09/15] Add hipMemsetD32 and hipMemsetD32Async Add 2 extra memset functions which fills memory with integer-typed data Also change the parameters of ihipMemset to better explain the semantic [ROCm/clr commit: 392271f4dbffeedd68488271c5fe2d0a075bd207] --- .../include/hip/hcc_detail/hip_runtime_api.h | 28 +++++++++++ projects/clr/hipamd/src/hip_memory.cpp | 48 +++++++++++++++---- 2 files changed, 66 insertions(+), 10 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index b6ae88729a..9c50ca4755 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1504,6 +1504,17 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes); */ hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes); +/** + * @brief Fills the memory area pointed to by dest with the constant integer + * value for specified number of times. + * + * @param[out] dst Data being filled + * @param[in] constant value to be set + * @param[in] number of values to be set + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized + */ +hipError_t hipMemsetD32(hipDeviceptr_t dest, int value, size_t count); + /** * @brief Fills the first sizeBytes bytes of the memory area pointed to by dev with the constant * byte value value. @@ -1521,6 +1532,23 @@ hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeByte */ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream __dparm(0)); +/** + * @brief Fills the memory area pointed to by dev with the constant integer + * value for specified number of times. + * + * hipMemsetD32Async() is asynchronous with respect to the host, so the call may return before the + * memset is complete. The operation can optionally be associated to a stream by passing a non-zero + * stream argument. If stream is non-zero, the operation may overlap with operations in other + * streams. + * + * @param[out] dst Pointer to device memory + * @param[in] value - Value to set for each byte of specified memory + * @param[in] count - number of values to be set + * @param[in] stream - Stream identifier + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + */ +hipError_t hipMemsetD32Async(void* dst, int value, size_t count, hipStream_t stream __dparm(0)); + /** * @brief Fills the memory area pointed to by dst with the constant value. * diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 8e29bee68e..b194faa1ca 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1508,13 +1508,13 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, } // namespace template -void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { +void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t count) { static constexpr uint32_t block_dim = 256; - const uint32_t grid_dim = clamp_integer(sizeBytes / block_dim, 1, UINT32_MAX); + const uint32_t grid_dim = clamp_integer(count / block_dim, 1, UINT32_MAX); hipLaunchKernelGGL(hip_fill_n, dim3(grid_dim), dim3{block_dim}, 0u, stream, ptr, - sizeBytes, std::move(val)); + count, std::move(val)); } template @@ -1533,20 +1533,20 @@ typedef enum ihipMemsetDataType { ihipMemsetDataTypeInt = 2 }ihipMemsetDataType; -hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetDataType copyDataType ) +hipError_t ihipMemset(void* dst, int value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType ) { hipError_t e = hipSuccess; - if (sizeBytes == 0) return e; + if (count == 0) return e; if (stream && (dst != NULL)) { if(copyDataType == ihipMemsetDataTypeChar){ - if ((sizeBytes & 0x3) == 0) { + if ((count & 0x3) == 0) { // use a faster dword-per-workitem copy: try { value = value & 0xff; uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; - ihipMemsetKernel (stream, static_cast (dst), value32, sizeBytes/sizeof(uint32_t)); + ihipMemsetKernel (stream, static_cast (dst), value32, count/sizeof(uint32_t)); } catch (std::exception &ex) { e = hipErrorInvalidValue; @@ -1554,7 +1554,7 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea } else { // use a slow byte-per-workitem copy: try { - ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + ihipMemsetKernel (stream, static_cast (dst), value, count); } catch (std::exception &ex) { e = hipErrorInvalidValue; @@ -1563,14 +1563,14 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea } else { if(copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value try { - ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + ihipMemsetKernel (stream, static_cast (dst), value, count); } catch (std::exception &ex) { e = hipErrorInvalidValue; } } else if(copyDataType == ihipMemsetDataTypeShort) { try { value = value & 0xffff; - ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + ihipMemsetKernel (stream, static_cast (dst), value, count); } catch (std::exception &ex) { e = hipErrorInvalidValue; } @@ -1719,6 +1719,18 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st return ihipLogStatus(e); }; +hipError_t hipMemsetD32Async(void* dst, int value, size_t count, hipStream_t stream) { + HIP_INIT_SPECIAL_API(hipMemsetD32Async, (TRACE_MCMD), dst, value, count, stream); + + hipError_t e = hipSuccess; + + stream = ihipSyncAndResolveStream(stream); + + e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeInt); + + return ihipLogStatus(e); +}; + hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes); @@ -1787,6 +1799,22 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes return ihipLogStatus(e); } +hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) { + HIP_INIT_SPECIAL_API(hipMemsetD32, (TRACE_MCMD), dst, value, count); + + hipError_t e = hipSuccess; + + hipStream_t stream = hipStreamNull; + stream = ihipSyncAndResolveStream(stream); + if (stream) { + e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeInt); + stream->locked_wait(); + } else { + e = hipErrorInvalidValue; + } + return ihipLogStatus(e); +} + hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ) { HIP_INIT_SPECIAL_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent); From 21cf5b0ae46d85183780956c2c72bd85dd2c2909 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Mon, 4 Mar 2019 17:11:54 +0000 Subject: [PATCH 10/15] Add direct test for hipMemsetD32 and hipMemsetD32Async [ROCm/clr commit: 365d08535ba9a9f5ce1a7152ed48c30409784bfe] --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 61 ++++++++++++++++++- projects/clr/hipamd/tests/src/test_common.cpp | 7 +++ projects/clr/hipamd/tests/src/test_common.h | 1 + 3 files changed, 66 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index bcb0adbeef..8ac5beaf55 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp @@ -26,11 +26,11 @@ THE SOFTWARE. * BUILD: %t %s ../../test_common.cpp * RUN: %t * //Small copy - * RUN: %t -N 10 --memsetval 0x42 + * RUN: %t -N 10 --memsetval 0x42 --memsetD32val 0x101 * // Oddball size - * RUN: %t -N 10013 --memsetval 0x5a + * RUN: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF * // Big copy - * RUN: %t -N 256M --memsetval 0xa6 + * RUN: %t -N 256M --memsetval 0xa6 --memsetD32val 0xCAFEBABE * HIT_END */ @@ -62,6 +62,30 @@ bool testhipMemset(int memsetval,int p_gpuDevice) return testResult; } +bool testhipMemsetD32(int memsetD32val,int p_gpuDevice) +{ + size_t Nbytes = N*sizeof(int); + printf ("testhipMemsetD32 N=%zu memsetD32val=%8x device=%d\n", N, memsetD32val, p_gpuDevice); + int *A_d; + int *A_h; + bool testResult = true; + + HIPCHECK ( hipMalloc(&A_d, Nbytes) ); + A_h = (int*)malloc(Nbytes); + HIPCHECK ( hipMemsetD32(A_d, memsetD32val, N) ); + HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + + for (int i=0; i= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetD32val argument"); + } + memsetD32val = ex; } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) { failed("Bad iterations argument"); diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h index bacbf35a22..8820381826 100644 --- a/projects/clr/hipamd/tests/src/test_common.h +++ b/projects/clr/hipamd/tests/src/test_common.h @@ -98,6 +98,7 @@ THE SOFTWARE. // standard command-line variables: extern size_t N; extern char memsetval; +extern int memsetD32val; extern int iterations; extern unsigned blocksPerCU; extern unsigned threadsPerBlock; From a4b654a5af81406f746061d9420971f0cefd3eb2 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Mon, 4 Mar 2019 20:11:12 -0800 Subject: [PATCH 11/15] Add implementation for NVCC path [ROCm/clr commit: b46e684d2ecf83ebdfe4e2e1d2a8c969b3053326] --- .../clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 02c4b7ee61..abce9a5fc6 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -655,11 +655,20 @@ inline static hipError_t hipMemset(void* devPtr, int value, size_t count) { return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count)); } +inline static hipError_t hipMemsetD32(void* devPtr, int value, size_t count) { + return hipCUDAErrorTohipError(cuMemsetD32(devPtr, value, count)); +} + inline static hipError_t hipMemsetAsync(void* devPtr, int value, size_t count, hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream)); } +inline static hipError_t hipMemsetD32Async(void* devPtr, int value, size_t count, + hipStream_t stream __dparm(0)) { + return hipCUDAErrorTohipError(cuMemsetD32Async(devPtr, value, count, stream)); +} + inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes) { return hipCUResultTohipError(cuMemsetD8(dest, value, sizeBytes)); } From e34b0ccd48d2cdf100ad4d112da07a91e6506495 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Tue, 5 Mar 2019 05:51:05 +0000 Subject: [PATCH 12/15] Address code review comments to use hipDeviceptr_t [ROCm/clr commit: 8b7baa0bd9185f3b09d05ee7b3141106b5c95cb1] --- .../clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h | 3 ++- .../clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 6 +++--- projects/clr/hipamd/src/hip_memory.cpp | 2 +- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 9c50ca4755..73996982d1 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1547,7 +1547,8 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st * @param[in] stream - Stream identifier * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree */ -hipError_t hipMemsetD32Async(void* dst, int value, size_t count, hipStream_t stream __dparm(0)); +hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, + hipStream_t stream __dparm(0)); /** * @brief Fills the memory area pointed to by dst with the constant value. diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index abce9a5fc6..a463d17e6d 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -655,7 +655,7 @@ inline static hipError_t hipMemset(void* devPtr, int value, size_t count) { return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count)); } -inline static hipError_t hipMemsetD32(void* devPtr, int value, size_t count) { +inline static hipError_t hipMemsetD32(hipDeviceptr_t devPtr, int value, size_t count) { return hipCUDAErrorTohipError(cuMemsetD32(devPtr, value, count)); } @@ -664,8 +664,8 @@ inline static hipError_t hipMemsetAsync(void* devPtr, int value, size_t count, return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream)); } -inline static hipError_t hipMemsetD32Async(void* devPtr, int value, size_t count, - hipStream_t stream __dparm(0)) { +inline static hipError_t hipMemsetD32Async(hipDeviceptr_t devPtr, int value, size_t count, + hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cuMemsetD32Async(devPtr, value, count, stream)); } diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index b194faa1ca..d8a9bd5708 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1719,7 +1719,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st return ihipLogStatus(e); }; -hipError_t hipMemsetD32Async(void* dst, int value, size_t count, hipStream_t stream) { +hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) { HIP_INIT_SPECIAL_API(hipMemsetD32Async, (TRACE_MCMD), dst, value, count, stream); hipError_t e = hipSuccess; From 34554f573004c9b877434c235c8f8e415c506938 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 5 Mar 2019 12:10:01 +0530 Subject: [PATCH 13/15] Update hip_runtime_api.h Use hipCUResultTohipError instead of hipCUDAErrorTohipError in hipMemsetD32 & hipMemsetD32Async. [ROCm/clr commit: 38b7a43b43278117f24f2ef57778b9add5fbb12a] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index a463d17e6d..add4c3f238 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -656,7 +656,7 @@ inline static hipError_t hipMemset(void* devPtr, int value, size_t count) { } inline static hipError_t hipMemsetD32(hipDeviceptr_t devPtr, int value, size_t count) { - return hipCUDAErrorTohipError(cuMemsetD32(devPtr, value, count)); + return hipCUResultTohipError(cuMemsetD32(devPtr, value, count)); } inline static hipError_t hipMemsetAsync(void* devPtr, int value, size_t count, @@ -666,7 +666,7 @@ inline static hipError_t hipMemsetAsync(void* devPtr, int value, size_t count, inline static hipError_t hipMemsetD32Async(hipDeviceptr_t devPtr, int value, size_t count, hipStream_t stream __dparm(0)) { - return hipCUDAErrorTohipError(cuMemsetD32Async(devPtr, value, count, stream)); + return hipCUResultTohipError(cuMemsetD32Async(devPtr, value, count, stream)); } inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes) { From 5c939d7f27cd16a7c01dd6fd66d5b52350fbd513 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 5 Mar 2019 12:11:11 +0530 Subject: [PATCH 14/15] Update hipMemset.cpp Address build issues on nvcc path. [ROCm/clr commit: 8af4e2b5e48ebf8304f91c0e1ee8e1020b2faabe] --- projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 8ac5beaf55..abff987437 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp @@ -72,7 +72,7 @@ bool testhipMemsetD32(int memsetD32val,int p_gpuDevice) HIPCHECK ( hipMalloc(&A_d, Nbytes) ); A_h = (int*)malloc(Nbytes); - HIPCHECK ( hipMemsetD32(A_d, memsetD32val, N) ); + HIPCHECK ( hipMemsetD32((hipDeviceptr_t)A_d, memsetD32val, N) ); HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); for (int i=0; i Date: Tue, 5 Mar 2019 18:13:18 +0300 Subject: [PATCH 15/15] [HIPIFY] Change CUDA Driver's functions' cuMemsetD32(Async) mapping cuMemsetD32(Async) -> hipMemsetD32(Async) (was hipMemset(Async)) based on: [#933] https://github.com/ROCm-Developer-Tools/HIP/pull/933 [ROCm/clr commit: dfc631fb4483484076e90124052a81c958806eee] --- .../markdown/CUDA_Driver_API_functions_supported_by_HIP.md | 4 ++-- .../hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index dc22286d84..89ca2d74e2 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -891,8 +891,8 @@ | `cuMemsetD2D32Async` | | | `cuMemsetD2D8` | | | `cuMemsetD2D8Async` | | -| `cuMemsetD32` | `hipMemset` | -| `cuMemsetD32Async` | `hipMemsetAsync` | +| `cuMemsetD32` | `hipMemsetD32` | +| `cuMemsetD32Async` | `hipMemsetD32Async` | | `cuMemsetD8` | `hipMemsetD8` | | `cuMemsetD8Async` | | | `cuMipmappedArrayCreate` | | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 8ad95cb537..8dfc429c4b 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -291,10 +291,10 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // no analogue {"cuMemsetD2D8Async", {"hipMemsetD2D8Async", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, // cudaMemset - {"cuMemsetD32", {"hipMemset", "", CONV_MEMORY, API_DRIVER}}, - {"cuMemsetD32_v2", {"hipMemset", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemsetD32", {"hipMemsetD32", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemsetD32_v2", {"hipMemsetD32", "", CONV_MEMORY, API_DRIVER}}, // cudaMemsetAsync - {"cuMemsetD32Async", {"hipMemsetAsync", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemsetD32Async", {"hipMemsetD32Async", "", CONV_MEMORY, API_DRIVER}}, // no analogue {"cuMemsetD8", {"hipMemsetD8", "", CONV_MEMORY, API_DRIVER}}, {"cuMemsetD8_v2", {"hipMemsetD8", "", CONV_MEMORY, API_DRIVER}},