From 8b25f7b4cf19bf2b4a0bb76eba61e659a8d89be9 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 --- bin/hipconfig | 23 ++++++++++++++++++----- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/bin/hipconfig b/bin/hipconfig index e70e6ae970..520a58cc2c 100755 --- a/bin/hipconfig +++ b/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 026dc9457e2c910c6a0dbfd732af31d72c5e1092 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/tests/src/runtimeApi/memory/hipMemset.cpp b/tests/src/runtimeApi/memory/hipMemset.cpp index 2d63133039..1617c0fd9e 100644 --- a/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 2cd34f3c01039cf4e790687c5fd31b16d5c29ca0. --- bin/hipcc | 6 ------ 1 file changed, 6 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 4c6a5ded07..b649bb6e7f 100755 --- a/bin/hipcc +++ b/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 1bd8d24278033a3ad8f1bf189c7ffca280c85598 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 1c6365fbe77c4d15b83f08f3a32b7bd3f1ee2640. --- bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index b649bb6e7f..d6eea8457b 100755 --- a/bin/hipcc +++ b/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 e8556388f9acfe8103576e8e120cac45315ef005 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 eb33507774444a984e6f8bed9671a0b0e78c7a07. --- bin/hipcc | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index d6eea8457b..d227eea267 100755 --- a/bin/hipcc +++ b/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 77c6d5451fd156e62542d88a2d2dec7f3be0c800 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 --- tests/src/runtimeApi/memory/hipMemset.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/src/runtimeApi/memory/hipMemset.cpp b/tests/src/runtimeApi/memory/hipMemset.cpp index 1617c0fd9e..5c8b1a18c1 100644 --- a/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 8d92d1ebd72a9bd4b0e72a3ca4fed4ae3d4dc8d9 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. --- src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index ef5f80aee4..8e29bee68e 100644 --- a/src/hip_memory.cpp +++ b/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 131c0ea3261a1db3c6031ecb2394d59f644c6770 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/tests/src/runtimeApi/memory/hipMemset.cpp b/tests/src/runtimeApi/memory/hipMemset.cpp index 5c8b1a18c1..bcb0adbeef 100644 --- a/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 ++++++++++++++ src/hip_memory.cpp | 48 +++++++++++++++++++----- 2 files changed, 66 insertions(+), 10 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index b6ae88729a..9c50ca4755 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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/src/hip_memory.cpp b/src/hip_memory.cpp index 8e29bee68e..b194faa1ca 100644 --- a/src/hip_memory.cpp +++ b/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 bfde8a7fab262f23079a13a7639986332644d88a 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 +++++++++++++++++++++-- tests/src/test_common.cpp | 7 +++ tests/src/test_common.h | 1 + 3 files changed, 66 insertions(+), 3 deletions(-) diff --git a/tests/src/runtimeApi/memory/hipMemset.cpp b/tests/src/runtimeApi/memory/hipMemset.cpp index bcb0adbeef..8ac5beaf55 100644 --- a/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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/tests/src/test_common.h b/tests/src/test_common.h index bacbf35a22..8820381826 100644 --- a/tests/src/test_common.h +++ b/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 0337b686ef68b701717af05278b701c3a3916fa4 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 --- include/hip/nvcc_detail/hip_runtime_api.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 02c4b7ee61..abce9a5fc6 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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 5cbd28f29b7e55695023c9110d71eee0f2b4b834 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 --- include/hip/hcc_detail/hip_runtime_api.h | 3 ++- include/hip/nvcc_detail/hip_runtime_api.h | 6 +++--- src/hip_memory.cpp | 2 +- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 9c50ca4755..73996982d1 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index abce9a5fc6..a463d17e6d 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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/src/hip_memory.cpp b/src/hip_memory.cpp index b194faa1ca..d8a9bd5708 100644 --- a/src/hip_memory.cpp +++ b/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 8db717c769b591bed5fb279e3259922ce964f977 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. --- include/hip/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index a463d17e6d..add4c3f238 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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 51f29b9cee9633b8f7e90f5eb616a9adae188d29 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. --- tests/src/runtimeApi/memory/hipMemset.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/src/runtimeApi/memory/hipMemset.cpp b/tests/src/runtimeApi/memory/hipMemset.cpp index 8ac5beaf55..abff987437 100644 --- a/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/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 --- docs/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/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index dc22286d84..89ca2d74e2 100644 --- a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/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/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 8ad95cb537..8dfc429c4b 100644 --- a/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/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}},