diff --git a/bin/hipcc b/bin/hipcc index 4c6a5ded07..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,16 +423,6 @@ 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' ) { $arg = "--cuda-device-only"; } @@ -767,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"; -} elsif ($HIP_PLATFORM eq 'clang') { - $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"; 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"; } } 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}}, diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index b6ae88729a..73996982d1 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,24 @@ 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(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 02c4b7ee61..add4c3f238 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(hipDeviceptr_t devPtr, int value, size_t count) { + return hipCUResultTohipError(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(hipDeviceptr_t devPtr, int value, size_t count, + hipStream_t stream __dparm(0)) { + return hipCUResultTohipError(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)); } diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index ef5f80aee4..d8a9bd5708 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(); @@ -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(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; + + 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); diff --git a/tests/src/runtimeApi/memory/hipMemset.cpp b/tests/src/runtimeApi/memory/hipMemset.cpp index 2d63133039..abff987437 100644 --- a/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/tests/src/runtimeApi/memory/hipMemset.cpp @@ -26,43 +26,133 @@ 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 */ #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= 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;