Merge 'master' into 'amd-master'

Change-Id: Ic8372fe0f6b4f6edaae468e01f6c4583dc028d3c

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;
 }
 

hipamd/hipify-clang/src/ArgParse.cpp

@@ -138,4 +138,12 @@ cl::opt<bool> SkipExcludedPPConditionalBlocks("skip-excluded-preprocessor-condit
   cl::value_desc("skip-excluded-preprocessor-conditional-blocks"),
   cl::cat(ToolTemplateCategory));
 
+cl::opt<std::string> 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);

hipamd/hipify-clang/src/ArgParse.h

@@ -52,3 +52,4 @@ extern cl::extrahelp CommonHelp;
 extern cl::opt<bool> TranslateToRoc;
 extern cl::opt<bool> DashDash;
 extern cl::opt<bool> SkipExcludedPPConditionalBlocks;
+extern cl::opt<std::string> CudaGpuArch;

hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp

@@ -253,7 +253,8 @@ namespace perl {
 
   void generateCubNamespace(unique_ptr<ostream>& 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<ostream>& streamPtr) {

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));

hipamd/include/hip/hcc_detail/hip_vector_types.h

@@ -58,10 +58,19 @@ THE SOFTWARE.
                     return &reinterpret_cast<const T*>(p)[idx];
                 }
                 __host__ __device__
+                operator const T*() const volatile noexcept {
+                    return &reinterpret_cast<const T*>(p)[idx];
+                }
+                __host__ __device__
                 operator T*() noexcept {
                     return &reinterpret_cast<T*>(
                         const_cast<Scalar_accessor*>(p))[idx];
                 }
+                __host__ __device__
+                operator T*() volatile noexcept {
+                    return &reinterpret_cast<T*>(
+                        const_cast<Scalar_accessor*>(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 {