From b7d8b4d148717e32e86a23ae75c0c913c8909a56 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 28 Oct 2019 23:22:27 +0300 Subject: [PATCH 01/23] [HIP][cmake] Simplify `UNIX` related code (the beginning) [REASONS] 1. Make OS-dependent code more clear and readable 2. To ease Windows support [ROCm/hip commit: d58b70d8a00513aa6b7752a08fb7c6f73cffca94] --- projects/hip/CMakeLists.txt | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 57369a9039..dc73d3c88c 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -163,11 +163,14 @@ if(NOT CMAKE_BUILD_TYPE) endif() # Determine HIP install path -if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT AND CMAKE_INSTALL_PREFIX MATCHES "/usr/local") +if (UNIX) + set(HIP_DEFAULT_INSTALL_PREFIX "/opt/rocm/hip") +endif() +if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) if(CMAKE_BUILD_TYPE MATCHES Debug) set(CMAKE_INSTALL_PREFIX ${CMAKE_CURRENT_SOURCE_DIR} CACHE PATH "Installation path for HIP" FORCE) elseif(CMAKE_BUILD_TYPE MATCHES Release) - set(CMAKE_INSTALL_PREFIX "/opt/rocm/hip" CACHE PATH "Installation path for HIP" FORCE) + set(CMAKE_INSTALL_PREFIX ${HIP_DEFAULT_INSTALL_PREFIX} CACHE PATH "Installation path for HIP" FORCE) else() message(FATAL_ERROR "Invalid CMAKE_BUILD_TYPE specified. Valid values are Debug and Release") endif() From cdf59d1de0a7fe4aec68246d9da8aa5ca56220df Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 4 Nov 2019 17:45:11 +0000 Subject: [PATCH 02/23] fix race condition in hipEventRecord [ROCm/hip commit: 6fcff0118ea4cbcacf13183bf320d321a3839770] --- projects/hip/src/hip_event.cpp | 65 +++++++++++----------------------- 1 file changed, 21 insertions(+), 44 deletions(-) diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index 73e296dae3..b74efeb4f1 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -109,51 +109,28 @@ hipError_t hipEventCreate(hipEvent_t* event) { hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream); - - hipError_t status; - if (event){ - auto ecd = event->locked_copyCrit(); - if( ecd._state != hipEventStatusUnitialized) { - stream = ihipSyncAndResolveStream(stream); - - if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) { - // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0 - // - // If default stream , then wait on all queues. - ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); - ctx->locked_syncDefaultStream(true, true); - - { - LockedAccessor_EventCrit_t eCrit(event->criticalData()); - eCrit->_eventData.marker(hc::completion_future()); // reset event - eCrit->_eventData._stream = stream; - eCrit->_eventData._timestamp = hc::get_system_ticks(); - eCrit->_eventData._state = hipEventStatusComplete; - } - status = hipSuccess; - } else { - // Record the event in the stream: - // Keep a copy outside the critical section so we lock stream first, then event - to - // avoid deadlock - hc::completion_future cf = stream->locked_recordEvent(event); - - { - LockedAccessor_EventCrit_t eCrit(event->criticalData()); - eCrit->_eventData.marker(cf); - eCrit->_eventData._stream = stream; - eCrit->_eventData._timestamp = 0; - eCrit->_eventData._state = hipEventStatusRecording; - } - - status = hipSuccess; - } - } else { - status = hipErrorInvalidResourceHandle; - } - } else { - status = hipErrorInvalidResourceHandle; + if (!event) return ihipLogStatus(hipErrorInvalidResourceHandle); + stream = ihipSyncAndResolveStream(stream); + LockedAccessor_EventCrit_t eCrit(event->criticalData()); + if (eCrit->_eventData._state == hipEventStatusUnitialized) return ihipLogStatus(hipErrorInvalidResourceHandle); + if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) { + // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0 + // If default stream , then wait on all queues. + ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true, true); + eCrit->_eventData.marker(hc::completion_future()); // reset event + eCrit->_eventData._stream = stream; + eCrit->_eventData._timestamp = hc::get_system_ticks(); + eCrit->_eventData._state = hipEventStatusComplete; } - return ihipLogStatus(status); + else { + // Record the event in the stream: + eCrit->_eventData.marker(stream->locked_recordEvent(event)); + eCrit->_eventData._stream = stream; + eCrit->_eventData._timestamp = 0; + eCrit->_eventData._state = hipEventStatusRecording; + } + return ihipLogStatus(hipSuccess); } From 2ad99a6326b48ec7f300ed54a1c9d357c7a3e20e Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Tue, 5 Nov 2019 02:02:46 -0500 Subject: [PATCH 03/23] [hip][tests] Add two more workgroup sizes for testing hipLaunchCooperativeKernel (#1613) [ROCm/hip commit: cf92fae9e6999e9c0e0c935151c00e49d83a3747] --- .../src/runtimeApi/module/hipLaunchCooperativeKernel.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp index 89d003ea94..c76685fa89 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp @@ -108,11 +108,11 @@ int main() { dim3 dimBlock = dim3(1); dim3 dimGrid = dim3(1); int numBlocks = 0; - uint workgroups[2] = {32, 64}; + uint workgroups[4] = {32, 64, 128, 256}; system_clock::time_point start = system_clock::now(); - for (uint i = 0; i < 2; ++i) { + for (uint i = 0; i < 4; ++i) { dimBlock.x = workgroups[i]; // Calculate the device occupancy to know how many blocks can be run concurrently From 1b8973d6bee55bf06c7ea18ce688a200492e9759 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 4 Nov 2019 23:02:59 -0800 Subject: [PATCH 04/23] Deprecate HIP Markers (#1622) * Deprecate HIP markers * Deprecate profiler start/stop [ROCm/hip commit: 54fab7c35c0cafe35d7045038d2d6af556e70ff1] --- projects/hip/CMakeLists.txt | 3 +++ projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- projects/hip/include/hip/hip_profile.h | 3 +++ 3 files changed, 8 insertions(+), 2 deletions(-) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 7aeb415bcc..5d1f785f1b 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -184,7 +184,10 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER) set(COMPILE_HIP_ATP_MARKER 0) else() set(COMPILE_HIP_ATP_MARKER $ENV{COMPILE_HIP_ATP_MARKER}) + message(WARNING "HIP Markers are deprecated, please use roctracer/rocTX marker APIs.") endif() +else() + message(WARNING "HIP Markers are deprecated, please use roctracer/rocTX marker APIs.") endif() add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER) 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 5a61219f4a..7c11eaf0bf 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2994,7 +2994,7 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, * When using this API, start the profiler with profiling disabled. (--startdisabled) * @warning : hipProfilerStart API is under development. */ -hipError_t hipProfilerStart(); +hipError_t hipProfilerStart() __attribute__((deprecated("use roctracer/rocTX instead"))); /** @@ -3002,7 +3002,7 @@ hipError_t hipProfilerStart(); * When using this API, start the profiler with profiling disabled. (--startdisabled) * @warning : hipProfilerStop API is under development. */ -hipError_t hipProfilerStop(); +hipError_t hipProfilerStop() __attribute__((deprecated("use roctracer/rocTX instead"))); /** diff --git a/projects/hip/include/hip/hip_profile.h b/projects/hip/include/hip/hip_profile.h index 7474839258..95224af4a3 100644 --- a/projects/hip/include/hip/hip_profile.h +++ b/projects/hip/include/hip/hip_profile.h @@ -23,11 +23,14 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_PROFILE_H #define HIP_INCLUDE_HIP_HIP_PROFILE_H +#warning "HIP Profiling through markers is deprecated, please check roctrace/rocTX support." + #if not defined(ENABLE_HIP_PROFILE) #define ENABLE_HIP_PROFILE 1 #endif #if defined(__HIP_PLATFORM_HCC__) and (ENABLE_HIP_PROFILE == 1) +#warning "HIP Markers are deprecated and would be removed soon." #include #define HIP_SCOPED_MARKER(markerName, group) \ amdtScopedMarker __scopedMarker(markerName, group, nullptr); From 5a9c8168c101f2f0f208f4dfbe936d558346706f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 5 Nov 2019 14:00:13 +0300 Subject: [PATCH 05/23] [HIPIFY][#1409] Fix for kernel launch macro expansion + Add a corresponding test kernel_launch_01.cu + Add isBefore() check to avoid crash on Replacement with negative length TODO: + Compatibility with former LLVM versions + More complicated kernel launch tests [ROCm/hip commit: 976f8e8bf6da407a260950879e66d023362cf388] --- .../hip/hipify-clang/src/HipifyAction.cpp | 28 +++++++---- .../kernel_launch/kernel_launch_01.cu | 46 +++++++++++++++++++ 2 files changed, 64 insertions(+), 10 deletions(-) create mode 100644 projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 510d91978a..0e776df63e 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -198,7 +198,7 @@ clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRa // If the range is contained within a macro, update the macro definition. // Otherwise, use the file location and hope for the best. if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) { - return {SM.getFileLoc(begin), SM.getFileLoc(end)}; + return {SM.getExpansionLoc(begin), SM.getExpansionLoc(end)}; } return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; } @@ -390,16 +390,24 @@ bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result) OS << readSourceText(*SM, {argStart, argEnd}); } OS << ")"; - clang::SourceRange replacementRange = getWriteRange(*SM, {llcompat::getBeginLoc(launchKernel), llcompat::getEndLoc(launchKernel)}); - clang::SourceLocation launchStart = replacementRange.getBegin(); + clang::SourceLocation launchKernelExprLocBeg = launchKernel->getExprLoc(); + clang::SourceLocation launchKernelExprLocEnd = launchKernelExprLocBeg.isMacroID() ? SM->getExpansionRange(launchKernelExprLocBeg).getEnd() : llcompat::getEndLoc(launchKernel); + clang::SourceLocation launchKernelEnd = llcompat::getEndLoc(launchKernel); + clang::BeforeThanCompare isBefore(*SM); + launchKernelExprLocEnd = isBefore(launchKernelEnd, launchKernelExprLocEnd) ? launchKernelExprLocEnd : launchKernelEnd; + clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernelExprLocBeg, launchKernelExprLocEnd}); + clang::SourceLocation launchBeg = replacementRange.getBegin(); clang::SourceLocation launchEnd = replacementRange.getEnd(); - size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart); - ct::Replacement Rep(*SM, launchStart, length, OS.str()); - clang::FullSourceLoc fullSL(launchStart, *SM); - insertReplacement(Rep, fullSL); - hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str()); - return true; + if (isBefore(launchBeg, launchEnd)) { + size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchBeg); + ct::Replacement Rep(*SM, launchBeg, length, OS.str()); + clang::FullSourceLoc fullSL(launchBeg, *SM); + insertReplacement(Rep, fullSL); + hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str()); + return true; + } + return false; } bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) { diff --git a/projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu b/projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu new file mode 100644 index 0000000000..3795d3c799 --- /dev/null +++ b/projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu @@ -0,0 +1,46 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones. +// ToDo: change lit testing in order to parse the output. + +#define LEN 1024 +#define SIZE LEN * sizeof(float) +#define ITER 1024*1024 + +// CHECK: #include +#include + +#define CUDA_LAUNCH(cuda_call,dimGrid,dimBlock, ...) \ + cuda_call<<>>(__VA_ARGS__); + +__global__ void Inc1(float *Ad, float *Bd) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if (tx < 1) { + for (int i = 0; i < ITER; ++i) { + Ad[tx] = Ad[tx] + 1.0f; + for (int j = 0; j < 256; ++j) { + Bd[tx] = Ad[tx]; + } + } + } +} + +int main() { + float *A, *Ad, *Bd; + A = new float[LEN]; + for (int i = 0; i < LEN; ++i) { + A[i] = 0.0f; + } + // CHECK: hipError_t status; + cudaError_t status; + // CHECK: status = hipHostRegister(A, SIZE, hipHostRegisterMapped); + status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped); + // CHECK: hipHostGetDevicePointer(&Ad, A, 0); + cudaHostGetDevicePointer(&Ad, A, 0); + // CHECK: hipMalloc((void**)&Bd, SIZE); + cudaMalloc((void**)&Bd, SIZE); + dim3 dimGrid(LEN / 512, 1, 1); + dim3 dimBlock(512, 1, 1); + + // CHECK: hipLaunchKernelGGL(Inc1, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd); + CUDA_LAUNCH(Inc1, dimGrid, dimBlock, Ad, Bd); +} From e250f761838061eaf0a3a771f58dfae69691857e Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Tue, 5 Nov 2019 11:51:00 -0500 Subject: [PATCH 06/23] Use portable macro for deprecation message. [ROCm/hip commit: a7f311cc145f3553c4f607dd49a2182bcd3aca53] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 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 7c11eaf0bf..d955c078dd 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2994,7 +2994,8 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, * When using this API, start the profiler with profiling disabled. (--startdisabled) * @warning : hipProfilerStart API is under development. */ -hipError_t hipProfilerStart() __attribute__((deprecated("use roctracer/rocTX instead"))); +DEPRECATED("use roctracer/rocTX instead") +hipError_t hipProfilerStart(); /** @@ -3002,7 +3003,8 @@ hipError_t hipProfilerStart() __attribute__((deprecated("use roctracer/rocTX ins * When using this API, start the profiler with profiling disabled. (--startdisabled) * @warning : hipProfilerStop API is under development. */ -hipError_t hipProfilerStop() __attribute__((deprecated("use roctracer/rocTX instead"))); +DEPRECATED("use roctracer/rocTX instead") +hipError_t hipProfilerStop(); /** From 1baaf44179ebc2b20cbef7bb3416077b10e26721 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 6 Nov 2019 02:43:04 +0200 Subject: [PATCH 07/23] __half2 should walk like CUDA and talk like CUDA [ROCm/hip commit: e5bd00d06b4bdcbc6912a78ef913cef7ae51ed95] --- projects/hip/include/hip/hcc_detail/hip_fp16.h | 18 ++++++------------ 1 file changed, 6 insertions(+), 12 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 18ad566121..23ad461065 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -1126,8 +1126,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __half2_raw{__builtin_convertvector(-r, _Float16_2)}; } inline __device__ @@ -1135,8 +1134,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data != static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1144,8 +1142,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data <= static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1153,8 +1150,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data >= static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1162,8 +1158,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ @@ -1171,8 +1166,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data; - return __half2_raw{_Float16_2{ - static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ From f8506bb669d964d9b09293ddeea23af00b0dec78 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 6 Nov 2019 02:46:21 +0200 Subject: [PATCH 08/23] Remove leftover noise. [ROCm/hip commit: b9faa9f8ae72962cd0ea48bbccf65c81134f34ea] --- projects/hip/include/hip/hcc_detail/hip_fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 23ad461065..52abc1a004 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -1126,7 +1126,7 @@ THE SOFTWARE. { auto r = static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data; - return __half2_raw{__builtin_convertvector(-r, _Float16_2)}; + return __builtin_convertvector(-r, _Float16_2); } inline __device__ From 1776ff55d365553727beaf5aac7e1e5666e3c54e Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 6 Nov 2019 14:30:39 +0300 Subject: [PATCH 09/23] [HIPIFY] Add a compatibility for CharSourceRange if LLVM < 7 [ROCm/hip commit: c87ab05fc6dee07fa0d9d9bb7345fd838c973fb2] --- projects/hip/hipify-clang/src/HipifyAction.cpp | 2 +- projects/hip/hipify-clang/src/LLVMCompat.cpp | 8 ++++++++ projects/hip/hipify-clang/src/LLVMCompat.h | 2 ++ 3 files changed, 11 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 0e776df63e..032bfef2aa 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -391,7 +391,7 @@ bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result) } OS << ")"; clang::SourceLocation launchKernelExprLocBeg = launchKernel->getExprLoc(); - clang::SourceLocation launchKernelExprLocEnd = launchKernelExprLocBeg.isMacroID() ? SM->getExpansionRange(launchKernelExprLocBeg).getEnd() : llcompat::getEndLoc(launchKernel); + clang::SourceLocation launchKernelExprLocEnd = launchKernelExprLocBeg.isMacroID() ? llcompat::getEndOfExpansionRangeForLoc(*SM, launchKernelExprLocBeg) : llcompat::getEndLoc(launchKernel); clang::SourceLocation launchKernelEnd = llcompat::getEndLoc(launchKernel); clang::BeforeThanCompare isBefore(*SM); launchKernelExprLocEnd = isBefore(launchKernelEnd, launchKernelExprLocEnd) ? launchKernelExprLocEnd : launchKernelEnd; diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index daebd214d1..1d1119f206 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp @@ -151,4 +151,12 @@ bool CheckCompatibility() { return true; } +clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager& SM, const clang::SourceLocation& loc) { +#if LLVM_VERSION_MAJOR > 6 + return SM.getExpansionRange(loc).getEnd(); +#else + return SM.getExpansionRange(loc).second; +#endif +} + } // namespace llcompat diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h index c1d671f513..8ad2862d9f 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.h +++ b/projects/hip/hipify-clang/src/LLVMCompat.h @@ -91,4 +91,6 @@ void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI); bool CheckCompatibility(); +clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager& SM, const clang::SourceLocation& loc); + } // namespace llcompat From 1b6e4b495abcf063119a7aba77863c6766affa49 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 6 Nov 2019 14:43:22 +0300 Subject: [PATCH 10/23] [HIPIFY][format] Apply clang style formatting in LLVMCompat [ROCm/hip commit: 3b76fd0b1cce975dbcafaedb297915ca7033ba9a] --- projects/hip/hipify-clang/src/LLVMCompat.cpp | 16 ++++++++-------- projects/hip/hipify-clang/src/LLVMCompat.h | 16 ++++++++-------- 2 files changed, 16 insertions(+), 16 deletions(-) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index 1d1119f206..b5aea95dca 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp @@ -40,7 +40,7 @@ void PrintStackTraceOnErrorSignal() { #endif } -ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file) { +ct::Replacements &getReplacements(ct::RefactoringTool &Tool, StringRef file) { #if LLVM_VERSION_MAJOR > 3 // getReplacements() now returns a map from filename to Replacements - so create an entry // for this source file and return a reference to it. @@ -50,7 +50,7 @@ ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file) { #endif } -void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep) { +void insertReplacement(ct::Replacements &replacements, const ct::Replacement &rep) { #if LLVM_VERSION_MAJOR > 3 // New clang added error checking to Replacements, and *insists* that you explicitly check it. llvm::consumeError(replacements.add(rep)); @@ -60,7 +60,7 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re #endif } -void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) { +void EnterPreprocessorTokenStream(clang::Preprocessor &_pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) { #if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) _pp.EnterTokenStream(start, len, false, DisableMacroExpansion); #else @@ -72,7 +72,7 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token * #endif } -clang::SourceLocation getBeginLoc(const clang::Stmt* stmt) { +clang::SourceLocation getBeginLoc(const clang::Stmt *stmt) { #if LLVM_VERSION_MAJOR < 8 return stmt->getLocStart(); #else @@ -80,7 +80,7 @@ clang::SourceLocation getBeginLoc(const clang::Stmt* stmt) { #endif } -clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc) { +clang::SourceLocation getBeginLoc(const clang::TypeLoc &typeLoc) { #if LLVM_VERSION_MAJOR < 8 return typeLoc.getLocStart(); #else @@ -88,7 +88,7 @@ clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc) { #endif } -clang::SourceLocation getEndLoc(const clang::Stmt* stmt) { +clang::SourceLocation getEndLoc(const clang::Stmt *stmt) { #if LLVM_VERSION_MAJOR < 8 return stmt->getLocEnd(); #else @@ -96,7 +96,7 @@ clang::SourceLocation getEndLoc(const clang::Stmt* stmt) { #endif } -clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc) { +clang::SourceLocation getEndLoc(const clang::TypeLoc &typeLoc) { #if LLVM_VERSION_MAJOR < 8 return typeLoc.getLocEnd(); #else @@ -151,7 +151,7 @@ bool CheckCompatibility() { return true; } -clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager& SM, const clang::SourceLocation& loc) { +clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager &SM, const clang::SourceLocation &loc) { #if LLVM_VERSION_MAJOR > 6 return SM.getExpansionRange(loc).getEnd(); #else diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h index 8ad2862d9f..e857023752 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.h +++ b/projects/hip/hipify-clang/src/LLVMCompat.h @@ -49,11 +49,11 @@ namespace llcompat { #define LLVM_DEBUG(X) DEBUG(X) #endif -clang::SourceLocation getBeginLoc(const clang::Stmt* stmt); -clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc); +clang::SourceLocation getBeginLoc(const clang::Stmt *stmt); +clang::SourceLocation getBeginLoc(const clang::TypeLoc &typeLoc); -clang::SourceLocation getEndLoc(const clang::Stmt* stmt); -clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc); +clang::SourceLocation getEndLoc(const clang::Stmt *stmt); +clang::SourceLocation getEndLoc(const clang::TypeLoc &typeLoc); void PrintStackTraceOnErrorSignal(); @@ -65,17 +65,17 @@ using namespace llvm; * Older LLVM versions don't actually support multiple filenames, so everything all gets * smushed together. It is the caller's responsibility to cope with this. */ -ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file); +ct::Replacements &getReplacements(ct::RefactoringTool &Tool, StringRef file); /** * Add a Replacement to a Replacements. */ -void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep); +void insertReplacement(ct::Replacements &replacements, const ct::Replacement &rep); /** * Version-agnostic version of Preprocessor::EnterTokenStream(). */ -void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, +void EnterPreprocessorTokenStream(clang::Preprocessor &_pp, const clang::Token *start, size_t len, bool DisableMacroExpansion); @@ -91,6 +91,6 @@ void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI); bool CheckCompatibility(); -clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager& SM, const clang::SourceLocation& loc); +clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager &SM, const clang::SourceLocation &loc); } // namespace llcompat From 8f033af43477a7728702e12a59da64bc26638cbc Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 5 Nov 2019 16:21:59 +0000 Subject: [PATCH 11/23] remove ihipEvent_t::refreshEventStatus(), new hipEventElapsedTime This fixes possible races in hipEventElapsedTime. [ROCm/hip commit: 7986090d9c3e5c6018537110b6ec30d8d110cdd9] --- projects/hip/src/hip_event.cpp | 115 ++++++++++++---------------- projects/hip/src/hip_hcc_internal.h | 2 - 2 files changed, 49 insertions(+), 68 deletions(-) diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index b74efeb4f1..7e84971813 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -44,28 +44,22 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future* cf, hipS } -std::pair ihipEvent_t::refreshEventStatus() { - auto ecd = locked_copyCrit(); - if (ecd._state == hipEventStatusRecording) { - bool isReady1 = ecd._stream->locked_eventIsReady(this); - if (isReady1) { - LockedAccessor_EventCrit_t eCrit(_criticalData); - - if ((eCrit->_eventData._type == hipEventTypeIndependent) || - (eCrit->_eventData._type == hipEventTypeStopCommand)) { - eCrit->_eventData._timestamp = eCrit->_eventData.marker().get_end_tick(); - } else if (eCrit->_eventData._type == hipEventTypeStartCommand) { - eCrit->_eventData._timestamp = eCrit->_eventData.marker().get_begin_tick(); - } else { - eCrit->_eventData._timestamp = 0; - assert(0); // TODO - move to debug assert - } - - eCrit->_eventData._state = hipEventStatusComplete; - - return std::pair(eCrit->_eventData._state, - eCrit->_eventData._timestamp); +static std::pair refreshEventStatus(ihipEventData_t &ecd) { + if (ecd._state == hipEventStatusRecording && ecd.marker().is_ready()) { + if ((ecd._type == hipEventTypeIndependent) || + (ecd._type == hipEventTypeStopCommand)) { + ecd._timestamp = ecd.marker().get_end_tick(); + } else if (ecd._type == hipEventTypeStartCommand) { + ecd._timestamp = ecd.marker().get_begin_tick(); + } else { + ecd._timestamp = 0; + assert(0); // TODO - move to debug assert } + + ecd._state = hipEventStatusComplete; + + return std::pair(ecd._state, + ecd._timestamp); } // Not complete path here: @@ -180,60 +174,49 @@ hipError_t hipEventSynchronize(hipEvent_t event) { hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) { HIP_INIT_API(hipEventElapsedTime, ms, start, stop); - hipError_t status = hipSuccess; + if (ms == nullptr) return ihipLogStatus(hipErrorInvalidValue); + if ((start == nullptr) || (stop == nullptr)) return ihipLogStatus(hipErrorInvalidResourceHandle); - if (ms == nullptr) { - status = hipErrorInvalidValue; + *ms = 0.0f; + auto startEcd = start->locked_copyCrit(); + auto stopEcd = stop->locked_copyCrit(); + + if ((start->_flags & hipEventDisableTiming) || + (startEcd._state == hipEventStatusUnitialized) || + (startEcd._state == hipEventStatusCreated) || + (stop->_flags & hipEventDisableTiming) || + (stopEcd._state == hipEventStatusUnitialized) || + (stopEcd._state == hipEventStatusCreated)) { + // Both events must be at least recorded else return hipErrorInvalidResourceHandle + return ihipLogStatus(hipErrorInvalidResourceHandle); } - else if ((start == nullptr) || (stop == nullptr)) { - status = hipErrorInvalidResourceHandle; - } else { - *ms = 0.0f; - auto startEcd = start->locked_copyCrit(); - auto stopEcd = stop->locked_copyCrit(); - if ((start->_flags & hipEventDisableTiming) || - (startEcd._state == hipEventStatusUnitialized) || - (startEcd._state == hipEventStatusCreated) || (stop->_flags & hipEventDisableTiming) || - (stopEcd._state == hipEventStatusUnitialized) || - (stopEcd._state == hipEventStatusCreated)) { - // Both events must be at least recorded else return hipErrorInvalidResourceHandle + // Refresh status, if still recording... - status = hipErrorInvalidResourceHandle; + auto startStatus = refreshEventStatus(startEcd); // pair < state, timestamp > + auto stopStatus = refreshEventStatus(stopEcd); // pair < state, timestamp > + if ((startStatus.first == hipEventStatusComplete) && + (stopStatus.first == hipEventStatusComplete)) { + // Common case, we have good information for both events. 'second' is the timestamp: + int64_t tickDiff = (stopStatus.second - startStatus.second); + uint64_t freqHz; + hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz); + if (freqHz) { + *ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f; + return ihipLogStatus(hipSuccess); } else { - // Refresh status, if still recording... - - auto startStatus = start->refreshEventStatus(); // pair < state, timestamp > - auto stopStatus = stop->refreshEventStatus(); // pair < state, timestamp > - - if ((startStatus.first == hipEventStatusComplete) && - (stopStatus.first == hipEventStatusComplete)) { - // Common case, we have good information for both events. 'second" is the - // timestamp: - int64_t tickDiff = (stopStatus.second - startStatus.second); - - uint64_t freqHz; - hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz); - if (freqHz) { - *ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f; - status = hipSuccess; - } else { - *ms = 0.0f; - status = hipErrorInvalidValue; - } - - - } else if ((startStatus.first == hipEventStatusRecording) || - (stopStatus.first == hipEventStatusRecording)) { - status = hipErrorNotReady; - } else { - assert(0); - } + *ms = 0.0f; + return ihipLogStatus(hipErrorInvalidValue); } + } else if ((startStatus.first == hipEventStatusRecording) || + (stopStatus.first == hipEventStatusRecording)) { + return ihipLogStatus(hipErrorNotReady); + } else { + assert(0); // TODO should we return hipErrorUnknown ? } - return ihipLogStatus(status); + return ihipLogStatus(hipSuccess); } hipError_t hipEventQuery(hipEvent_t event) { diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 658cfbf576..8c1827535c 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -724,8 +724,6 @@ class ihipEvent_t { explicit ihipEvent_t(unsigned flags); void attachToCompletionFuture(const hc::completion_future* cf, hipStream_t stream, ihipEventType_t eventType); - std::pair refreshEventStatus(); // returns pair - // Return a copy of the critical state. The critical data is locked during the copy. ihipEventData_t locked_copyCrit() { From f9ad5643808091c14529a7fbb73c301f4808db76 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 6 Nov 2019 15:56:32 +0000 Subject: [PATCH 12/23] hipEventRecord only needs one lock; remove locked_eventIsReady [ROCm/hip commit: 85080905c0594170834ecf6b37654f73f4d546b3] --- projects/hip/src/hip_event.cpp | 32 ++++++++++++----------------- projects/hip/src/hip_hcc.cpp | 8 -------- projects/hip/src/hip_hcc_internal.h | 2 -- 3 files changed, 13 insertions(+), 29 deletions(-) diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index 7e84971813..dcab059774 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -222,25 +222,19 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) { hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event); - hipError_t status = hipSuccess; + if (!event) return ihipLogStatus(hipErrorInvalidResourceHandle); - if ( NULL == event) - { - status = hipErrorInvalidResourceHandle; - } else { - if (!(event->_flags & hipEventReleaseToSystem)) { - tprintf(DB_WARN, - "hipEventQuery on event without system-scope fence ; consider creating with " - "hipEventReleaseToSystem\n"); - } - - auto ecd = event->locked_copyCrit(); - - if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) { - status = hipErrorNotReady; - } else { - status = hipSuccess; - } + if (!(event->_flags & hipEventReleaseToSystem)) { + tprintf(DB_WARN, + "hipEventQuery on event without system-scope fence ; consider creating with " + "hipEventReleaseToSystem\n"); } - return ihipLogStatus(status); + + auto ecd = event->locked_copyCrit(); + + if (ecd._state == hipEventStatusRecording && !ecd.marker().is_ready()) { + return ihipLogStatus(hipErrorNotReady); + } + + return ihipLogStatus(hipSuccess); } diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 7150140e12..5c64eee317 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -325,14 +325,6 @@ void ihipStream_t::locked_streamWaitEvent(ihipEventData_t& ecd) { } -// Causes current stream to wait for specified event to complete: -// Note this does not provide any kind of host serialization. -bool ihipStream_t::locked_eventIsReady(hipEvent_t event) { - LockedAccessor_EventCrit_t ecrit(event->criticalData()); - - return (ecrit->_eventData.marker().is_ready()); -} - // Create a marker in this stream. // Save state in the event so it can track the status of the event. hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event) { diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 8c1827535c..8ee47eba4a 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -583,8 +583,6 @@ class ihipStream_t { void locked_streamWaitEvent(ihipEventData_t& event); hc::completion_future locked_recordEvent(hipEvent_t event); - bool locked_eventIsReady(hipEvent_t event); - ihipStreamCritical_t& criticalData() { return _criticalData; }; //--- From 7b216ca51c61fc47b6a95e772d136b005e088377 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 6 Nov 2019 19:18:13 +0300 Subject: [PATCH 13/23] [HIPIFY][doc] Update README.md + Supported versions, testing [ROCm/hip commit: 96483d0de561b5f803a906a83e7242a18536469a] --- projects/hip/hipify-clang/README.md | 156 ++++++++++++++-------------- 1 file changed, 79 insertions(+), 77 deletions(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index da5abc19da..e8c3d0807b 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -146,9 +146,9 @@ To run it: * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option: - - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38` + - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.5.32` - - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38` + - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.5.32` 5. Ensure [`CUB`](https://github.com/NVlabs/cub) of the version corresponding to CUDA's version is installed. @@ -190,9 +190,9 @@ To run it: On Linux the following configurations are tested: -Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.4.38 +Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.5.32 -Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.4.38 +Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.5.32 Minimum build system requirements for the above configurations: @@ -207,7 +207,7 @@ cmake -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \ -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \ - -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38 \ + -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.5.32 \ -DCUDA_CUB_ROOT_DIR=/srv/git/CUB \ -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \ .. @@ -264,74 +264,76 @@ Linux 5.2.0 - Platform OS 64 - hipify-clang binary bitness 64 - python 2.7.12 binary bitness ======================================== --- Testing: 65 tests, 12 threads -- -PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 65) -PASS: hipify :: unit_tests/device/math_functions.cu (2 of 65) -PASS: hipify :: unit_tests/device/atomics.cu (3 of 65) -PASS: hipify :: unit_tests/device/device_symbols.cu (4 of 65) -PASS: hipify :: unit_tests/headers/headers_test_02.cu (5 of 65) -PASS: hipify :: unit_tests/headers/headers_test_03.cu (6 of 65) -PASS: hipify :: unit_tests/headers/headers_test_01.cu (7 of 65) -PASS: hipify :: unit_tests/headers/headers_test_04.cu (8 of 65) -PASS: hipify :: unit_tests/headers/headers_test_05.cu (9 of 65) -PASS: hipify :: unit_tests/headers/headers_test_07.cu (10 of 65) -PASS: hipify :: unit_tests/headers/headers_test_06.cu (11 of 65) -PASS: hipify :: unit_tests/headers/headers_test_11.cu (12 of 65) -PASS: hipify :: unit_tests/headers/headers_test_08.cu (13 of 65) -PASS: hipify :: unit_tests/headers/headers_test_10.cu (14 of 65) -PASS: hipify :: unit_tests/headers/headers_test_09.cu (15 of 65) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (16 of 65) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (17 of 65) -PASS: hipify :: unit_tests/libraries/CUB/cub_01.cu (18 of 65) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (19 of 65) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (20 of 65) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (21 of 65) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (22 of 65) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (23 of 65) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (24 of 65) -PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (25 of 65) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (26 of 65) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (27 of 65) -PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (28 of 65) -PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (29 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (30 of 65) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (31 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (32 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (33 of 65) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (34 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (35 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (36 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (37 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (38 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (39 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (40 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (41 of 65) -PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (42 of 65) -PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (43 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (44 of 65) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (45 of 65) -PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (46 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (47 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (48 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (49 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (50 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (52 of 65) -PASS: hipify :: unit_tests/samples/MallocManaged.cpp (53 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (54 of 65) -PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (55 of 65) -PASS: hipify :: unit_tests/samples/allocators.cu (56 of 65) -PASS: hipify :: unit_tests/samples/coalescing.cu (57 of 65) -PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (58 of 65) -PASS: hipify :: unit_tests/samples/axpy.cu (59 of 65) -PASS: hipify :: unit_tests/samples/cudaRegister.cu (60 of 65) -PASS: hipify :: unit_tests/samples/intro.cu (61 of 65) -PASS: hipify :: unit_tests/samples/square.cu (62 of 65) -PASS: hipify :: unit_tests/samples/static_shared_memory.cu (63 of 65) -PASS: hipify :: unit_tests/samples/vec_add.cu (64 of 65) -PASS: hipify :: unit_tests/libraries/CUB/cub_02.cu (18 of 65) -Testing Time: 3.01s - Expected Passes : 65 +-- Testing: 67 tests, 12 threads -- +PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 67) +PASS: hipify :: unit_tests/device/math_functions.cu (2 of 67) +PASS: hipify :: unit_tests/device/atomics.cu (3 of 67) +PASS: hipify :: unit_tests/device/device_symbols.cu (4 of 67) +PASS: hipify :: unit_tests/headers/headers_test_01.cu (5 of 67) +PASS: hipify :: unit_tests/headers/headers_test_02.cu (6 of 67) +PASS: hipify :: unit_tests/headers/headers_test_03.cu (7 of 67) +PASS: hipify :: unit_tests/headers/headers_test_05.cu (8 of 67) +PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 67) +PASS: hipify :: unit_tests/headers/headers_test_06.cu (10 of 67) +PASS: hipify :: unit_tests/headers/headers_test_07.cu (11 of 67) +PASS: hipify :: unit_tests/headers/headers_test_10.cu (12 of 67) +PASS: hipify :: unit_tests/headers/headers_test_11.cu (13 of 67) +PASS: hipify :: unit_tests/headers/headers_test_08.cu (14 of 67) +PASS: hipify :: unit_tests/kernel_launch/kernel_launch_01.cu (15 of 67) +PASS: hipify :: unit_tests/headers/headers_test_09.cu (16 of 67) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (17 of 67) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (18 of 67) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (19 of 67) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (20 of 67) +PASS: hipify :: unit_tests/libraries/CUB/cub_03.cu (21 of 67) +PASS: hipify :: unit_tests/libraries/CUB/cub_01.cu (22 of 67) +PASS: hipify :: unit_tests/libraries/CUB/cub_02.cu (23 of 67) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (24 of 67) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (25 of 67) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (26 of 67) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (27 of 67) +PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (28 of 67) +PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (29 of 67) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (30 of 67) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (31 of 67) +PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (32 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (33 of 67) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (34 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (35 of 67) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (36 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (37 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (38 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (39 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (40 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (41 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (42 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (43 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (44 of 67) +PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (45 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (46 of 67) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (47 of 67) +PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (48 of 67) +PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (49 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (50 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (51 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (52 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (53 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (54 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (55 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (56 of 67) +PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (57 of 67) +PASS: hipify :: unit_tests/samples/MallocManaged.cpp (58 of 67) +PASS: hipify :: unit_tests/samples/allocators.cu (59 of 67) +PASS: hipify :: unit_tests/samples/coalescing.cu (60 of 67) +PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (61 of 67) +PASS: hipify :: unit_tests/samples/axpy.cu (62 of 67) +PASS: hipify :: unit_tests/samples/intro.cu (63 of 67) +PASS: hipify :: unit_tests/samples/cudaRegister.cu (64 of 67) +PASS: hipify :: unit_tests/samples/square.cu (65 of 67) +PASS: hipify :: unit_tests/samples/static_shared_memory.cu (66 of 67) +PASS: hipify :: unit_tests/samples/vec_add.cu (67 of 67) +Testing Time: 3.07s + Expected Passes : 67 [100%] Built target test-hipify ``` ### Windows @@ -340,13 +342,13 @@ On Windows 10 the following configurations are tested: LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn 5.1.10 - 7.1.4.18 -LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn 7.0.5.15 - 7.6.4.38 +LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn 7.0.5.15 - 7.6.5.32 -LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn 7.0.5.15 - 7.6.4.38 +LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn 7.0.5.15 - 7.6.5.32 Build system requirements for the latest configuration LLVM 9.0.0/CUDA 10.1 Update 2: -Python 3.6.0 - 3.8.0, cmake 3.5.1 - 3.15.5, Visual Studio 2017 (15.5.2) - 2019 (16.3.5). +Python 3.6.0 - 3.8.0, cmake 3.5.1 - 3.16.0, Visual Studio 2017 (15.5.2) - 2019 (16.3.8). Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`: @@ -360,7 +362,7 @@ cmake -DCMAKE_PREFIX_PATH=f:/LLVM/9.0.0/dist \ -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \ -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \ - -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38 \ + -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.5.32 \ -DCUDA_CUB_ROOT_DIR=f:/GIT/cub \ -DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \ -Thost=x64 From 282c76c26f2acfc03f915830f46954d5f6faf22d Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 6 Nov 2019 19:25:42 +0300 Subject: [PATCH 14/23] [HIPIFY][Win][fix] canCompileHostAndDeviceInOneJob is true only for LLVM >= 10 [ROCm/hip commit: 9ca4e966415cf2f8bb0e62bb6132c7057179a774] --- projects/hip/hipify-clang/src/LLVMCompat.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index b5aea95dca..b524c41613 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp @@ -128,7 +128,7 @@ bool pragma_once_outside_header() { } bool canCompileHostAndDeviceInOneJob() { -#if LLVM_VERSION_MAJOR >= 9 && defined(_WIN32) +#if LLVM_VERSION_MAJOR > 9 && defined(_WIN32) return true; #else return false; From 6968362d996a376f5570502c7cb31bc4a86ec867 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 6 Nov 2019 23:47:10 -0800 Subject: [PATCH 15/23] Rename hip/hip_hcc.h to hip/hip_ext.h (#1341) * Rename hip/hip_hcc.h to hip/hip_ext.h * Deprecate hip_hcc.h [ROCm/hip commit: 579a4f36fa765b3eecb5cc04074a613fa6924840] --- .../hcc_detail/macro_based_grid_launch.hpp | 4 +- projects/hip/include/hip/hip_ext.h | 118 ++++++++++++++++++ projects/hip/include/hip/hip_hcc.h | 99 +-------------- .../0_Intro/module_api/launchKernelHcc.cpp | 2 +- .../0_Intro/module_api_global/runKernel.cpp | 2 +- .../11_texture_driver/texture2dDrv.cpp | 2 - projects/hip/src/hip_hcc.cpp | 2 +- projects/hip/src/hip_module.cpp | 2 +- projects/hip/tests/src/hipHcc.cpp | 2 +- .../runtimeApi/module/hipModuleGetGlobal.cpp | 2 +- .../module/hipModuleTexture2dDrv.cpp | 2 - 11 files changed, 128 insertions(+), 109 deletions(-) create mode 100644 projects/hip/include/hip/hip_ext.h diff --git a/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp index ca52e5b614..96d449b213 100644 --- a/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp +++ b/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "helpers.hpp" #include "hc.hpp" -#include "hip/hip_hcc.h" +#include "hip/hip_ext.h" #include "hip_runtime.h" #include @@ -795,4 +795,4 @@ requires(Domain == {Ts...}) inline std::enable_if_t< hipLaunchKernelGGL(kernel_name, num_blocks, dim_blocks, group_mem_bytes, stream, \ hipLaunchParm{}, ##__VA_ARGS__); \ } while (0) -} // namespace hip_impl \ No newline at end of file +} // namespace hip_impl diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h new file mode 100644 index 0000000000..3df911ee7a --- /dev/null +++ b/projects/hip/include/hip/hip_ext.h @@ -0,0 +1,118 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_INCLUDE_HIP_HIP_EXT_H +#define HIP_INCLUDE_HIP_HIP_EXT_H +#include "hip/hip_runtime_api.h" +#ifdef __HCC__ + +// Forward declarations: +namespace hc { +class accelerator; +class accelerator_view; +}; // namespace hc + + +/** + *------------------------------------------------------------------------------------------------- + *------------------------------------------------------------------------------------------------- + * @defgroup HCC-specific features + * @warning These APIs provide access to special features of HCC compiler and are not available + *through the CUDA path. + * @{ + */ + + +/** + * @brief Return hc::accelerator associated with the specified deviceId + * @return #hipSuccess, #hipErrorInvalidDevice + */ +HIP_PUBLIC_API +hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc); + +/** + * @brief Return hc::accelerator_view associated with the specified stream + * + * If stream is 0, the accelerator_view for the default stream is returned. + * @return #hipSuccess + */ +HIP_PUBLIC_API +hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av); + +#endif // #ifdef __HCC__ + + +/** + * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed + to kernelparams or extra + * + * @param [in[ f Kernel to launch. + * @param [in] gridDimX X grid dimension specified in work-items + * @param [in] gridDimY Y grid dimension specified in work-items + * @param [in] gridDimZ Z grid dimension specified in work-items + * @param [in] blockDimX X block dimensions specified in work-items + * @param [in] blockDimY Y grid dimension specified in work-items + * @param [in] blockDimZ Z grid dimension specified in work-items + * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The + kernel can access this with HIP_DYNAMIC_SHARED. + * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th + default stream is used with associated synchronization rules. + * @param [in] kernelParams + * @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and + must be in the memory layout and alignment expected by the kernel. + * @param [in] startEvent If non-null, specified event will be updated to track the start time of + the kernel launch. The event must be created before calling this API. + * @param [in] stopEvent If non-null, specified event will be updated to track the stop time of + the kernel launch. The event must be created before calling this API. + * + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + * + * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please + refer to hip_porting_driver_api.md for sample usage. + + * HIP/ROCm actually updates the start event when the associated kernel completes. + */ +HIP_PUBLIC_API +hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, + uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, uint32_t localWorkSizeY, + uint32_t localWorkSizeZ, size_t sharedMemBytes, + hipStream_t hStream, void** kernelParams, void** extra, + hipEvent_t startEvent = nullptr, + hipEvent_t stopEvent = nullptr, + uint32_t flags = 0); + +HIP_PUBLIC_API +hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, + uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, uint32_t localWorkSizeY, + uint32_t localWorkSizeZ, size_t sharedMemBytes, + hipStream_t hStream, void** kernelParams, void** extra, + hipEvent_t startEvent = nullptr, + hipEvent_t stopEvent = nullptr) + __attribute__((deprecated("use hipExtModuleLaunchKernel instead"))); + +// doxygen end AMD-specific features +/** + * @} + */ +#endif // #ifdef HIP_INCLUDE_HIP_HIP_EXT_H diff --git a/projects/hip/include/hip/hip_hcc.h b/projects/hip/include/hip/hip_hcc.h index a82a8c622b..e7e27fc721 100644 --- a/projects/hip/include/hip/hip_hcc.h +++ b/projects/hip/include/hip/hip_hcc.h @@ -1,16 +1,13 @@ /* Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -22,98 +19,6 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_HCC_H #define HIP_INCLUDE_HIP_HIP_HCC_H - -#ifdef __HCC__ - -#include "hip/hip_runtime_api.h" - -// Forward declarations: -namespace hc { -class accelerator; -class accelerator_view; -}; // namespace hc - - -/** - *------------------------------------------------------------------------------------------------- - *------------------------------------------------------------------------------------------------- - * @defgroup HCC-specific features - * @warning These APIs provide access to special features of HCC compiler and are not available - *through the CUDA path. - * @{ - */ - - -/** - * @brief Return hc::accelerator associated with the specified deviceId - * @return #hipSuccess, #hipErrorInvalidDevice - */ -HIP_PUBLIC_API -hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc); - -/** - * @brief Return hc::accelerator_view associated with the specified stream - * - * If stream is 0, the accelerator_view for the default stream is returned. - * @return #hipSuccess - */ -HIP_PUBLIC_API -hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av); - - -/** - * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed - to kernelparams or extra - * - * @param [in[ f Kernel to launch. - * @param [in] gridDimX X grid dimension specified in work-items - * @param [in] gridDimY Y grid dimension specified in work-items - * @param [in] gridDimZ Z grid dimension specified in work-items - * @param [in] blockDimX X block dimensions specified in work-items - * @param [in] blockDimY Y grid dimension specified in work-items - * @param [in] blockDimZ Z grid dimension specified in work-items - * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The - kernel can access this with HIP_DYNAMIC_SHARED. - * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th - default stream is used with associated synchronization rules. - * @param [in] kernelParams - * @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and - must be in the memory layout and alignment expected by the kernel. - * @param [in] startEvent If non-null, specified event will be updated to track the start time of - the kernel launch. The event must be created before calling this API. - * @param [in] stopEvent If non-null, specified event will be updated to track the stop time of - the kernel launch. The event must be created before calling this API. - * - * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue - * - * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please - refer to hip_porting_driver_api.md for sample usage. - - * HIP/ROCm actually updates the start event when the associated kernel completes. - */ -HIP_PUBLIC_API -hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, - uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, - uint32_t localWorkSizeX, uint32_t localWorkSizeY, - uint32_t localWorkSizeZ, size_t sharedMemBytes, - hipStream_t hStream, void** kernelParams, void** extra, - hipEvent_t startEvent = nullptr, - hipEvent_t stopEvent = nullptr, - uint32_t flags = 0); - -HIP_PUBLIC_API -hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, - uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, - uint32_t localWorkSizeX, uint32_t localWorkSizeY, - uint32_t localWorkSizeZ, size_t sharedMemBytes, - hipStream_t hStream, void** kernelParams, void** extra, - hipEvent_t startEvent = nullptr, - hipEvent_t stopEvent = nullptr) - __attribute__((deprecated("use hipExtModuleLaunchKernel instead"))); - -// doxygen end HCC-specific features -/** - * @} - */ -#endif // #ifdef __HCC__ +#warning "hip/hip_hcc.h is deprecated, please use hip/hip_ext.h" +#include "hip/hip_ext.h" #endif // #ifdef HIP_INCLUDE_HIP_HIP_HCC_H diff --git a/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp index 4ff000d1af..38cf0d414c 100644 --- a/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -27,7 +27,7 @@ THE SOFTWARE. #include #ifdef __HIP_PLATFORM_HCC__ -#include +#include #endif #define LEN 64 diff --git a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp index 7aa7f7ba2d..3a2804b7a2 100644 --- a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include #include #include -#include +#include #define LEN 64 #define SIZE LEN * sizeof(float) diff --git a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 7c22e19d5d..b42ac86ad1 100755 --- a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -21,11 +21,9 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -//#include "hip/hip_runtime_api.h" #include #include #include -//#include #define fileName "tex2dKernel.code" diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 5c64eee317..afc57e4b64 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -46,7 +46,7 @@ THE SOFTWARE. #include "hsa/hsa_ext_image.h" #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" -#include "hip/hip_hcc.h" +#include "hip/hip_ext.h" #include "trace_helper.h" #include "env.h" diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 2afbabf0a8..498ee5bb03 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include "hip/hcc_detail/hsa_helpers.hpp" #include "hip/hcc_detail/program_state.hpp" #include "hip_hcc_internal.h" -#include "hip/hip_hcc.h" +#include "hip/hip_ext.h" #include "program_state.inl" #include "trace_helper.h" diff --git a/projects/hip/tests/src/hipHcc.cpp b/projects/hip/tests/src/hipHcc.cpp index e0a843d4f5..2d4523ecd6 100644 --- a/projects/hip/tests/src/hipHcc.cpp +++ b/projects/hip/tests/src/hipHcc.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" -#include "hip/hip_hcc.h" +#include "hip/hip_ext.h" #include "test_common.h" #define CHECK(error) \ diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp index 956ba4e50a..f98d5e4ec4 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #include #include #include -#include +#include #define LEN 64 #define SIZE LEN * sizeof(float) diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp index e678c35d6e..9ae5883608 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp @@ -27,11 +27,9 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -//#include "hip/hip_runtime_api.h" #include #include #include -//#include #define fileName "tex2d_kernel.code" From 92dcba11ace58dbcff9252a7d1afd907304037b6 Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Thu, 7 Nov 2019 13:17:46 +0530 Subject: [PATCH 16/23] Fixed texture 2D mapping for pitched arrays & 3D Texture read (#1415) Texture 2D image mapping for pitched arrays: github issue: Texture Object's Buffer seems to be Misaligned #886 JIRA ticket: SWDEV-199313 SWDEV-151670 : Fixed issue with 3D texture with 4 components SWDEV-151671 : Issue with 2D layered texture with 4 components [ROCm/hip commit: e07926ce0f485b3cc4c95ad6492206c749dc922d] --- .../include/hip/hcc_detail/hip_runtime_api.h | 2 +- projects/hip/src/hip_memory.cpp | 595 ++++++------------ projects/hip/src/hip_texture.cpp | 20 +- projects/hip/tests/src/test_common.h | 30 + .../tests/src/texture/hipBindTex2DPitch.cpp | 87 +++ .../texture/hipNormalizedFloatValueTex.cpp | 33 +- .../hip/tests/src/texture/hipTexObjPitch.cpp | 106 ++++ .../src/texture/simpleTexture2DLayered.cpp | 106 ++++ .../hip/tests/src/texture/simpleTexture3D.cpp | 133 ++++ 9 files changed, 698 insertions(+), 414 deletions(-) create mode 100644 projects/hip/tests/src/texture/hipBindTex2DPitch.cpp create mode 100644 projects/hip/tests/src/texture/hipTexObjPitch.cpp create mode 100644 projects/hip/tests/src/texture/simpleTexture2DLayered.cpp create mode 100644 projects/hip/tests/src/texture/simpleTexture3D.cpp 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 d955c078dd..ba7f2f4ade 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -3315,7 +3315,7 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, - size_t width, size_t height, textureReference* tex); + size_t width, size_t height, textureReference* tex, size_t pitch); template hipError_t hipBindTexture2D(size_t* offset, struct texture& tex, diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 009c2db505..3f82246cf9 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -429,58 +429,62 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { return hipHostMalloc(ptr, sizeBytes, flags); }; +hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width, int height, int depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType,void ** ptr, hsa_ext_image_data_info_t &imageInfo, int array_size __dparm(0)) { + auto ctx = ihipGetTlsDefaultCtx(); + if (ctx) { + hc::accelerator acc = ctx->getDevice()->_acc; + hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); + if (!agent) + return hipErrorInvalidResourceHandle; + size_t allocGranularity = 0; + hsa_amd_memory_pool_t* allocRegion = static_cast(acc.get_hsa_am_region()); + hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); + + hsa_ext_image_descriptor_t imageDescriptor; + imageDescriptor.geometry = geometry; + imageDescriptor.width = width; + imageDescriptor.height = height; + imageDescriptor.depth = depth; + imageDescriptor.array_size = array_size; + imageDescriptor.format.channel_order = channelOrder; + imageDescriptor.format.channel_type = channelType; + + hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + hsa_status_t status = + hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo); + if(imageInfo.size == 0 || HSA_STATUS_SUCCESS != status){ + return hipErrorRuntimeOther; + } + size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; + const unsigned am_flags = 0; + *ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx, + false /*shareWithAll*/, am_flags, 0, alignment); + if (*ptr == NULL) { + return hipErrorMemoryAllocation; + } + return hipSuccess; + } + else { + return hipErrorMemoryAllocation; + } +} + // width in bytes hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) { hipError_t hip_status = hipSuccess; - if(ptr==NULL || pitch == NULL) - { - hip_status=hipErrorInvalidValue; - return hip_status; - } - // hardcoded 128 bytes - *pitch = ((((int)width - 1) / 128) + 1) * 128; - const size_t sizeBytes = (*pitch) * height * ((depth==0) ? 1 : depth); - - auto ctx = ihipGetTlsDefaultCtx(); - - if (ctx) { - hc::accelerator acc = ctx->getDevice()->_acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - - size_t allocGranularity = 0; - hsa_amd_memory_pool_t* allocRegion = - static_cast(acc.get_hsa_am_region()); - hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &allocGranularity); - - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = *pitch; - imageDescriptor.height = height; - imageDescriptor.depth = depth; - imageDescriptor.array_size = 0; - if (depth == 0) - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; - else - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.format.channel_order = HSA_EXT_IMAGE_CHANNEL_ORDER_R; - imageDescriptor.format.channel_type = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; - - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_ext_image_data_info_t imageInfo; - hsa_status_t status = - hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); - size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; - - const unsigned am_flags = 0; - *ptr = hip_internal::allocAndSharePtr("device_pitch", sizeBytes, ctx, - false /*shareWithAll*/, am_flags, 0, alignment); - - if (sizeBytes && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - } else { - hip_status = hipErrorMemoryAllocation; + if(ptr==NULL || pitch == NULL){ + return hipErrorInvalidValue; } + hsa_ext_image_data_info_t imageInfo; + if (depth == 0) + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,width,height,0,HSA_EXT_IMAGE_CHANNEL_ORDER_R, + HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo); + else + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,width,height,depth,HSA_EXT_IMAGE_CHANNEL_ORDER_R, + HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo); + + if(hip_status == hipSuccess) + *pitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth); return hip_status; } @@ -541,18 +545,74 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, hsa_ext_image_channel_order_t* channelOrder, hsa_ext_image_channel_type_t* channelType); +hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height, int depth, hipChannelFormatDesc desc, hsa_ext_image_data_info_t &imageInfo,int array_size __dparm(0)) +{ + hsa_ext_image_descriptor_t imageDescriptor; + imageDescriptor.geometry = geometry; + imageDescriptor.width = width; + imageDescriptor.height = height; + imageDescriptor.depth = depth; + imageDescriptor.array_size = array_size; + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(desc, hipReadModeElementType, &channelOrder, &channelType); + imageDescriptor.format.channel_order = channelOrder; + imageDescriptor.format.channel_type = channelType; + + hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + // Get the current device agent. + hc::accelerator acc; + hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); + if (!agent) + return hipErrorInvalidResourceHandle; + hsa_status_t status = + hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo); + if(HSA_STATUS_SUCCESS != status){ + return hipErrorRuntimeOther; + } + + return hipSuccess; +} + +hipError_t ihipArrayToImageFormat(hipArray_Format format,hsa_ext_image_channel_type_t &channelType) { + switch (format) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; + break; + case HIP_AD_FORMAT_UNSIGNED_INT16: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; + break; + case HIP_AD_FORMAT_UNSIGNED_INT32: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; + break; + case HIP_AD_FORMAT_SIGNED_INT8: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; + break; + case HIP_AD_FORMAT_SIGNED_INT16: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; + break; + case HIP_AD_FORMAT_SIGNED_INT32: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32; + break; + case HIP_AD_FORMAT_HALF: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT; + break; + case HIP_AD_FORMAT_FLOAT: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT; + break; + default: + return hipErrorUnknown; + break; + } + return hipSuccess; +} + hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_SPECIAL_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; if (pAllocateArray->Width > 0) { - auto ctx = ihipGetTlsDefaultCtx(); *array = (hipArray*)malloc(sizeof(hipArray)); - HIP_ARRAY3D_DESCRIPTOR array3D; - array3D.Width = pAllocateArray->Width; - array3D.Height = pAllocateArray->Height; - array3D.Format = pAllocateArray->Format; - array3D.NumChannels = pAllocateArray->NumChannels; array[0]->width = pAllocateArray->Width; array[0]->height = pAllocateArray->Height; array[0]->Format = pAllocateArray->Format; @@ -560,100 +620,26 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat array[0]->isDrv = true; array[0]->textureType = hipTextureType2D; void** ptr = &array[0]->data; - if (ctx) { - const unsigned am_flags = 0; - size_t size = pAllocateArray->Width; - if (pAllocateArray->Height > 0) { - size = size * pAllocateArray->Height; - } - hsa_ext_image_channel_type_t channelType; - size_t allocSize = 0; - switch (pAllocateArray->Format) { - case HIP_AD_FORMAT_UNSIGNED_INT8: - allocSize = size * sizeof(uint8_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; - break; - case HIP_AD_FORMAT_UNSIGNED_INT16: - allocSize = size * sizeof(uint16_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; - break; - case HIP_AD_FORMAT_UNSIGNED_INT32: - allocSize = size * sizeof(uint32_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; - break; - case HIP_AD_FORMAT_SIGNED_INT8: - allocSize = size * sizeof(int8_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; - break; - case HIP_AD_FORMAT_SIGNED_INT16: - allocSize = size * sizeof(int16_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; - break; - case HIP_AD_FORMAT_SIGNED_INT32: - allocSize = size * sizeof(int32_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32; - break; - case HIP_AD_FORMAT_HALF: - allocSize = size * sizeof(int16_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT; - break; - case HIP_AD_FORMAT_FLOAT: - allocSize = size * sizeof(float); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT; - break; - default: - hip_status = hipErrorUnknown; - break; - } - hc::accelerator acc = ctx->getDevice()->_acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); + hsa_ext_image_channel_type_t channelType; + hsa_ext_image_channel_order_t channelOrder; - size_t allocGranularity = 0; - hsa_amd_memory_pool_t* allocRegion = - static_cast(acc.get_hsa_am_region()); - hsa_amd_memory_pool_get_info( - *allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); + hip_status = ihipArrayToImageFormat(pAllocateArray->Format,channelType); + if(hipSuccess != hip_status) + return ihipLogStatus(hip_status); - hsa_ext_image_descriptor_t imageDescriptor; - - imageDescriptor.width = pAllocateArray->Width; - imageDescriptor.height = pAllocateArray->Height; - imageDescriptor.depth = 0; - imageDescriptor.array_size = 0; - - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; - - hsa_ext_image_channel_order_t channelOrder; - - if (pAllocateArray->NumChannels == 4) { - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; - } else if (pAllocateArray->NumChannels == 2) { - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; - } else if (pAllocateArray->NumChannels == 1) { - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; - } - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; - - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_ext_image_data_info_t imageInfo; - hsa_status_t status = - hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); - size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; - - *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, - false /*shareWithAll*/, am_flags, 0, alignment); - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - } else { - hip_status = hipErrorMemoryAllocation; + if (pAllocateArray->NumChannels == 4) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; + } else if (pAllocateArray->NumChannels == 2) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; + } else if (pAllocateArray->NumChannels == 1) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; } + hsa_ext_image_data_info_t imageInfo; + return ihipLogStatus(allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,pAllocateArray->Width, + pAllocateArray->Height,0,channelOrder,channelType,ptr,imageInfo)); } else { - hip_status = hipErrorInvalidValue; + return ihipLogStatus(hipErrorInvalidValue); } - - return ihipLogStatus(hip_status); } hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, @@ -662,8 +648,6 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; if (width > 0) { - auto ctx = ihipGetTlsDefaultCtx(); - *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = flags; array[0]->width = width; @@ -674,67 +658,25 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si array[0]->textureType = hipTextureType2D; void** ptr = &array[0]->data; - if (ctx) { - const unsigned am_flags = 0; - size_t size = width; - if (height > 0) { - size = size * height; - } - - const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8); - - hc::accelerator acc = ctx->getDevice()->_acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - - size_t allocGranularity = 0; - hsa_amd_memory_pool_t* allocRegion = - static_cast(acc.get_hsa_am_region()); - hsa_amd_memory_pool_get_info( - *allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); - - hsa_ext_image_descriptor_t imageDescriptor; - - imageDescriptor.width = width; - imageDescriptor.height = height; - imageDescriptor.depth = 0; - imageDescriptor.array_size = 0; - switch (flags) { - case hipArrayLayered: - case hipArrayCubemap: - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - assert(0); - break; - case hipArrayDefault: - default: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; - break; - } - hsa_ext_image_channel_order_t channelOrder; - hsa_ext_image_channel_type_t channelType; - getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; - - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_ext_image_data_info_t imageInfo; - hsa_status_t status = - hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); - size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; - - *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, - false /*shareWithAll*/, am_flags, 0, alignment); - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - - } else { - hip_status = hipErrorMemoryAllocation; + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); + hsa_ext_image_data_info_t imageInfo; + switch (flags) { + case hipArrayLayered: + case hipArrayCubemap: + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + default: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,width,height,0,channelOrder,channelType,ptr,imageInfo); + break; } } else { hip_status = hipErrorInvalidValue; } - return ihipLogStatus(hip_status); } @@ -742,8 +684,6 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll HIP_INIT_SPECIAL_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray); hipError_t hip_status = hipSuccess; - auto ctx = ihipGetTlsDefaultCtx(); - *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = pAllocateArray->Flags; array[0]->width = pAllocateArray->Width; @@ -752,111 +692,37 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll array[0]->Format = pAllocateArray->Format; array[0]->NumChannels = pAllocateArray->NumChannels; array[0]->isDrv = true; - array[0]->textureType = hipTextureType3D; void** ptr = &array[0]->data; - if (ctx) { - const unsigned am_flags = 0; - const size_t size = pAllocateArray->Width * pAllocateArray->Height * pAllocateArray->Depth; + hsa_ext_image_channel_order_t channelOrder; + if (pAllocateArray->NumChannels == 4) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; + } else if (pAllocateArray->NumChannels == 2) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; + } else if (pAllocateArray->NumChannels == 1) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; + } - size_t allocSize = 0; - hsa_ext_image_channel_type_t channelType; - switch (pAllocateArray->Format) { - case HIP_AD_FORMAT_UNSIGNED_INT8: - allocSize = size * sizeof(uint8_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; - break; - case HIP_AD_FORMAT_UNSIGNED_INT16: - allocSize = size * sizeof(uint16_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; - break; - case HIP_AD_FORMAT_UNSIGNED_INT32: - allocSize = size * sizeof(uint32_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; - break; - case HIP_AD_FORMAT_SIGNED_INT8: - allocSize = size * sizeof(int8_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; - break; - case HIP_AD_FORMAT_SIGNED_INT16: - allocSize = size * sizeof(int16_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; - break; - case HIP_AD_FORMAT_SIGNED_INT32: - allocSize = size * sizeof(int32_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32; - break; - case HIP_AD_FORMAT_HALF: - allocSize = size * sizeof(int16_t); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT; - break; - case HIP_AD_FORMAT_FLOAT: - allocSize = size * sizeof(float); - channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT; - break; - default: - hip_status = hipErrorUnknown; - break; - } - - hc::accelerator acc = ctx->getDevice()->_acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - - size_t allocGranularity = 0; - hsa_amd_memory_pool_t* allocRegion = - static_cast(acc.get_hsa_am_region()); - hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &allocGranularity); - - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = pAllocateArray->Width; - imageDescriptor.height = pAllocateArray->Height; - imageDescriptor.depth = 0; - imageDescriptor.array_size = 0; - switch (pAllocateArray->Flags) { - case hipArrayLayered: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA; - imageDescriptor.array_size = pAllocateArray->Depth; - break; - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - case hipArrayDefault: - assert(0); - break; - case hipArrayCubemap: - default: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.depth = pAllocateArray->Depth; - break; - } - hsa_ext_image_channel_order_t channelOrder; - - // getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); - if (pAllocateArray->NumChannels == 4) { - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; - } else if (pAllocateArray->NumChannels == 2) { - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; - } else if (pAllocateArray->NumChannels == 1) { - channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; - } - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; - - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_ext_image_data_info_t imageInfo; - hsa_status_t status = - hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); - size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; - - *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0, - alignment); - - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - - } else { - hip_status = hipErrorMemoryAllocation; + hsa_ext_image_channel_type_t channelType; + hip_status = ihipArrayToImageFormat(pAllocateArray->Format,channelType); + hsa_ext_image_data_info_t imageInfo; + switch (pAllocateArray->Flags) { + case hipArrayLayered: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2DA,pAllocateArray->Width,pAllocateArray->Height,0, + channelOrder,channelType,ptr,imageInfo,pAllocateArray->Depth); + array[0]->textureType = hipTextureType2DLayered; + break; + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + case hipArrayCubemap: + default: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,pAllocateArray->Width,pAllocateArray->Height, + pAllocateArray->Depth,channelOrder,channelType,ptr,imageInfo); + array[0]->textureType = hipTextureType3D; + break; } return ihipLogStatus(hip_status); @@ -865,19 +731,13 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, struct hipExtent extent, unsigned int flags) { - - HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; - if(array==NULL ) - { - hip_status=hipErrorInvalidValue; - return ihipLogStatus(hip_status); + if(array==NULL ){ + return ihipLogStatus(hipErrorInvalidValue); } - auto ctx = ihipGetTlsDefaultCtx(); - *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = flags; array[0]->width = extent.width; @@ -885,68 +745,27 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* array[0]->depth = extent.depth; array[0]->desc = *desc; array[0]->isDrv = false; - array[0]->textureType = hipTextureType3D; void** ptr = &array[0]->data; - - if (ctx) { - const unsigned am_flags = 0; - const size_t size = extent.width * extent.height * extent.depth; - - const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8); - - hc::accelerator acc = ctx->getDevice()->_acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - - size_t allocGranularity = 0; - hsa_amd_memory_pool_t* allocRegion = - static_cast(acc.get_hsa_am_region()); - hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &allocGranularity); - - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = extent.width; - imageDescriptor.height = extent.height; - imageDescriptor.depth = extent.depth; - imageDescriptor.array_size = 0; - switch (flags) { - case hipArrayLayered: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA; - imageDescriptor.array_size = extent.depth; - break; - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - case hipArrayDefault: - assert(0); - break; - case hipArrayCubemap: - default: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.depth = extent.depth; - break; - } - hsa_ext_image_channel_order_t channelOrder; - hsa_ext_image_channel_type_t channelType; - getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); - imageDescriptor.format.channel_order = channelOrder; - imageDescriptor.format.channel_type = channelType; - - hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; - hsa_ext_image_data_info_t imageInfo; - hsa_status_t status = - hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); - size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; - - *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0, - alignment); - - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - - } else { - hip_status = hipErrorMemoryAllocation; + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); + hsa_ext_image_data_info_t imageInfo; + switch (flags) { + case hipArrayLayered: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2DA,extent.width,extent.height,0,channelOrder,channelType,ptr,imageInfo,extent.depth); + array[0]->textureType = hipTextureType2DLayered; + break; + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + case hipArrayCubemap: + default: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,extent.width,extent.height,extent.depth,channelOrder,channelType,ptr,imageInfo); + array[0]->textureType = hipTextureType3D; + break; } - return ihipLogStatus(hip_status); } @@ -1283,7 +1102,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); hc::completion_future marker; - try { stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false); } catch (ihipException& ex) { @@ -1452,15 +1270,9 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) { hipError_t e = hipSuccess; if(p) { - size_t byteSize; - size_t depth; - size_t height; - size_t widthInBytes; - size_t srcPitch; - size_t dstPitch; - void* srcPtr; - void* dstPtr; - size_t ySize; + size_t byteSize, width, height, depth, widthInBytes, srcPitch, dstPitch, ySize; + hipChannelFormatDesc desc; + void* srcPtr;void* dstPtr; if (p->dstArray != nullptr) { if (p->dstArray->isDrv == false) { switch (p->dstArray->desc.f) { @@ -1482,22 +1294,32 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo } depth = p->extent.depth; height = p->extent.height; + width = p->extent.width; widthInBytes = p->extent.width * byteSize; srcPitch = p->srcPtr.pitch; srcPtr = p->srcPtr.ptr; ySize = p->srcPtr.ysize; - dstPitch = p->dstArray->width * byteSize; + desc = p->dstArray->desc; dstPtr = p->dstArray->data; } else { depth = p->Depth; height = p->Height; widthInBytes = p->WidthInBytes; - dstPitch = p->dstArray->width * 4; + width = p->dstArray->width; + desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned); srcPitch = p->srcPitch; srcPtr = (void*)p->srcHost; ySize = p->srcHeight; dstPtr = p->dstArray->data; } + hsa_ext_image_data_info_t imageInfo; + if(hipTextureType2DLayered == p->dstArray->textureType) + GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth); + else + GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo); + + dstPitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth); + } else { // Non array destination depth = p->extent.depth; @@ -1509,6 +1331,7 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo ySize = p->srcPtr.ysize; dstPitch = p->dstPtr.pitch; } + stream = ihipSyncAndResolveStream(stream); hc::completion_future marker; try { diff --git a/projects/hip/src/hip_texture.cpp b/projects/hip/src/hip_texture.cpp index 8820878fc0..7ea8da8624 100644 --- a/projects/hip/src/hip_texture.cpp +++ b/projects/hip/src/hip_texture.cpp @@ -222,7 +222,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; void* devPtr = nullptr; - + size_t pitch = 0; switch (pResDesc->resType) { case hipResourceTypeArray: devPtr = pResDesc->res.array.array->data; @@ -279,6 +279,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou imageDescriptor.depth = 0; imageDescriptor.array_size = 0; imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; + pitch = pResDesc->res.pitch2D.pitchInBytes; getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, &channelOrder, &channelType); break; @@ -296,7 +297,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout( *agent, &imageDescriptor, devPtr, permission, - HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) || + HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) || HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { return ihipLogStatus(hipErrorRuntimeOther); @@ -456,7 +457,7 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, - size_t width, size_t height, textureReference* tex) { + size_t width, size_t height, textureReference* tex, size_t pitch) { hipError_t hip_status = hipSuccess; enum hipTextureAddressMode addressMode = tex->addressMode[0]; enum hipTextureFilterMode filterMode = tex->filterMode; @@ -504,7 +505,7 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout( *agent, &imageDescriptor, devPtr, permission, - HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) || + HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) || HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { return hipErrorRuntimeOther; @@ -522,8 +523,12 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d size_t pitch) { HIP_INIT_API(hipBindTexture2D, offset, tex, devPtr, desc, width, height, pitch); hipError_t hip_status = hipSuccess; + + //TODO: Fix when HSA accepts user defined pitch + if(pitch % 64) pitch =0; + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, offset, devPtr, - desc, width, height, tex); + desc, width, height, tex, pitch); return ihipLogStatus(hip_status); } @@ -797,7 +802,10 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT size_t offset; hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. + //TODO: Fix when HSA accepts user defined pitch + if(pitch % 64) pitch =0; + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, &offset, devPtr, - NULL, desc->Width, desc->Height, tex); + NULL, desc->Width, desc->Height, tex, pitch); return ihipLogStatus(hip_status); } diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index 67a8e5e60a..92528e38e3 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -153,6 +153,36 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg); unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N); +template // pointer type +void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) +{ + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + int offset = i*width*height + j*width + k; + if (hData[offset] != hOutputData[offset]) { + std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" << hOutputData[offset]<<" "; + failed("mistmatch at:%d %d %d",i,j,k); + } + } + } + } +} + +template +void checkArray(T input, T output, size_t height, size_t width) +{ + for(int i=0; i __global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { diff --git a/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp b/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp new file mode 100644 index 0000000000..f463942329 --- /dev/null +++ b/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp @@ -0,0 +1,87 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/*HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "test_common.h" + +#define SIZE_H 8 +#define SIZE_W 12 +#define TYPE_t float +texture tex; + +// texture object is a kernel argument +__global__ void texture2dCopyKernel( TYPE_t* dst) { + + int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; + if ( (x< SIZE_W) && (y< SIZE_H) ){ + dst[SIZE_W*y+x] = tex2D(tex, x, y); + } +} + +int main (void) +{ + TYPE_t* B; + TYPE_t* A; + TYPE_t* devPtrB; + TYPE_t* devPtrA; + + B = new TYPE_t[SIZE_H*SIZE_W]; + A = new TYPE_t[SIZE_H*SIZE_W]; + for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){ + A[i-1] = i; + } + + size_t devPitchA, tex_ofs; + HIPCHECK(hipMallocPitch((void**)&devPtrA, &devPitchA ,SIZE_W*sizeof(TYPE_t), SIZE_H)) ; + HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); + + tex.normalized = false; + #if defined(__HIP_PLATFORM_NVCC__) + + cudaError_t status = cudaBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, + SIZE_W, SIZE_H, devPitchA); + if (status != cudaSuccess) { + printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, cudaGetErrorString(status), + status,__FILE__, __LINE__, KNRM); + failed("API returned error code."); + } + #else + HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, + SIZE_W, SIZE_H, devPitchA)); + #endif + HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB); + hipDeviceSynchronize(); + HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost)); + + HipTest::checkArray(A, B, SIZE_H, SIZE_W); + delete []A; + delete []B; + hipFree(devPtrA); + hipFree(devPtrB); + passed(); +} diff --git a/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp b/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp index fd2e8ccf0f..0a9c879376 100644 --- a/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -26,15 +26,6 @@ THE SOFTWARE. * HIT_END */ -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if(error != hipSuccess) {\ - fprintf(stderr, "error: '%s' (%d) at %s:%d\n", hipGetErrorString(error), error, __FILE__, __LINE__);\ - exit(EXIT_FAILURE);\ - }\ -} -#include "hip/hip_runtime.h" #include "test_common.h" #define SIZE 10 @@ -54,31 +45,31 @@ bool textureTest(enum hipArray_Format texFormat) { T hData[] = {65, 66, 67, 68, 69, 70, 71, 72,73,74}; T *dData = NULL; - CHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE)); - CHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE)); + HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE)); + HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE)); textureReference* texRef = &textureNormalizedVal_1D; - CHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp)); - CHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp)); - CHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint)); - CHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES)); - CHECK(hipTexRefSetFormat(texRef, texFormat, 1)); + HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp)); + HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp)); + HIPCHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint)); + HIPCHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES)); + HIPCHECK(hipTexRefSetFormat(texRef, texFormat, 1)); HIP_ARRAY_DESCRIPTOR desc; desc.Width = SIZE; desc.Height = 1; desc.Format = texFormat; desc.NumChannels = 1; - CHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE)); + HIPCHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE)); bool testResult = true; float *dOutputData = NULL; - CHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE)); + HIPCHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE)); hipLaunchKernelGGL(HIP_KERNEL_NAME(normalizedValTextureTest), dim3(1,1,1), dim3(SIZE,1,1), 0, 0, SIZE, dOutputData); float *hOutputData = new float[SIZE]; - CHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE))); + HIPCHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE))); for(int i = 0; i < SIZE; i++) { @@ -100,9 +91,9 @@ int main(int argc, char** argv) { int device = 0; bool status = true; - CHECK(hipSetDevice(device)); + HIPCHECK(hipSetDevice(device)); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, device)); + HIPCHECK(hipGetDeviceProperties(&props, device)); std::cout << "Device :: " << props.name << std::endl; #ifdef __HIP_PLATFORM_HCC__ std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl; diff --git a/projects/hip/tests/src/texture/hipTexObjPitch.cpp b/projects/hip/tests/src/texture/hipTexObjPitch.cpp new file mode 100644 index 0000000000..47648d5e73 --- /dev/null +++ b/projects/hip/tests/src/texture/hipTexObjPitch.cpp @@ -0,0 +1,106 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/*HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "test_common.h" + +#define SIZE_H 20 +#define SIZE_W 179 +// texture object is a kernel argument +template +__global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) { + + for(int i =0;i(texObj, j, i); + __syncthreads(); +} + +template +void texture2Dtest() +{ + TYPE_t* B; + TYPE_t* A; + TYPE_t* devPtrB; + TYPE_t* devPtrA; + + B = new TYPE_t[SIZE_H*SIZE_W]; + A = new TYPE_t[SIZE_H*SIZE_W]; + for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){ + A[i-1] = i; + } + + size_t devPitchA; + HIPCHECK(hipMallocPitch((void**)&devPtrA, &devPitchA ,SIZE_W*sizeof(TYPE_t), SIZE_H)) ; + HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); + + // Use the texture object + hipResourceDesc texRes; + hipMemset(&texRes, 0, sizeof(texRes)); + texRes.resType = hipResourceTypePitch2D; + texRes.res.pitch2D.devPtr = devPtrA; + texRes.res.pitch2D.height = SIZE_H; + texRes.res.pitch2D.width = SIZE_W; + texRes.res.pitch2D.pitchInBytes = devPitchA; + texRes.res.pitch2D.desc = hipCreateChannelDesc(); + + hipTextureDesc texDescr; + hipMemset(&texDescr, 0, sizeof(texDescr)); + texDescr.normalizedCoords = false; + texDescr.filterMode = hipFilterModePoint; + texDescr.mipmapFilterMode = hipFilterModePoint; + texDescr.addressMode[0] = hipAddressModeClamp; + texDescr.addressMode[1] = hipAddressModeClamp; + texDescr.addressMode[2] = hipAddressModeClamp; + texDescr.readMode = hipReadModeElementType; + + hipTextureObject_t texObj; + hipResourceViewDesc resDesc; + HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, &resDesc)); + + HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(1,1,1), dim3(1,1,1), 0, 0, + texObj, devPtrB, devPtrA); + + HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost)); + + HipTest::checkArray(A, B, SIZE_H, SIZE_W); + delete []A; + delete []B; + hipFree(devPtrA); + hipFree(devPtrB); +} + +int main() +{ + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + passed(); +} diff --git a/projects/hip/tests/src/texture/simpleTexture2DLayered.cpp b/projects/hip/tests/src/texture/simpleTexture2DLayered.cpp new file mode 100644 index 0000000000..e3ba6d9afe --- /dev/null +++ b/projects/hip/tests/src/texture/simpleTexture2DLayered.cpp @@ -0,0 +1,106 @@ +/* +Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "test_common.h" +typedef float T; +// Texture reference for 2D Layered texture +texture tex2DL; +__global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer) +{ + unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; + unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; + outputData[layer*width*height + y*width + x] = tex2DLayered(tex2DL, x, y, layer); +} + +//////////////////////////////////////////////////////////////////////////////// +void runTest(int width,int height,int num_layers,texture *tex) +{ + unsigned int size = width * height * num_layers * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + + for (unsigned int layer = 0; layer < num_layers; layer++){ + for (int i = 0; i < (int)(width * height); i++){ + hData[layer*width*height + i] = i; + } + } + hipChannelFormatDesc channelDesc; + // Allocate array and copy image data + channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindFloat); + hipArray *arr; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, num_layers), hipArrayLayered)); + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width, height, num_layers); + //myparms.kind = hipMemcpyHostToDevice; + HIPCHECK(hipMemcpy3D(&myparms)); + + // set texture parameters + tex->addressMode[0] = hipAddressModeWrap; + tex->addressMode[1] = hipAddressModeWrap; + tex->filterMode = hipFilterModePoint; + tex->normalized = false; + + // Bind the array to the texture + HIPCHECK(hipBindTextureToArray(*tex, arr, channelDesc)); + + // Allocate device memory for result + T* dData = NULL; + hipMalloc((void **) &dData, size); + dim3 dimBlock(8, 8, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + for (unsigned int layer = 0; layer < num_layers; layer++) + hipLaunchKernelGGL(simpleKernelLayeredArray, dimGrid, dimBlock, 0, 0, dData, width, height, layer); + + HIPCHECK(hipDeviceSynchronize()); + // Allocate mem for the result on host side + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + + // copy result from device to host + HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + HipTest::checkArray(hData,hOutputData,width,height,num_layers); + + hipFree(dData); + hipFreeArray(arr); + free(hData); + free(hOutputData); +} + +//////////////////////////////////////////////////////////////////////////////// +// Program main +//////////////////////////////////////////////////////////////////////////////// +int main(int argc, char **argv) +{ + runTest(512,512,5,&tex2DL); + passed(); +} + diff --git a/projects/hip/tests/src/texture/simpleTexture3D.cpp b/projects/hip/tests/src/texture/simpleTexture3D.cpp new file mode 100644 index 0000000000..741b1671af --- /dev/null +++ b/projects/hip/tests/src/texture/simpleTexture3D.cpp @@ -0,0 +1,133 @@ +/* +Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ +#include "test_common.h" + +//typedef char T; +const char *sampleName = "simpleTexture3D"; + +// Texture reference for 3D texture +texture texf; +texture texi; +texture texc; + +template +__global__ void simpleKernel3DArray(T* outputData, + int width, + int height,int depth) +{ + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + if(std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i); + else if(std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i); + else if(std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i); + } + } + } +} + +//////////////////////////////////////////////////////////////////////////////// +//! Run a simple test for tex3D +//////////////////////////////////////////////////////////////////////////////// +template +void runTest(int width,int height,int depth,texture *tex) +{ + unsigned int size = width * height * depth * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindSigned); + hipArray *arr; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayCubemap)); + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyHostToDevice; + HIPCHECK(hipMemcpy3D(&myparms)); + + // set texture parameters + tex->addressMode[0] = hipAddressModeWrap; + tex->addressMode[1] = hipAddressModeWrap; + tex->filterMode = hipFilterModePoint; + tex->normalized = false; + + // Bind the array to the texture + HIPCHECK(hipBindTextureToArray(*tex, arr, channelDesc)); + + // Allocate device memory for result + T* dData = NULL; + hipMalloc((void **) &dData, size); + + hipLaunchKernelGGL(simpleKernel3DArray, dim3(1,1,1), dim3(1,1,1), 0, 0, dData, width, height, depth); + HIPCHECK(hipDeviceSynchronize()); + + // Allocate mem for the result on host side + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + + // copy result from device to host + HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + HipTest::checkArray(hData,hOutputData,width,height,depth); + + hipFree(dData); + hipFreeArray(arr); + free(hData); + free(hOutputData); +} + +//////////////////////////////////////////////////////////////////////////////// +// Program main +//////////////////////////////////////////////////////////////////////////////// +int main(int argc, char **argv) +{ + printf("%s starting...\n", sampleName); + for(int i=1;i<25;i++) + { + runTest(i,i,i,&texf); + runTest(i+1,i,i,&texi); + runTest(i,i+1,i,&texc); + } + passed(); +} + From fa1e44aa0e201c34744da84348d949e3f85bf052 Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Thu, 7 Nov 2019 13:18:12 +0530 Subject: [PATCH 17/23] [HIP] Fixed hipStreamAddCallback [SWDEV#165185] (#1425) Fixed hipStreamAddCallback() as requested in SWDEV#165185 Added unit test to test the behavior [ROCm/hip commit: 45613311d72ad4c678d1e799aa0ac7dab8023e4d] --- projects/hip/src/hip_hcc.cpp | 15 +- projects/hip/src/hip_hcc_internal.h | 4 +- projects/hip/src/hip_stream.cpp | 127 ++++++++++++- .../runtimeApi/stream/StreamAddCallback.cpp | 145 ++++++++++++++ .../hipStreamAddCallbackCrossStream.cpp | 133 +++++++++++++ .../hipStreamAddCallbackMultiThread.cpp | 178 ++++++++++++++++++ 6 files changed, 593 insertions(+), 9 deletions(-) create mode 100644 projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp create mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp create mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index afc57e4b64..377ce0aad4 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -1497,18 +1497,21 @@ hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) { return e; } -void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) { +bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs) { hipError_t e = hipSuccess; - // Synchronize stream - tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n", - ToString(cb->_stream).c_str()); - GET_TLS(); - e = ihipStreamSynchronize(tls, cb->_stream); + ihipStreamCallback_t* cb = static_cast (cbArgs); + + if(cb->comFuture.valid()) + cb->comFuture.wait(); // Call registered callback function cb->_callback(cb->_stream, e, cb->_userData); + + hsa_signal_store_screlease(cb->_signal,0); + delete cb; + return false; } //--- diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 8ee47eba4a..6337edd370 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -650,7 +650,9 @@ class ihipStreamCallback_t { : _stream(stream), _callback(callback), _userData(userData) { }; hipStream_t _stream; + hsa_signal_t _signal; hipStreamCallback_t _callback; + hc::completion_future comFuture; void* _userData; }; @@ -968,7 +970,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0); hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream); -void ihipStreamCallbackHandler(ihipStreamCallback_t* cb); +bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs); // Stream printf functions: inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) { diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index e3e4975b7e..bf3289bb49 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -254,14 +254,137 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { //--- +void setCallbackPacket(hsa_queue_t* queue, + uint64_t& index, uint64_t& nextIndex, + hsa_barrier_and_packet_t** barrier1, + hsa_barrier_and_packet_t** barrier2){ + + uint64_t tempIndex = 0; + uint32_t mask = queue->size - 1; + hsa_barrier_and_packet_t* tempBarrier; + + // Check for empty packets + do{ + tempIndex = hsa_queue_load_write_index_scacquire(queue); + tempBarrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[tempIndex & mask]); + }while(!(tempBarrier->header & HSA_PACKET_TYPE_INVALID)); + + // Reserve two packets for two barriers + index = hsa_queue_add_write_index_scacquire(queue, 2); + + if(index > mask) + { + index = 0; + nextIndex = 1; + } + else if(index == mask) + nextIndex = 0; + else + nextIndex = index + 1; + + tempBarrier = new hsa_barrier_and_packet_t; + memset(tempBarrier, 0, sizeof(hsa_barrier_and_packet_t)); + tempBarrier->header = HSA_PACKET_TYPE_INVALID; + + // Barrier 1 + *barrier1 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index & mask]); + memcpy(*barrier1,tempBarrier,sizeof(hsa_barrier_and_packet_t)); + + // Barrier 2 + *barrier2 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[nextIndex & mask]); + memcpy(*barrier2,tempBarrier,sizeof(hsa_barrier_and_packet_t)); + + delete tempBarrier; +} + hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { + HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); hipError_t e = hipSuccess; - // Create a thread in detached mode to handle callback + if(stream == hipStreamNull) + { + ihipCtx_t* device = ihipGetTlsDefaultCtx(); + stream = device->_defaultStream; + } + + stream = ihipSyncAndResolveStream(stream); + + // Lock the stream + LockedAccessor_StreamCrit_t crit(stream->criticalData()); + + // Device synchronization + hc::completion_future marker = crit->_av.create_marker(hc::system_scope); + + // 1. Lock the queue + hsa_queue_t* lockedQ = static_cast (crit->_av.acquire_locked_hsa_queue()); + + if(lockedQ == nullptr) + { + // No queue attached to stream hence exiting early + return ihipLogStatus(hipErrorMissingConfiguration); + } + + // 2. Allocate a singals + hsa_signal_t signal; + hsa_status_t status = hsa_signal_create(1, 0, NULL, &signal); + + if(status != HSA_STATUS_SUCCESS) + { + crit->_av.release_locked_hsa_queue(); + return ihipLogStatus(hipErrorInvalidValue); + } + + hsa_signal_t depSignal; + status = hsa_signal_create(1, 0, NULL, &depSignal); + + if(status != HSA_STATUS_SUCCESS) + { + crit->_av.release_locked_hsa_queue(); + return ihipLogStatus(hipErrorInvalidValue); + } + + // 3. Store callback details, will destroy allocation in callback handler ihipStreamCallback_t* cb = new ihipStreamCallback_t(stream, callback, userData); - std::thread(ihipStreamCallbackHandler, cb).detach(); + if(cb == nullptr) + { + crit->_av.release_locked_hsa_queue(); + return ihipLogStatus(hipErrorMemoryAllocation); + } + cb->_signal = depSignal; + cb->comFuture = marker ; + + // 4. Create barrier packets + uint64_t index ; + uint64_t nextIndex; + + hsa_barrier_and_packet_t* barrier; + hsa_barrier_and_packet_t* depBarrier; + + setCallbackPacket(lockedQ, index, nextIndex, &barrier, &depBarrier); + + barrier->completion_signal = signal; + + depBarrier->dep_signal[0] = depSignal; + + uint16_t header = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE)| 1 << HSA_PACKET_HEADER_BARRIER; + + // 5. Update packet header, + // Intentionally updated second barrier header before first in order to avoid race + depBarrier->header = header; + barrier->header = header; + + // 6. Trigger the doorbell + nextIndex = nextIndex + 1; + hsa_queue_store_write_index_screlease(lockedQ, nextIndex); + hsa_signal_store_relaxed(lockedQ->doorbell_signal, index+1); + + // 7. Release queue + crit->_av.release_locked_hsa_queue(); + + // 8. Register signal callback + hsa_amd_signal_async_handler(signal, HSA_SIGNAL_CONDITION_EQ, 0, ihipStreamCallbackHandler, cb); return ihipLogStatus(e); } diff --git a/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp b/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp new file mode 100644 index 0000000000..e6492c7ce2 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp @@ -0,0 +1,145 @@ +#include +#include +#include +#include "test_common.h" +#include + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + +enum class ExecState +{ + EXEC_NOT_STARTED, + EXEC_STARTED, + EXEC_CB_STARTED, + EXEC_CB_FINISHED, + EXEC_FINISHED +}; + +struct UserData +{ + size_t size; + int* ptr; +}; + +// Global variable to check exection order +std::atomic gData(ExecState::EXEC_NOT_STARTED); + + +void myCallback(hipStream_t stream, hipError_t status, void* user_data) +{ + if(gData.load() != ExecState::EXEC_STARTED) + return; // Error hence return early + + gData.store(ExecState::EXEC_CB_STARTED); + + UserData* data = reinterpret_cast(user_data); + printf("Callback started\n"); + + sleep(1); + + printf("Callback ending.\n"); + gData.store(ExecState::EXEC_CB_FINISHED); +} + +bool test(int count) +{ + printf("\n============ Test iteration %d =============\n",count); + // Stream + hipStream_t stream; + bool result = true; + + gData.store(ExecState::EXEC_STARTED); + + HIPCHECK(hipStreamCreate(&stream)); + + // Array size + size_t size = 10000; + + // Device array + int *data = NULL; + HIPCHECK(hipMalloc((void**)&data, sizeof(int) * size)); + + // Initialize device array to -1 + HIPCHECK(hipMemset(data, -1, sizeof(int) * size)); + + // Host array + int *host = NULL; + HIPCHECK(hipHostMalloc((void**)&host, sizeof(int) * size)); + + // Print host ptr address + printf("In main thread\n"); + + // Initialize user_data for callback + UserData arg; + arg.size = size; + arg.ptr = host; + + // Synchronize device + HIPCHECK(hipDeviceSynchronize()); + + // Asynchronous copy from device to host + HIPCHECK(hipMemcpyAsync(host, data, sizeof(int) * size, hipMemcpyDeviceToHost, stream)); + + // Asynchronous memset on device + HIPCHECK(hipMemsetAsync(data, 0, sizeof(int) * size, stream)); + + // Add callback - should happen after hipMemsetAsync() + HIPCHECK(hipStreamAddCallback(stream, myCallback, &arg, 0)); + + printf("Will wait in main thread until callback completes\n"); + + //This should synchronize the stream (including the callback) + HIPCHECK(hipStreamSynchronize(stream)); + + if(gData.load() != ExecState::EXEC_CB_FINISHED) + { + std::cout<<"Callback is not finished\n"; + return false; + } + printf("Callback completed will resume main thread execution\n"); + + if(host[size/2] != -1) + { + // Print some host data that just got copied + printf("Pseudo host data printing (should be -1): %d\n", host[size/2]); + result = false; + } + + HIPCHECK(hipMemcpy(host, data, sizeof(int)*size, hipMemcpyDeviceToHost)); + + if(host[size-1] != 0) + { + printf("Pseudo host data printing (should be 0): %d\n", host[size-1]); + result = false; + } + + HIPCHECK(hipFree(data)); + HIPCHECK(hipHostFree(host)); + HIPCHECK(hipStreamDestroy(stream)); + + gData.store(ExecState::EXEC_FINISHED); + return result; +} + +int main() +{ + // Test involves multithreading hence running multiple times + // to make sure consitency in the behavior + bool status = true; + + for(int i=0; i < 10; i++){ + status = test(i+1); + if(status == false) + { + failed("Test Failed!\n"); + break; + } + } + + if(status == true) passed(); + return 0; +} diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp new file mode 100644 index 0000000000..b2fb168fbc --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp @@ -0,0 +1,133 @@ +/* + Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END +*/ +#include "test_common.h" +#include +#include +#include + + +//Globals +const int workloadCount = 1000000; +std::mutex gMutx; +bool callbackCompleted = false; +std::condition_variable condVar; + +// Device function +__global__ void increment(int *data,int N) +{ + int i = blockIdx.x*blockDim.x + threadIdx.x; + + if(i < N) + data[i] = 1 + data[i]; +} + +struct USER_DATA +{ + int* result; // Data received from device + int* copyOfOriginalData; // Copy of initial data which will be used for validation +}; + +// Callback +void callback(hipStream_t event, hipError_t status, void *userData) +{ + USER_DATA *data = (USER_DATA *)userData; + + if(!(data == nullptr || data->result == nullptr || data->copyOfOriginalData == nullptr)) + { + for(int i=0;iresult[i] != data->copyOfOriginalData[i]+1) + { + std::cout<<"Error value : "<result[i]<<"| Expected value :"<copyOfOriginalData[i]+1<result = hResultData; + inputParam->copyOfOriginalData = hData; + + HIPCHECK(hipStreamCreate(&stream)); + HIPCHECK(hipStreamAddCallback(stream,callback,inputParam,0)); + + // Wait for stream add callback to complete + std::unique_lock l(gMutx); + + while(!callbackCompleted) + condVar.wait(l); + + // Will destroy device memory hence explicite hipFree is not needed + HIPCHECK(hipDeviceReset()); + + free(hData); + free(hResultData); + passed(); +} diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp new file mode 100644 index 0000000000..3babf5869e --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp @@ -0,0 +1,178 @@ +/* + Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END +*/ + +#include +#include +#include +#include +#include "test_common.h" + +// Will indicate completion of callback added as part of hipStreamAddCallback +struct signal +{ + int completedThreads; + std::mutex mu; + std::condition_variable cv; +}; + +struct workload +{ + int _workloadId; + int _deviceID; + + int *copyOf_hData; // copy of host data which will be used to validation + int *hData; // will contain host data + int *dData; // device data will be stored + hipStream_t _stream;// stream on which data will be processed + + bool success; // start will be stored +}; + +// Global data +int numWorkloads = 8; +const int perWorkloadSize = 1000000; +signal completionSignal; + +// Kernel run on device +__global__ void increment(int *data, int N) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < N) + data[i] = data[i]+1; +} + +/* + * Method validates processed data array with saved copy and notifies via conditional variable to all waiting threads +*/ +void Analyze(hipStream_t event, hipError_t status, void *data) +{ + HIPCHECK(status); + + workload *W = (workload *) data; + + if(W != NULL) + { + W->success = true; + + for (int i=0; i< perWorkloadSize; ++i) + { + W->success &= (W->copyOf_hData[i] == (W->hData[i]+1)); + if(!W->success) + { + std::cout<<"\nExpected Data :"<<(W->hData[i]+1)<<" Current Data :"<copyOf_hData[i]< guard(completionSignal.mu); + completionSignal.completedThreads += 1; + + completionSignal.cv.notify_all(); +} + +/* + * Thread routine to launch workloads into separate stream + */ +void LaunchWorkload(void *args) +{ + workload *W = (workload *) args; + + if (W == nullptr) return; + + std::srand(std::time(nullptr)); + + HIPCHECK(hipSetDevice(W->_deviceID)); + + // Allocate memory + HIPCHECK(hipMalloc(&W->dData,perWorkloadSize*sizeof(int))); + + size_t s = perWorkloadSize*sizeof(int); + HIPCHECK(hipMemset(W->dData,0,s)) + + W->hData = (int *) malloc(perWorkloadSize*sizeof(int)); + W->copyOf_hData = (int *) malloc(perWorkloadSize*sizeof(int)); + + // Initialize host array + for(int i =0;ihData[i] = W->copyOf_hData[i] = std::rand() % perWorkloadSize; + } + + HIPCHECK(hipStreamCreate(&W->_stream)); + + dim3 block(256); + dim3 grid((perWorkloadSize + block.x-1) / block.x); + + HIPCHECK(hipMemcpyAsync(W->dData,W->hData,perWorkloadSize*sizeof(int),hipMemcpyHostToDevice,W->_stream)); + + hipLaunchKernelGGL(increment, grid, block,0, W->_stream,W->dData, perWorkloadSize); + + HIPCHECK(hipMemcpyAsync(W->copyOf_hData,W->dData,perWorkloadSize*sizeof(int),hipMemcpyDeviceToHost,W->_stream)); + + HIPCHECK(hipStreamAddCallback(W->_stream, Analyze, W,0)) +} + +int main(int argc, char* argv[]) { + + int numDevice = 0; + + HIPCHECK(hipGetDeviceCount(&numDevice)); + + std::thread workerThread[numWorkloads]; + + workload *workloads; + workloads = (workload *) malloc(numWorkloads * sizeof(workload)); + + for(int i =0;i l(completionSignal.mu); + while(completionSignal.completedThreads != numWorkloads) + { + completionSignal.cv.wait(l); + } + + //clean-up + hipDeviceReset(); + free(workloads); + + passed(); + return 0; +} From 68b4fbd0436edd44c0c0b9f2cd8f7f0e9207d5e2 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 7 Nov 2019 09:49:14 +0200 Subject: [PATCH 18/23] Remove native vector support from the GCC case, since it never worked (#1637) [ROCm/hip commit: 5530c15cc340e35f2bf6a668130e18eeb5f821ae] --- .../include/hip/hcc_detail/hip_vector_types.h | 49 ++++--------------- 1 file changed, 10 insertions(+), 39 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 75c79d422c..3a10a06372 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -37,11 +37,6 @@ THE SOFTWARE. #if !defined(_MSC_VER) || __clang__ #if defined(__clang__) #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) -#elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax. - #define __ROUND_UP_TO_NEXT_POT__(x) \ - (1 << (31 - __builtin_clz(x) + (x > (1 << (31 - __builtin_clz(x)))))) - #define __NATIVE_VECTOR__(n, T) \ - __attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T)))) #endif #if defined(__cplusplus) && defined(__clang__) @@ -932,47 +927,23 @@ THE SOFTWARE. using CUDA_name##4 = HIP_vector_type; #else #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ - typedef T CUDA_name##_impl1 __NATIVE_VECTOR__(1, T);\ - typedef T CUDA_name##_impl2 __NATIVE_VECTOR__(2, T);\ - typedef T CUDA_name##_impl3 __NATIVE_VECTOR__(3, T);\ - typedef T CUDA_name##_impl4 __NATIVE_VECTOR__(4, T);\ typedef struct {\ - union {\ - CUDA_name##_impl1 data;\ - struct {\ - T x;\ - };\ - };\ + T x;\ } CUDA_name##1;\ typedef struct {\ - union {\ - CUDA_name##_impl2 data;\ - struct {\ - T x;\ - T y;\ - };\ - };\ + T x;\ + T y;\ } CUDA_name##2;\ typedef struct {\ - union {\ - T data[3];\ - struct {\ - T x;\ - T y;\ - T z;\ - };\ - };\ + T x;\ + T y;\ + T z;\ } CUDA_name##3;\ typedef struct {\ - union {\ - CUDA_name##_impl4 data;\ - struct {\ - T x;\ - T y;\ - T z;\ - T w;\ - };\ - };\ + T x;\ + T y;\ + T z;\ + T w;\ } CUDA_name##4; #endif From d4fe8ff82290140d38c837be8b0232cdd5e3d9e8 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 6 Nov 2019 23:49:54 -0800 Subject: [PATCH 19/23] General hipMemset improvements (#1495) * hipMemset et al can use HSA API directly for synchronous cases * lock and flush stream in hipMemset, hold lock until complete * move hipMemset async check to front of conditional * use hsa_amd_memory_fill for additional sync memset cases code cleanup/review for all memset calls * Fix inversion of execution mutating value. * ihipMemsetSync fall back to kernel if HSA memset fails * Never fallback, never surrender. * Allow NULL stream. * Optimise memset kernel. Remove deadwood. * Update hip_memory.cpp * Clean up stream logic in sync memset * Revert "Clean up stream logic in sync memset" This reverts commit 6117dedf673367f44cc704192573a117a3d92477. [ROCm/hip commit: e31e0ca12ea65b4dbd48a78ddfed220fab1dcce7] --- projects/hip/src/hip_memory.cpp | 420 ++++++++++++++------------------ 1 file changed, 187 insertions(+), 233 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 3f82246cf9..266f9b51d6 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -41,22 +41,23 @@ hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKi // Return success if number of bytes to copy is 0 if (sizeBytes == 0) return e; + if (!dst || !src) return hipErrorInvalidValue; - stream = ihipSyncAndResolveStream(stream); - - if ((dst == NULL) || (src == NULL)) { - e = hipErrorInvalidValue; - } else if (stream) { - try { - stream->locked_copyAsync(dst, src, sizeBytes, kind); - } catch (ihipException& ex) { - e = ex._code; - } - } else { - e = hipErrorInvalidValue; + if (!(stream = ihipSyncAndResolveStream(stream))) { + return hipErrorInvalidValue; } - return e; + try { + stream->locked_copyAsync(dst, src, sizeBytes, kind); + } + catch (ihipException& ex) { + e = ex._code; + } + catch (...) { + return hipErrorUnknown; + } + + return hipSuccess; } // return 0 on success or -1 on error: @@ -504,9 +505,9 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height, unsigned int elementSizeBytes){ HIP_INIT_SPECIAL_API(hipMemAllocPitch, (TRACE_MEM), dptr, pitch, widthInBytes, height,elementSizeBytes); HIP_SET_DEVICE(); - + if (widthInBytes == 0 || height == 0) return ihipLogStatus(hipErrorInvalidValue); - + return ihipLogStatus(ihipMallocPitch(tls, dptr, pitch, widthInBytes, height, 0)); } @@ -1027,7 +1028,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { if(dst==NULL || src==NULL){ return ihipLogStatus(hipErrorInvalidValue); } - + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); hc::completion_future marker; @@ -1065,7 +1066,6 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { return ihipLogStatus(e); } - hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes); @@ -1094,7 +1094,7 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { hipError_t e = hipSuccess; if (sizeBytes == 0) return ihipLogStatus(e); - + if(dst==NULL || src==NULL){ return ihipLogStatus(hipErrorInvalidValue); } @@ -1146,8 +1146,6 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; - hipError_t e = hipSuccess; size_t byteSize; @@ -1198,8 +1196,6 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; - hipError_t e = hipSuccess; try { @@ -1217,8 +1213,6 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; - hipError_t e = hipSuccess; try { @@ -1235,8 +1229,6 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; - hipError_t e = hipSuccess; try { stream->locked_copySync((char*)dstArray->data + dstOffset, srcHost, count, @@ -1253,8 +1245,6 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; - hipError_t e = hipSuccess; try { @@ -1333,7 +1323,6 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo } stream = ihipSyncAndResolveStream(stream); - hc::completion_future marker; try { if((widthInBytes == dstPitch) && (widthInBytes == srcPitch)) { if(isAsync) @@ -1379,16 +1368,23 @@ hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms* p, hipStream_t stream } namespace { -template +template __global__ void hip_fill_n(RandomAccessIterator f, N n, T value) { - const uint32_t grid_dim = gridDim.x * blockDim.x; + const auto grid_dim = gridDim.x * blockDim.x * items_per_lane; + const auto gidx = blockIdx.x * block_dim + threadIdx.x; - size_t idx = blockIdx.x * block_dim + threadIdx.x; - while (idx < n) { - __builtin_memcpy(reinterpret_cast(&f[idx]), reinterpret_cast(&value), - sizeof(T)); + size_t idx = gidx * items_per_lane; + while (idx + items_per_lane <= n) { + for (auto i = 0u; i != items_per_lane; ++i) { + __builtin_nontemporal_store(value, &f[idx + i]); + } idx += grid_dim; } + + if (gidx < n % grid_dim) { + __builtin_nontemporal_store(value, &f[n - gidx - 1]); + } } template {}>::type* = nullptr> @@ -1443,11 +1439,14 @@ hipError_t ihipMemPtrGetInfo(void* ptr, size_t* size) { template void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t count) { static constexpr uint32_t block_dim = 256; + static constexpr uint32_t max_write_width = 4 * sizeof(std::uint32_t); // 4 DWORDs + static constexpr uint32_t items_per_lane = max_write_width / sizeof(T); - const uint32_t grid_dim = clamp_integer(count / block_dim, 1, UINT32_MAX); + const uint32_t grid_dim = clamp_integer( + count / (block_dim * items_per_lane), 1, UINT32_MAX); - hipLaunchKernelGGL(hip_fill_n, dim3(grid_dim), dim3{block_dim}, 0u, stream, ptr, - count, std::move(val)); + hipLaunchKernelGGL(hip_fill_n, dim3(grid_dim), + dim3{block_dim}, 0u, stream, ptr, count, std::move(val)); } template @@ -1466,62 +1465,133 @@ typedef enum ihipMemsetDataType { ihipMemsetDataTypeInt = 2 }ihipMemsetDataType; -hipError_t ihipMemset(void* dst, int value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType) -{ - hipError_t e = hipSuccess; +hipError_t ihipMemsetAsync(void* dst, int value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType) { + if (count == 0) return hipSuccess; + if (!dst) return hipErrorInvalidValue; - if (count == 0) return e; - - size_t allocSize = 0; - bool isInbound = (ihipMemPtrGetInfo(dst, &allocSize) == hipSuccess); - isInbound &= (allocSize >= count); - - if (stream && (dst != NULL) && isInbound) { - if(copyDataType == ihipMemsetDataTypeChar){ + try { + if (copyDataType == ihipMemsetDataTypeChar) { 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, count/sizeof(uint32_t)); - } - catch (std::exception &ex) { - e = hipErrorInvalidValue; - } - } else { + value = value & 0xff; + uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; + ihipMemsetKernel (stream, static_cast (dst), value32, count/sizeof(uint32_t)); + } else { // use a slow byte-per-workitem copy: - try { - ihipMemsetKernel (stream, static_cast (dst), value, count); - } - catch (std::exception &ex) { - e = hipErrorInvalidValue; - } - } - } else { - if(copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value - try { - 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, count); - } catch (std::exception &ex) { - e = hipErrorInvalidValue; - } + ihipMemsetKernel (stream, static_cast (dst), value, count); } + } else if (copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value + ihipMemsetKernel (stream, static_cast (dst), value, count); + } else if (copyDataType == ihipMemsetDataTypeShort) { + value = value & 0xffff; + ihipMemsetKernel (stream, static_cast (dst), value, count); } - if (HIP_API_BLOCKING) { - tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str()); - stream->locked_wait(); - } - } else { - e = hipErrorInvalidValue; + } catch (...) { + return hipErrorInvalidValue; } - return e; -}; + + if (HIP_API_BLOCKING) { + tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str()); + stream->locked_wait(); + } + + return hipSuccess; +} + +namespace { + template + void handleHeadTail(T* dst, std::size_t n_head, std::size_t n_body, + std::size_t n_tail, hipStream_t stream, int value) { + struct Cleaner { + static + __global__ + void clean(T* p, std::size_t nh, std::size_t nb, int x) noexcept { + p[(threadIdx.x < nh) ? threadIdx.x : (threadIdx.x - nh + nb)] = x; + } + }; + + hipLaunchKernelGGL(Cleaner::clean, 1, n_head + n_tail, 0, stream, + dst, n_head, + n_body * sizeof(std::uint32_t) / sizeof(T), value); + + } +} // Anonymous namespace. + +hipError_t ihipMemsetSync(void* dst, int value, size_t count, hipStream_t stream, ihipMemsetDataType copyDataType) { + if (count == 0) return hipSuccess; + if (!dst) return hipErrorInvalidValue; + + try { + size_t n = count; + auto aligned_dst{(copyDataType == ihipMemsetDataTypeInt) ? dst : + reinterpret_cast( + hip_impl::round_up_to_next_multiple_nonnegative( + reinterpret_cast(dst), 4ul))}; + size_t n_head{}; + size_t n_tail{}; + int original_value = value; + + switch (copyDataType) { + case ihipMemsetDataTypeChar: + value &= 0xff; + value = (value << 24) | (value << 16) | (value << 8) | value; + n_head = static_cast(aligned_dst) - + static_cast(dst); + n -= n_head; + n /= sizeof(std::uint32_t); + n_tail = count % sizeof(std::uint32_t); + break; + case ihipMemsetDataTypeShort: + value &= 0xffff; + value = (value << 16) | value; + n_head = static_cast(aligned_dst) - + static_cast(dst); + n = (count - n_head) * + sizeof(std::uint16_t) / sizeof(std::uint32_t); + n_tail = ((count - n_head) * + sizeof(std::uint16_t)) % sizeof(std::uint32_t); + break; + default: break; + } + + // queue the memset kernel for the remainder of the buffer before the HSA call below + if (aligned_dst != dst || n_tail != 0) { + switch (copyDataType) { + case ihipMemsetDataTypeChar: + handleHeadTail(static_cast(dst), n_head, n, + n_tail, stream, value & 0xff); + break; + case ihipMemsetDataTypeShort: + handleHeadTail(static_cast(dst), n_head, n, + n_tail, stream, value & 0xffff); + break; + default: break; + } + } + + // The stream must be locked from all other op insertions to guarantee + // that the following HSA call can complete before any other ops. + // Flush the stream while locked. Once the stream is empty, we can safely perform + // the out-of-band HSA call. Lastly, the stream will unlock via RAII. + if (!stream) stream = ihipSyncAndResolveStream(stream); + if (!stream) return hipErrorInvalidValue; + + LockedAccessor_StreamCrit_t crit(stream->criticalData()); + crit->_av.wait(stream->waitMode()); + const auto s = hsa_amd_memory_fill(aligned_dst, value, n); + if (s != HSA_STATUS_SUCCESS) return hipErrorInvalidValue; + } + catch (...) { + return hipErrorInvalidValue; + } + + if (HIP_API_BLOCKING) { + tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetSync.\n", ToString(stream).c_str()); + stream->locked_wait(); + } + + return hipSuccess; +} hipError_t getLockedPointer(void *hostPtr, size_t dataLen, void **devicePtrPtr) { @@ -1539,7 +1609,7 @@ hipError_t getLockedPointer(void *hostPtr, size_t dataLen, void **devicePtrPtr) return(hipSuccess); }; return(hipErrorHostMemoryNotRegistered); -}; +} // TODO - review and optimize hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, @@ -1778,182 +1848,66 @@ hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { HIP_INIT_SPECIAL_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream); - - hipError_t e = hipSuccess; - - stream = ihipSyncAndResolveStream(stream); - - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); - - return ihipLogStatus(e); -}; + return ihipLogStatus(ihipMemsetAsync(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar)); +} 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); -}; + return ihipLogStatus(ihipMemsetAsync(dst, value, count, stream, ihipMemsetDataTypeInt)); +} hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes); - - hipError_t e = hipSuccess; - - hipStream_t stream = hipStreamNull; - stream = ihipSyncAndResolveStream(stream); - if (stream) { - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); - stream->locked_wait(); - } else { - e = hipErrorInvalidValue; - } - return ihipLogStatus(e); + return ihipLogStatus(ihipMemsetSync(dst, value, sizeBytes, nullptr, ihipMemsetDataTypeChar)); } hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) { HIP_INIT_SPECIAL_API(hipMemset2D, (TRACE_MCMD), dst, pitch, value, width, height); - - hipError_t e = hipSuccess; - - hipStream_t stream = hipStreamNull; - stream = ihipSyncAndResolveStream(stream); - if (stream) { - size_t sizeBytes = pitch * height; - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); - stream->locked_wait(); - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); + size_t sizeBytes = pitch * height; + return ihipLogStatus(ihipMemsetSync(dst, value, sizeBytes, nullptr, ihipMemsetDataTypeChar)); } -hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream ) -{ +hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream ) { HIP_INIT_SPECIAL_API(hipMemset2DAsync, (TRACE_MCMD), dst, pitch, value, width, height, stream); - - hipError_t e = hipSuccess; - - stream = ihipSyncAndResolveStream(stream); - - if (stream) { - size_t sizeBytes = pitch * height; - e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar); - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); -}; + size_t sizeBytes = pitch * height; + return ihipLogStatus(ihipMemsetAsync(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar)); +} hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) { HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, count); - - hipError_t e = hipSuccess; - - hipStream_t stream = hipStreamNull; - stream = ihipSyncAndResolveStream(stream); - if (stream) { - e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeChar); - stream->locked_wait(); - } else { - e = hipErrorInvalidValue; - } - return ihipLogStatus(e); + return ihipLogStatus(ihipMemsetSync(dst, value, count, nullptr, ihipMemsetDataTypeChar)); } hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t count , hipStream_t stream ) { HIP_INIT_SPECIAL_API(hipMemsetD8Async, (TRACE_MCMD), dst, value, count, stream); - - stream = ihipSyncAndResolveStream(stream); - if (stream) { - return ihipLogStatus(ihipMemset(dst, value, count, stream, ihipMemsetDataTypeChar)); - } else { - return ihipLogStatus(hipErrorInvalidValue); - } + return ihipLogStatus(ihipMemsetAsync(dst, value, count, stream, ihipMemsetDataTypeChar)); } hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t count){ HIP_INIT_SPECIAL_API(hipMemsetD16, (TRACE_MCMD), dst, value, count); - hipError_t e = hipSuccess; - hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - if (stream) { - e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeShort); - if(hipSuccess == e) - stream->locked_wait(); - } else { - e = hipErrorInvalidValue; - } - return ihipLogStatus(e); + return ihipLogStatus(ihipMemsetSync(dst, value, count, nullptr, ihipMemsetDataTypeShort)); } hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t count, hipStream_t stream ){ HIP_INIT_SPECIAL_API(hipMemsetD16Async, (TRACE_MCMD), dst, value, count, stream); - - stream = ihipSyncAndResolveStream(stream); - if (stream) { - return ihipLogStatus(ihipMemset(dst, value, count, stream, ihipMemsetDataTypeShort)); - } else { - return ihipLogStatus(hipErrorInvalidValue); - } + return ihipLogStatus(ihipMemsetAsync(dst, value, count, stream, ihipMemsetDataTypeShort)); } 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); + return ihipLogStatus(ihipMemsetSync(dst, value, count, nullptr, ihipMemsetDataTypeInt)); } -hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ) -{ +hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) { HIP_INIT_SPECIAL_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent); - hipError_t e = hipSuccess; - - hipStream_t stream = hipStreamNull; - // TODO - call an ihip memset so HIP_TRACE is correct. - stream = ihipSyncAndResolveStream(stream); - if (stream) { - size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; - e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar); - stream->locked_wait(); - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); + size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; + return ihipLogStatus(ihipMemsetSync(pitchedDevPtr.ptr, value, sizeBytes, nullptr, ihipMemsetDataTypeChar)); } -hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream ) -{ +hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream ) { HIP_INIT_SPECIAL_API(hipMemset3DAsync, (TRACE_MCMD), &pitchedDevPtr, value, &extent); - hipError_t e = hipSuccess; - - // TODO - call an ihip memset so HIP_TRACE is correct. - stream = ihipSyncAndResolveStream(stream); - if (stream) { - size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; - e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar); - } else { - e = hipErrorInvalidValue; - } - - return ihipLogStatus(e); + size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; + return ihipLogStatus(ihipMemsetAsync(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar)); } hipError_t hipMemGetInfo(size_t* free, size_t* total) { @@ -1969,20 +1923,20 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) { } else { e = hipErrorInvalidValue; } - + if (free) { if (!device->_driver_node_id) return ihipLogStatus(hipErrorInvalidDevice); - - std::string fileName = std::string("/sys/class/kfd/kfd/topology/nodes/") + std::to_string(device->_driver_node_id) + std::string("/mem_banks/0/used_memory"); + + std::string fileName = std::string("/sys/class/kfd/kfd/topology/nodes/") + std::to_string(device->_driver_node_id) + std::string("/mem_banks/0/used_memory"); std::ifstream file; file.open(fileName); if (!file) return ihipLogStatus(hipErrorFileNotFound); - - std::string deviceSize; + + std::string deviceSize; size_t deviceMemSize; - + file >> deviceSize; - file.close(); + file.close(); if ((deviceMemSize=strtol(deviceSize.c_str(),NULL,10))){ *free = device->_props.totalGlobalMem - deviceMemSize; // Deduct the amount of memory from the free memory reported from the system From 58fc750bf02b9a05c76d50bceee9abb33b522969 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 7 Nov 2019 11:21:06 +0300 Subject: [PATCH 20/23] [HIPIFY] Reorder options propagation to clang [Reason] Leave overriding opportunity for -D, -std=c++XX, etc. [ROCm/hip commit: 16a27213e43b86b1f6980f7d48a0244e466a2699] --- projects/hip/hipify-clang/src/main.cpp | 49 +++++++++++++------------- 1 file changed, 25 insertions(+), 24 deletions(-) diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 3c7424fbca..812ba48197 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -194,19 +194,17 @@ int main(int argc, const char **argv) { ct::RefactoringTool Tool(OptionsParser.getCompilations(), std::string(tmpFile.c_str())); ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str()); ReplacementsFrontendActionFactory actionFactory(&replacementsToUse); - std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str(); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); - if (llcompat::canCompileHostAndDeviceInOneJob()) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-compile-host-device", ct::ArgumentInsertPosition::BEGIN)); - } else { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); + if (!IncludeDirs.empty()) { + for (std::string s : IncludeDirs) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::BEGIN)); + } } - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN)); - if (!CudaPath.empty()) { - std::string sCudaPath = "--cuda-path=" + CudaPath; - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN)); + if (!MacroNames.empty()) { + for (std::string s : MacroNames) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN)); + } } // Includes for clang's CUDA wrappers for using by packaged hipify-clang Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN)); @@ -219,28 +217,31 @@ int main(int argc, const char **argv) { stdCpp = "-std=c++14"; #endif Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN)); + std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str(); #if defined(HIPIFY_CLANG_RES) - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES, ct::ArgumentInsertPosition::BEGIN)); #endif + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN)); if (llcompat::pragma_once_outside_header()) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN)); } + if (llcompat::canCompileHostAndDeviceInOneJob()) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-compile-host-device", ct::ArgumentInsertPosition::BEGIN)); + } + else { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); + } if (!CudaGpuArch.empty()) { std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch; Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN)); } - if (!MacroNames.empty()) { - for (std::string s : MacroNames) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::END)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END)); - } - } - if (!IncludeDirs.empty()) { - for (std::string s : IncludeDirs) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::END)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END)); - } + if (!CudaPath.empty()) { + std::string sCudaPath = "--cuda-path=" + CudaPath; + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN)); } + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); if (Verbose) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-v", ct::ArgumentInsertPosition::END)); } From 860e89df5e46a06c1fa37c9a6845c1d0749545e3 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 7 Nov 2019 11:30:40 +0300 Subject: [PATCH 21/23] [HIPIFY][fix] Abandon canCompileHostAndDeviceInOneJob check for --cuda-compile-host-device option [Reason] It turned out that it is not so: 2 jobs are always [ROCm/hip commit: c2c144ac04c47b0df85dbce4c29a763a05511b7b] --- projects/hip/hipify-clang/src/LLVMCompat.cpp | 8 -------- projects/hip/hipify-clang/src/main.cpp | 7 +------ 2 files changed, 1 insertion(+), 14 deletions(-) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index b524c41613..604841bcd9 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp @@ -127,14 +127,6 @@ bool pragma_once_outside_header() { #endif } -bool canCompileHostAndDeviceInOneJob() { -#if LLVM_VERSION_MAJOR > 9 && defined(_WIN32) - return true; -#else - return false; -#endif -} - void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI) { #if LLVM_VERSION_MAJOR > 9 clang::PreprocessorOptions &PPOpts = CI.getPreprocessorOpts(); diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 812ba48197..b4200bd109 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -226,12 +226,7 @@ int main(int argc, const char **argv) { if (llcompat::pragma_once_outside_header()) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN)); } - if (llcompat::canCompileHostAndDeviceInOneJob()) { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-compile-host-device", ct::ArgumentInsertPosition::BEGIN)); - } - else { - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); - } + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); if (!CudaGpuArch.empty()) { std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch; Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN)); From 0ae791586b4227ca9ff8751af67c0ad8cdd37e81 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 7 Nov 2019 11:32:53 +0300 Subject: [PATCH 22/23] [HIPIFY][fix] Delete canCompileHostAndDeviceInOneJob declaration as well [ROCm/hip commit: 0c07b6b421acb7e17002c0408ddd7758e6b65992] --- projects/hip/hipify-clang/src/LLVMCompat.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h index e857023752..48e008d40d 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.h +++ b/projects/hip/hipify-clang/src/LLVMCompat.h @@ -85,8 +85,6 @@ std::error_code real_path(const Twine &path, SmallVectorImpl &output, bool pragma_once_outside_header(); -bool canCompileHostAndDeviceInOneJob(); - void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI); bool CheckCompatibility(); From 9e34880005fb7ce08f54e8d6ee517aa106fe74b8 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 7 Nov 2019 12:48:05 +0300 Subject: [PATCH 23/23] [HIPIFY] Clang style formatting [ROCm/hip commit: 93bc9c91950465c5ed3443acf4a869de3ec52417] --- .../hip/hipify-clang/src/CUDA2HIP_Perl.cpp | 24 ++++++------- .../hip/hipify-clang/src/CUDA2HIP_Python.cpp | 6 ++-- .../hip/hipify-clang/src/HipifyAction.cpp | 2 +- .../src/ReplacementsFrontendActionFactory.h | 6 ++-- projects/hip/hipify-clang/src/Statistics.cpp | 34 +++++++++---------- projects/hip/hipify-clang/src/Statistics.h | 14 ++++---- projects/hip/hipify-clang/src/StringUtils.cpp | 8 ++--- projects/hip/hipify-clang/src/StringUtils.h | 8 ++--- 8 files changed, 51 insertions(+), 51 deletions(-) diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index cd5c952145..d74ba70f21 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -117,7 +117,7 @@ namespace perl { {sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam} }; - void generateHeader(unique_ptr& streamPtr) { + void generateHeader(unique_ptr &streamPtr) { *streamPtr.get() << "#!/usr/bin/perl -w" << endl_2; *streamPtr.get() << sCopyright << endl; *streamPtr.get() << sImportant << endl_2; @@ -148,7 +148,7 @@ namespace perl { *streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << endl_2; } - void generateStatFunctions(unique_ptr& streamPtr) { + void generateStatFunctions(unique_ptr &streamPtr) { *streamPtr.get() << endl << sub << "totalStats" << " {" << endl; *streamPtr.get() << tab << my << "%count = %{ shift() };" << endl; *streamPtr.get() << tab << my << "$total = 0;" << endl; @@ -175,11 +175,11 @@ namespace perl { } } - void generateSimpleSubstitutions(unique_ptr& streamPtr) { + void generateSimpleSubstitutions(unique_ptr &streamPtr) { *streamPtr.get() << endl << sub << "simpleSubstitutions" << " {" << endl; for (int i = 0; i < NUM_CONV_TYPES; ++i) { if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) { - for (auto& ma : CUDA_INCLUDE_MAP) { + for (auto &ma : CUDA_INCLUDE_MAP) { if (Statistics::isUnsupported(ma.second)) continue; if (i == ma.second.type) { string sCUDA = ma.first.str(); @@ -190,7 +190,7 @@ namespace perl { } } } else { - for (auto& ma : CUDA_RENAMES_MAP()) { + for (auto &ma : CUDA_RENAMES_MAP()) { if (Statistics::isUnsupported(ma.second)) continue; if (i == ma.second.type) { *streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << endl; @@ -201,7 +201,7 @@ namespace perl { *streamPtr.get() << "}" << endl; } - void generateExternShared(unique_ptr& streamPtr) { + void generateExternShared(unique_ptr &streamPtr) { *streamPtr.get() << endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << endl; *streamPtr.get() << sub << "transformExternShared" << " {" << endl; *streamPtr.get() << tab << no_warns << endl; @@ -210,7 +210,7 @@ namespace perl { *streamPtr.get() << tab << "$ft{'extern_shared'} += $k;" << endl << "}" << endl; } - void generateKernelLaunch(unique_ptr& streamPtr) { + void generateKernelLaunch(unique_ptr &streamPtr) { *streamPtr.get() << endl << "# CUDA Kernel Launch Syntax" << endl << sub << "transformKernelLaunch" << " {" << endl; *streamPtr.get() << tab << no_warns << endl; *streamPtr.get() << tab << my_k << endl_2; @@ -251,13 +251,13 @@ namespace perl { *streamPtr.get() << tab_2 << "$Tkernels{$1}++;" << endl_tab << "}" << endl << "}" << endl; } - void generateCubNamespace(unique_ptr& streamPtr) { + void generateCubNamespace(unique_ptr &streamPtr) { *streamPtr.get() << endl << sub << "transformCubNamespace" << " {" << endl_tab << my_k << endl; *streamPtr.get() << tab << "$k += s/using\\s*namespace\\s*cub/using namespace hipcub/g;" << endl; *streamPtr.get() << tab << "$k += s/\\bcub::\\b/hipcub::/g;" << endl << tab << return_k << "}" << endl; } - void generateHostFunctions(unique_ptr& streamPtr) { + void generateHostFunctions(unique_ptr &streamPtr) { *streamPtr.get() << endl << sub << "transformHostFunctions" << " {" << endl_tab << my_k << endl; set &funcSet = DeviceSymbolFunctions0; const string s0 = "$k += s/(?second.hipName.str() << "\""; @@ -291,12 +291,12 @@ namespace perl { *streamPtr.get() << tab << return_k << "}" << endl; } - void generateDeviceFunctions(unique_ptr& streamPtr) { + void generateDeviceFunctions(unique_ptr &streamPtr) { unsigned int countUnsupported = 0; unsigned int countSupported = 0; stringstream sSupported; stringstream sUnsupported; - for (auto& ma : CUDA_DEVICE_FUNC_MAP) { + for (auto &ma : CUDA_DEVICE_FUNC_MAP) { bool isUnsupported = Statistics::isUnsupported(ma.second); (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << tab_2 << "\"" << ma.first.str() << "\""; if (isUnsupported) countUnsupported++; diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp index 5ffd84e0b0..fec138915d 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp @@ -60,9 +60,9 @@ namespace python { *pythonStreamPtr.get() << "from pyHIPIFY.constants import *\n\n"; *pythonStreamPtr.get() << "CUDA_RENAMES_MAP = collections.OrderedDict([\n"; const std::string sHIP_UNS = ", HIP_UNSUPPORTED"; - for (int i = 0; i < NUM_CONV_TYPES; i++) { + for (int i = 0; i < NUM_CONV_TYPES; ++i) { if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) { - for (auto& ma : CUDA_INCLUDE_MAP) { + for (auto &ma : CUDA_INCLUDE_MAP) { if (i == ma.second.type) { std::string sUnsupported; if (Statistics::isUnsupported(ma.second)) { @@ -74,7 +74,7 @@ namespace python { } } else { - for (auto& ma : CUDA_RENAMES_MAP()) { + for (auto &ma : CUDA_RENAMES_MAP()) { if (i == ma.second.type) { std::string sUnsupported; if (Statistics::isUnsupported(ma.second)) { diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 032bfef2aa..dee9a25d49 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -150,7 +150,7 @@ void HipifyAction::RewriteToken(const clang::Token &t) { void HipifyAction::FindAndReplace(StringRef name, clang::SourceLocation sl, - const std::map &repMap, + const std::map &repMap, bool bReplace) { const auto found = repMap.find(name); if (found == repMap.end()) { diff --git a/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h b/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h index 77d23ff6b2..92d77655af 100644 --- a/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h +++ b/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h @@ -36,15 +36,15 @@ namespace ct = clang::tooling; */ template class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory { - ct::Replacements* replacements; + ct::Replacements *replacements; public: - explicit ReplacementsFrontendActionFactory(ct::Replacements* r): + explicit ReplacementsFrontendActionFactory(ct::Replacements *r): ct::FrontendActionFactory(), replacements(r) {} #if LLVM_VERSION_MAJOR < 10 - clang::FrontendAction* create() override { + clang::FrontendAction *create() override { return new T(replacements); } #else diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp index 8dff998dfa..686b51da41 100644 --- a/projects/hip/hipify-clang/src/Statistics.cpp +++ b/projects/hip/hipify-clang/src/Statistics.cpp @@ -163,51 +163,51 @@ void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string } // Anonymous namespace -void StatCounter::incrementCounter(const hipCounter& counter, const std::string& name) { +void StatCounter::incrementCounter(const hipCounter &counter, const std::string &name) { counters[name]++; apiCounters[(int) counter.apiType]++; convTypeCounters[(int) counter.type]++; } -void StatCounter::add(const StatCounter& other) { - for (const auto& p : other.counters) { +void StatCounter::add(const StatCounter &other) { + for (const auto &p : other.counters) { counters[p.first] += p.second; } - for (int i = 0; i < NUM_API_TYPES; i++) { + for (int i = 0; i < NUM_API_TYPES; ++i) { apiCounters[i] += other.apiCounters[i]; } - for (int i = 0; i < NUM_CONV_TYPES; i++) { + for (int i = 0; i < NUM_CONV_TYPES; ++i) { convTypeCounters[i] += other.convTypeCounters[i]; } } int StatCounter::getConvSum() { int acc = 0; - for (const int& i : convTypeCounters) { + for (const int &i : convTypeCounters) { acc += i; } return acc; } -void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix) { - for (int i = 0; i < NUM_CONV_TYPES; i++) { +void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string &prefix) { + for (int i = 0; i < NUM_CONV_TYPES; ++i) { if (convTypeCounters[i] > 0) { conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n"); break; } } - for (int i = 0; i < NUM_CONV_TYPES; i++) { + for (int i = 0; i < NUM_CONV_TYPES; ++i) { if (convTypeCounters[i] > 0) { printStat(csv, printOut, counterNames[i], convTypeCounters[i]); } } - for (int i = 0; i < NUM_API_TYPES; i++) { + for (int i = 0; i < NUM_API_TYPES; ++i) { if (apiCounters[i] > 0) { conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n"); break; } } - for (int i = 0; i < NUM_API_TYPES; i++) { + for (int i = 0; i < NUM_API_TYPES; ++i) { if (apiCounters[i] > 0) { printStat(csv, printOut, apiNames[i], apiCounters[i]); } @@ -220,7 +220,7 @@ void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const st } } -Statistics::Statistics(const std::string& name): fileName(name) { +Statistics::Statistics(const std::string &name): fileName(name) { // Compute the total bytes/lines in the input file. std::ifstream src_file(name, std::ios::binary | std::ios::ate); src_file.clear(); @@ -235,7 +235,7 @@ Statistics::Statistics(const std::string& name): fileName(name) { ///////// Counter update routines ////////// -void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) { +void Statistics::incrementCounter(const hipCounter &counter, const std::string &name) { if (Statistics::isUnsupported(counter)) { unsupported.incrementCounter(counter, name); } else { @@ -308,7 +308,7 @@ void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) Statistics globalStats = getAggregate(); // A file is considered "converted" if we made any changes to it. int convertedFiles = 0; - for (const auto& p : stats) { + for (const auto &p : stats) { if (p.second.touchedLines && p.second.totalBytes && p.second.totalLines && !p.second.hasErrors) { convertedFiles++; @@ -326,18 +326,18 @@ void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) Statistics Statistics::getAggregate() { Statistics globalStats("GLOBAL"); - for (const auto& p : stats) { + for (const auto &p : stats) { globalStats.add(p.second); } return globalStats; } -Statistics& Statistics::current() { +Statistics &Statistics::current() { assert(Statistics::currentStatistics); return *Statistics::currentStatistics; } -void Statistics::setActive(const std::string& name) { +void Statistics::setActive(const std::string &name) { stats.emplace(std::make_pair(name, Statistics{name})); Statistics::currentStatistics = &stats.at(name); } diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h index 9b9889d0e5..030ea6b2d6 100644 --- a/projects/hip/hipify-clang/src/Statistics.h +++ b/projects/hip/hipify-clang/src/Statistics.h @@ -172,11 +172,11 @@ private: int convTypeCounters[NUM_CONV_TYPES] = {}; public: - void incrementCounter(const hipCounter& counter, const std::string& name); + void incrementCounter(const hipCounter &counter, const std::string &name); // Add the counters from `other` onto the counters of this object. - void add(const StatCounter& other); + void add(const StatCounter &other); int getConvSum(); - void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix); + void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string &prefix); }; /** @@ -195,8 +195,8 @@ class Statistics { chr::steady_clock::time_point completionTime; public: - Statistics(const std::string& name); - void incrementCounter(const hipCounter &counter, const std::string& name); + Statistics(const std::string &name); + void incrementCounter(const hipCounter &counter, const std::string &name); // Add the counters from `other` onto the counters of this object. void add(const Statistics &other); void lineTouched(int lineNumber); @@ -226,12 +226,12 @@ public: * processing one file at a time, this allows us to simply expose the stats for the current file globally, * simplifying things. */ - static Statistics& current(); + static Statistics ¤t(); /** * Set the active Statistics object to the named one, creating it if necessary, and write the completion * timestamp into the currently active one. */ - static void setActive(const std::string& name); + static void setActive(const std::string &name); // Check the counter and option TranslateToRoc whether it should be translated to Roc or not. static bool isToRoc(const hipCounter &counter); // Check whether the counter is HIP_UNSUPPORTED or not. diff --git a/projects/hip/hipify-clang/src/StringUtils.cpp b/projects/hip/hipify-clang/src/StringUtils.cpp index 78d0f8113e..31ab331174 100644 --- a/projects/hip/hipify-clang/src/StringUtils.cpp +++ b/projects/hip/hipify-clang/src/StringUtils.cpp @@ -33,14 +33,14 @@ llvm::StringRef unquoteStr(llvm::StringRef s) { return s; } -void removePrefixIfPresent(std::string &s, const std::string& prefix) { +void removePrefixIfPresent(std::string &s, const std::string &prefix) { if (s.find(prefix) != 0) { return; } s.erase(0, prefix.size()); } -std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) { +std::string getAbsoluteFilePath(const std::string &sFile, std::error_code &EC) { if (sFile.empty()) { return sFile; } @@ -59,8 +59,8 @@ std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) { return fileAbsPath.c_str(); } -std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC, - const std::string& sDirType, bool bCreateDir) { +std::string getAbsoluteDirectoryPath(const std::string &sDir, std::error_code &EC, + const std::string &sDirType, bool bCreateDir) { if (sDir.empty()) { return sDir; } diff --git a/projects/hip/hipify-clang/src/StringUtils.h b/projects/hip/hipify-clang/src/StringUtils.h index 39ce726d1b..ecbca5e832 100644 --- a/projects/hip/hipify-clang/src/StringUtils.h +++ b/projects/hip/hipify-clang/src/StringUtils.h @@ -33,16 +33,16 @@ llvm::StringRef unquoteStr(llvm::StringRef s); /** * If `s` starts with `prefix`, remove it. Otherwise, does nothing. */ -void removePrefixIfPresent(std::string &s, const std::string& prefix); +void removePrefixIfPresent(std::string &s, const std::string &prefix); /** * Returns Absolute File Path based on filename, otherwise - error. */ -std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC); +std::string getAbsoluteFilePath(const std::string &sFile, std::error_code &EC); /** * Returns Absolute Directory Path based on directory name, otherwise - error; * by default the directory is temporary and created. */ -std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC, - const std::string& sDirType = "temporary", bool bCreateDir = true); +std::string getAbsoluteDirectoryPath(const std::string &sDir, std::error_code &EC, + const std::string &sDirType = "temporary", bool bCreateDir = true);