From 03adc12474cf322e48c6ec80ef707d01f98fc8a2 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 --- hipamd/bin/hipconfig | 23 ++++++++++++++++++----- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/hipamd/bin/hipconfig b/hipamd/bin/hipconfig index e70e6ae970..520a58cc2c 100755 --- a/hipamd/bin/hipconfig +++ b/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 0156388a6b57864e9034522622087539bd120763 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 --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 78 +++++++++++++------ 1 file changed, 56 insertions(+), 22 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 2d63133039..1617c0fd9e 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 3a57510fbbf68f3cfd970c4b849445bf14438f07. --- hipamd/bin/hipcc | 6 ------ 1 file changed, 6 deletions(-) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index 4c6a5ded07..b649bb6e7f 100755 --- a/hipamd/bin/hipcc +++ b/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 9002c7d09d2deecf0ad0caf4dd19274750cd6245 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 2a956540ee299a169114a293b7d2d950b5d9e9e6. --- hipamd/bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index b649bb6e7f..d6eea8457b 100755 --- a/hipamd/bin/hipcc +++ b/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 b40d9c7849cee4f21d9d7249629e216bbc384300 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 8062305f0b86dee730848b8bf10c14bfb7863157. --- hipamd/bin/hipcc | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index d6eea8457b..d227eea267 100755 --- a/hipamd/bin/hipcc +++ b/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 41afe4d9478a0934f69a68fec0c7c28d57491643 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 --- hipamd/tests/src/runtimeApi/memory/hipMemset.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 1617c0fd9e..5c8b1a18c1 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 99540373cf1459fcb5b48a055a8b7d549594b864 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. --- hipamd/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index ef5f80aee4..8e29bee68e 100644 --- a/hipamd/src/hip_memory.cpp +++ b/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 5900416629c1642396315b6242137e45b72aae24 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 --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 5c8b1a18c1..bcb0adbeef 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 --- .../include/hip/hcc_detail/hip_runtime_api.h | 28 +++++++++++ hipamd/src/hip_memory.cpp | 48 +++++++++++++++---- 2 files changed, 66 insertions(+), 10 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index b6ae88729a..9c50ca4755 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/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/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 8e29bee68e..b194faa1ca 100644 --- a/hipamd/src/hip_memory.cpp +++ b/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 365d08535ba9a9f5ce1a7152ed48c30409784bfe 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 --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 61 ++++++++++++++++++- hipamd/tests/src/test_common.cpp | 7 +++ hipamd/tests/src/test_common.h | 1 + 3 files changed, 66 insertions(+), 3 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index bcb0adbeef..8ac5beaf55 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index bacbf35a22..8820381826 100644 --- a/hipamd/tests/src/test_common.h +++ b/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 b46e684d2ecf83ebdfe4e2e1d2a8c969b3053326 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 --- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 02c4b7ee61..abce9a5fc6 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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 8b7baa0bd9185f3b09d05ee7b3141106b5c95cb1 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 --- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 3 ++- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 6 +++--- hipamd/src/hip_memory.cpp | 2 +- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 9c50ca4755..73996982d1 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/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/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index abce9a5fc6..a463d17e6d 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index b194faa1ca..d8a9bd5708 100644 --- a/hipamd/src/hip_memory.cpp +++ b/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 38b7a43b43278117f24f2ef57778b9add5fbb12a 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. --- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index a463d17e6d..add4c3f238 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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 8af4e2b5e48ebf8304f91c0e1ee8e1020b2faabe 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. --- hipamd/tests/src/runtimeApi/memory/hipMemset.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp index 8ac5beaf55..abff987437 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 --- .../markdown/CUDA_Driver_API_functions_supported_by_HIP.md | 4 ++-- hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index dc22286d84..89ca2d74e2 100644 --- a/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/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/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 8ad95cb537..8dfc429c4b 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/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}},