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/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/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) { 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)); diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index b203d942a8..75c79d422c 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -58,10 +58,19 @@ THE SOFTWARE. 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]; + } }; // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor @@ -69,6 +78,8 @@ THE SOFTWARE. __host__ __device__ operator T() const noexcept { return data[idx]; } + __host__ __device__ + operator T() const volatile noexcept { return data[idx]; } __host__ __device__ Address operator&() const noexcept { return Address{this}; } @@ -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 {