From c5a2a2daf24cca6de2b2fca24cc134be1b9692ad Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 1 Nov 2019 14:34:18 +0300 Subject: [PATCH 1/4] [HIPIFY][CUB][#1460][perl] Add "cub::" namespace prefix support in hipify-perl as well --- hipamd/bin/hipify-perl | 1 + hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp | 3 ++- 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/hipamd/bin/hipify-perl b/hipamd/bin/hipify-perl index 64a9e1dd5b..7b286d638d 100755 --- a/hipamd/bin/hipify-perl +++ b/hipamd/bin/hipify-perl @@ -1680,6 +1680,7 @@ sub transformKernelLaunch { sub transformCubNamespace { my $k = 0; $k += s/using\s*namespace\s*cub/using namespace hipcub/g; + $k += s/\bcub::\b/hipcub::/g; return $k; } diff --git a/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp b/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp index 8d59089d4f..cd5c952145 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -253,7 +253,8 @@ namespace perl { 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 << tab << return_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) { From 2d76dde05b43534a8ff397c198a2a595cee9c8f7 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 1 Nov 2019 22:18:01 +0200 Subject: [PATCH 2/4] Accessors should work even when oddly volatile. --- hipamd/include/hip/hcc_detail/hip_vector_types.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index b203d942a8..9764bc2a16 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -54,11 +54,11 @@ THE SOFTWARE. const Scalar_accessor* p; __host__ __device__ - operator const T*() const noexcept { + operator const T*() const volatile noexcept { return &reinterpret_cast(p)[idx]; } __host__ __device__ - operator T*() noexcept { + operator T*() volatile noexcept { return &reinterpret_cast( const_cast(p))[idx]; } @@ -68,7 +68,7 @@ THE SOFTWARE. Vector data; __host__ __device__ - operator T() const noexcept { return data[idx]; } + operator T() const volatile noexcept { return data[idx]; } __host__ __device__ Address operator&() const noexcept { return Address{this}; } From 7142b884ab8d731c117c637b7acc6a54029dacbf Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sat, 2 Nov 2019 14:19:31 +0300 Subject: [PATCH 3/4] [HIPIFY] Introduce --cuda-gpu-arch as hipify-clang's option + Pass it to clang if specified --- hipamd/hipify-clang/src/ArgParse.cpp | 8 ++++++++ hipamd/hipify-clang/src/ArgParse.h | 1 + hipamd/hipify-clang/src/main.cpp | 4 ++++ 3 files changed, 13 insertions(+) diff --git a/hipamd/hipify-clang/src/ArgParse.cpp b/hipamd/hipify-clang/src/ArgParse.cpp index 5ce92dfdca..4f648c996f 100644 --- a/hipamd/hipify-clang/src/ArgParse.cpp +++ b/hipamd/hipify-clang/src/ArgParse.cpp @@ -138,4 +138,12 @@ cl::opt SkipExcludedPPConditionalBlocks("skip-excluded-preprocessor-condit cl::value_desc("skip-excluded-preprocessor-conditional-blocks"), cl::cat(ToolTemplateCategory)); +cl::opt CudaGpuArch("cuda-gpu-arch", + cl::desc("CUDA GPU architecture (e.g. sm_35);\nmay be specified more than once"), + cl::value_desc("value"), + cl::ZeroOrMore, + cl::Prefix, + cl::cat(ToolTemplateCategory)); + + cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage); diff --git a/hipamd/hipify-clang/src/ArgParse.h b/hipamd/hipify-clang/src/ArgParse.h index 886a658d78..84053a036c 100644 --- a/hipamd/hipify-clang/src/ArgParse.h +++ b/hipamd/hipify-clang/src/ArgParse.h @@ -52,3 +52,4 @@ extern cl::extrahelp CommonHelp; extern cl::opt TranslateToRoc; extern cl::opt DashDash; extern cl::opt SkipExcludedPPConditionalBlocks; +extern cl::opt CudaGpuArch; diff --git a/hipamd/hipify-clang/src/main.cpp b/hipamd/hipify-clang/src/main.cpp index 27dcf0b1ac..3c7424fbca 100644 --- a/hipamd/hipify-clang/src/main.cpp +++ b/hipamd/hipify-clang/src/main.cpp @@ -225,6 +225,10 @@ 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 (!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)); From ed0d6ec51e858e380cf2bd07ff8add0c0940348a Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 2 Nov 2019 22:02:08 +0200 Subject: [PATCH 4/4] Separate volatile for clarity. Handle assignment. --- .../include/hip/hcc_detail/hip_vector_types.h | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 9764bc2a16..75c79d422c 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -53,11 +53,20 @@ THE SOFTWARE. struct Address { const Scalar_accessor* p; + __host__ __device__ + operator const T*() const noexcept { + return &reinterpret_cast(p)[idx]; + } __host__ __device__ operator const T*() const volatile noexcept { return &reinterpret_cast(p)[idx]; } __host__ __device__ + operator T*() noexcept { + return &reinterpret_cast( + const_cast(p))[idx]; + } + __host__ __device__ operator T*() volatile noexcept { return &reinterpret_cast( const_cast(p))[idx]; @@ -67,6 +76,8 @@ THE SOFTWARE. // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor Vector data; + __host__ __device__ + operator T() const noexcept { return data[idx]; } __host__ __device__ operator T() const volatile noexcept { return data[idx]; } @@ -79,6 +90,12 @@ THE SOFTWARE. return *this; } + __host__ __device__ + volatile Scalar_accessor& operator=(T x) volatile noexcept { + data[idx] = x; + + return *this; + } __host__ __device__ Scalar_accessor& operator++() noexcept {