From 55da014b793e04b2cc1438e327f309fb47c1d8d3 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/hip commit: 8b25f7b4cf19bf2b4a0bb76eba61e659a8d89be9] --- projects/hip/bin/hipconfig | 23 ++++++++++++++++++----- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/projects/hip/bin/hipconfig b/projects/hip/bin/hipconfig index e70e6ae970..520a58cc2c 100755 --- a/projects/hip/bin/hipconfig +++ b/projects/hip/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 fa0c9d8a73f28356a568bb243daa9450b0d78b3e 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/hip commit: 026dc9457e2c910c6a0dbfd732af31d72c5e1092] --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 78 +++++++++++++------ 1 file changed, 56 insertions(+), 22 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp index 2d63133039..1617c0fd9e 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/hip/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 b19e4b5b63efeaaf70f0a602fc2a8a57136f3bd3. [ROCm/hip commit: 737972d31c92bedb590376f639532134bcf56a92] --- projects/hip/bin/hipcc | 6 ------ 1 file changed, 6 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 4c6a5ded07..b649bb6e7f 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/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 e748fd2ae02071e89678b31266b9b7c955f3fb41 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 17a853c277c0c29a9bf41afad0f27e1b59ecdfd8. [ROCm/hip commit: 1bd8d24278033a3ad8f1bf189c7ffca280c85598] --- projects/hip/bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index b649bb6e7f..d6eea8457b 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/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 d9f59b18f82f1570d5ee4d650c7d95ccbab4d5bd 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 1a8b50e63c8aadd2293a931c7071d983881b854c. [ROCm/hip commit: e8556388f9acfe8103576e8e120cac45315ef005] --- projects/hip/bin/hipcc | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index d6eea8457b..d227eea267 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/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 4093c03dd742efa4f55490302f11e2b9f45b2da0 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/hip commit: 77c6d5451fd156e62542d88a2d2dec7f3be0c800] --- projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp index 1617c0fd9e..5c8b1a18c1 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/hip/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 4a0d68ba3f1b50bbdd6b013979f8c2a6d8694cde 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/hip commit: 8d92d1ebd72a9bd4b0e72a3ca4fed4ae3d4dc8d9] --- projects/hip/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index ef5f80aee4..8e29bee68e 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/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 7a9ebee08abfcffd720d6f4190b8c28ccad492b0 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/hip commit: 131c0ea3261a1db3c6031ecb2394d59f644c6770] --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp index 5c8b1a18c1..bcb0adbeef 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/hip/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/hip commit: 7ebbbd35254be59b305bbca10e8a4b472eaedf30] --- .../include/hip/hcc_detail/hip_runtime_api.h | 28 +++++++++++ projects/hip/src/hip_memory.cpp | 48 +++++++++++++++---- 2 files changed, 66 insertions(+), 10 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 b6ae88729a..9c50ca4755 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/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/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 8e29bee68e..b194faa1ca 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/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 085c6c15b261745b5a585f3bb1ea461f06b06798 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/hip commit: bfde8a7fab262f23079a13a7639986332644d88a] --- .../tests/src/runtimeApi/memory/hipMemset.cpp | 61 ++++++++++++++++++- projects/hip/tests/src/test_common.cpp | 7 +++ projects/hip/tests/src/test_common.h | 1 + 3 files changed, 66 insertions(+), 3 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp index bcb0adbeef..8ac5beaf55 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/hip/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/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index bacbf35a22..8820381826 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/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 ea689f63ca7d6804ba66881f618be8ce032cbceb 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/hip commit: 0337b686ef68b701717af05278b701c3a3916fa4] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 02c4b7ee61..abce9a5fc6 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/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 da589e38edbd50cf7c3684dc25a653821774d2ab 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/hip commit: 5cbd28f29b7e55695023c9110d71eee0f2b4b834] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 3 ++- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 6 +++--- projects/hip/src/hip_memory.cpp | 2 +- 3 files changed, 6 insertions(+), 5 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 9c50ca4755..73996982d1 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/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/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index abce9a5fc6..a463d17e6d 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/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/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index b194faa1ca..d8a9bd5708 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/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 19b6b2b570af0184d878903d936ed4a1e9e29c1f 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/hip commit: 8db717c769b591bed5fb279e3259922ce964f977] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index a463d17e6d..add4c3f238 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/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 54c72c1079ceb9b5f5e5c5a76e36cf9efb7700f2 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/hip commit: 51f29b9cee9633b8f7e90f5eb616a9adae188d29] --- projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp index 8ac5beaf55..abff987437 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/hip/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/hip commit: a87d798d90828da5414d65f265ae0fe1311a2828] --- .../markdown/CUDA_Driver_API_functions_supported_by_HIP.md | 4 ++-- .../hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index dc22286d84..89ca2d74e2 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/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/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 8ad95cb537..8dfc429c4b 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/hip/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}},