Merge 'master' into 'amd-master'

Change-Id: I9e11a66df0398a3794eaeaf0d3994da225649be0

bin/hipcc

@@ -22,8 +22,10 @@ use Cwd 'abs_path';
 # Other environment variable controls:
 # HIP_PATH  : Path to HIP directory, default is one dir level above location of this script
 # CUDA_PATH : Path to CUDA SDK (default /usr/local/cuda). Used on NVIDIA platforms only.
-# HCC_HOME  : Path to HCC SDK (default /opt/rocm/hcc).  Used on AMD platforms only.
+# HCC_HOME  : Path to HCC SDK (defaults to ../../hcc relative to this
-# HSA_PATH  : Path to HSA dir (default /opt/rocm/hsa).  Used on AMD platforms only.
+#             script's abs_path). Used on AMD platforms only.
+# HSA_PATH  : Path to HSA dir (defaults to ../../hsa relative to abs_path
+#             of this script). Used on AMD platforms only.
 # HIP_VDI_HOME : Path to HIP/VDI directory. Used on AMD platforms only.
 
 if(scalar @ARGV == 0){
@@ -58,7 +60,23 @@ $isWindows = $^O eq 'MSWin32';
 $HIPCC_COMPILE_FLAGS_APPEND=$ENV{'HIPCC_COMPILE_FLAGS_APPEND'};
 $HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'};
 
-$HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc
+#
+# TODO: Fix rpath LDFLAGS settings
+#
+# Since this hipcc script gets installed at two uneven hierarchical levels,
+# linked by symlink, the absolute path of this script should be used to
+# derive HIP_PATH, as dirname $0 could be /opt/rocm/bin or /opt/rocm/hip/bin
+# depending on how it gets invoked.
+# ROCM_PATH which points to <rocm_install_dir> is determined based on whether
+# we find .info/version in the parent of HIP_PATH or not. If it is found,
+# ROCM_PATH is defined relative to HIP_PATH else it is hardcoded to /opt/rocm.
+#
+$HIP_PATH=$ENV{'HIP_PATH'} // dirname(Cwd::abs_path("$0/../")); # use parent directory of hipcc
+if (-e "$HIP_PATH/../.info/version") {
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // dirname("$HIP_PATH"); # use parent directory of HIP_PATH
+} else {
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
+}
 $HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'};
 $HIP_LIB_PATH=$ENV{'HIP_LIB_PATH'};
 $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
@@ -68,7 +86,7 @@ $HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibilit
 if (defined $HIP_VDI_HOME) {
     $HIP_INFO_PATH= "$HIP_VDI_HOME/lib/.hipInfo";
 } else {
-    $HIP_INFO_PATH= "$HIP_PATH/lib/.hipInfo";
+    $HIP_INFO_PATH= "$HIP_PATH/lib/.hipInfo"; # use actual file
 }
 
 #---
@@ -112,7 +130,7 @@ if (defined $HIP_RUNTIME and $HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
     if (-e "$hipcc_dir/../lib/bitcode") {
         $HIP_VDI_HOME = abs_path($hipcc_dir . "/..");
     } else {
-        $HIP_VDI_HOME = "/opt/rocm/hip";
+        $HIP_VDI_HOME = $HIP_PATH; # use HIP_PATH
     }
 }
 
@@ -136,10 +154,10 @@ if (defined $HIP_VDI_HOME) {
 if (defined $HIP_COMPILER and $HIP_COMPILER eq "clang") {
   $HIP_PLATFORM = "clang";
   if (!defined $HIP_CLANG_PATH) {
-    $HIP_CLANG_PATH = "/opt/rocm/llvm/bin";
+    $HIP_CLANG_PATH = "$ROCM_PATH/llvm/bin";
   }
   if (!defined $DEVICE_LIB_PATH) {
-    $DEVICE_LIB_PATH = "/opt/rocm/lib";
+    $DEVICE_LIB_PATH = "$ROCM_PATH/lib";
   }
 }
 
@@ -163,7 +181,6 @@ $target_gfx1012 = 0;
 $default_amdgpu_target = 1;
 
 if ($HIP_PLATFORM eq "clang") {
-    $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
     $HIPCC="$HIP_CLANG_PATH/clang++";
 
     # If $HIPCC clang++ is not compiled, use clang instead
@@ -215,7 +232,7 @@ if ($HIP_PLATFORM eq "clang") {
     }
 
     if ($HIP_RUNTIME eq "HCC" ) {
-      $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
+      $HSA_PATH=$ENV{'HSA_PATH'} // "$ROCM_PATH/hsa";
       $HIPCXXFLAGS .= " -isystem $HSA_PATH/include";
     }
 
@@ -224,9 +241,9 @@ if ($HIP_PLATFORM eq "clang") {
     if (! defined $HIP_LIB_PATH) {
         $HIP_LIB_PATH = "$HIP_PATH/lib";
     }
-    $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
+    $HSA_PATH=$ENV{'HSA_PATH'} // "$ROCM_PATH/hsa";
 
-    $HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc";
+    $HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "$ROCM_PATH/hcc";
 
     $HCC_VERSION=`${HCC_HOME}/bin/hcc --version`;
     $HCC_VERSION=~/.*based on HCC ([^ ]+).*/;
@@ -234,8 +251,6 @@ if ($HIP_PLATFORM eq "clang") {
     $HCC_VERSION_MAJOR=$HCC_VERSION;
     $HCC_VERSION_MAJOR=~s/\..*//;
 
-    $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
-
     $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'} // 1;
     $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger";
 

bin/hipconfig

@@ -71,9 +71,17 @@ sub can_run {
     }
 }
 
+# Define HIP_PATH based on location of the script. Same as hipcc.
+# Derive ROCM_PATH same as hipcc does. Others are relative to ROCM_PATH. 
+$HIP_PATH=$ENV{'HIP_PATH'} // dirname(Cwd::abs_path("$0/../")); # use parent directory of hipcc
+if (-e "$HIP_PATH/../.info/version") {
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // dirname("$HIP_PATH"); # use parent directory of HIP_PATH
+} else {
+    $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
+}
 $CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda';
-$HCC_HOME=$ENV{'HCC_HOME'} // '/opt/rocm/hcc';
+$HCC_HOME=$ENV{'HCC_HOME'} // "$ROCM_PATH/hcc";
-$HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa';
+$HSA_PATH=$ENV{'HSA_PATH'} // "$ROCM_PATH/hsa";
 
 #---
 #HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
@@ -89,8 +97,6 @@ if (not defined $HIP_PLATFORM) {
     }
 }
 
-$HIP_PATH=$ENV{'HIP_PATH'} // Cwd::realpath (dirname (dirname $0)); # use parent directory of this tool
-
 if ($HIP_PLATFORM eq "hcc") {
     $CPP_CONFIG= " -D__HIP_PLATFORM_HCC__=   -I$HIP_PATH/include -I$HCC_HOME/include -I$HSA_PATH/include";
 }
@@ -181,7 +187,8 @@ if ($p_check) {
     print "\nCheck system installation:\n";
 
     printf ("%-70s", "check hipconfig in PATH...");
-    if (system ("hipconfig > /dev/null 2>&1") != 0)  {
+    # Safer to use which hipconfig instead of invoking hipconfig
+    if (system ("which hipconfig > /dev/null 2>&1") != 0)  {
         print "FAIL\n";
     } else {
         printf "good\n";

bin/hipify-perl

@@ -22,6 +22,8 @@
 # THE SOFTWARE.
 ##
 
+# IMPORTANT: Do not change this file manually: it is generated by hipify-clang --perl
+
 #usage hipify-perl [OPTIONS] INPUT_FILE
 
 use Getopt::Long;

hip_prof_gen.py

@@ -329,7 +329,6 @@ def generate_prof_header(f, api_map, opts_map):
     f.write('  HIP_API_ID_' + name + ' = ' + str(cb_id) + ',\n')
     cb_id += 1
   f.write('  HIP_API_ID_NUMBER = ' + str(cb_id) + ',\n')
-  f.write('  HIP_API_ID_ANY = ' + str(cb_id + 1) + ',\n')
   f.write('\n')
   f.write('  HIP_API_ID_NONE = HIP_API_ID_NUMBER,\n')
   for name in priv_lst:

hipify-clang/src/CUDA2HIP_Perl.cpp

@@ -60,6 +60,7 @@ namespace perl {
     "# THE SOFTWARE.\n"
     "##\n";
 
+  const string sImportant = "# IMPORTANT: Do not change this file manually: it is generated by hipify-clang --perl";
   const string tab = "    ";
   const string tab_2 = tab + tab;
   const string tab_3 = tab_2 + tab;
@@ -119,6 +120,7 @@ namespace perl {
   void generateHeader(unique_ptr<ostream>& streamPtr) {
     *streamPtr.get() << "#!/usr/bin/perl -w" << endl_2;
     *streamPtr.get() << sCopyright << endl;
+    *streamPtr.get() << sImportant << endl_2;
     *streamPtr.get() << "#usage " << hipify_perl << " [OPTIONS] INPUT_FILE" << endl_2;
     *streamPtr.get() << "use Getopt::Long;" << endl;
     *streamPtr.get() << my << "$whitelist = \"\";" << endl;

include/hip/hcc_detail/library_types.h

@@ -0,0 +1,42 @@
+/*
+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_HCC_DETAIL_LIBRARY_TYPES_H
+#define HIP_INCLUDE_HIP_HCC_DETAIL_LIBRARY_TYPES_H
+
+typedef enum hipDataType {
+  HIP_R_16F = 2,
+  HIP_R_32F = 0,
+  HIP_R_64F = 1,
+  HIP_C_16F = 6,
+  HIP_C_32F = 4,
+  HIP_C_64F = 5
+} hipDataType;
+
+typedef enum libraryPropertyType {
+  MAJOR_VERSION,
+  MINOR_VERSION,
+  PATCH_LEVEL
+} libraryPropertyType;
+
+#endif
+

include/hip/hip_runtime.h

@@ -63,5 +63,6 @@ THE SOFTWARE.
 
 #include <hip/hip_runtime_api.h>
 #include <hip/hip_vector_types.h>
+#include <hip/library_types.h>
 
 #endif

include/hip/library_types.h

@@ -0,0 +1,36 @@
+/*
+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_LIBRARY_TYPES_H
+#define HIP_INCLUDE_HIP_LIBRARY_TYPES_H
+
+#include <hip/hip_common.h>
+
+#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
+#include <hip/hcc_detail/library_types.h>
+#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__)
+#include "library_types.h"
+#else
+#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
+#endif
+
+#endif

include/hip/nvcc_detail/hip_runtime_api.h

@@ -50,6 +50,15 @@ typedef enum hipMemcpyKind {
     hipMemcpyDefault
 } hipMemcpyKind;
 
+// hipDataType
+#define hipDataType cudaDataType
+#define HIP_R_16F CUDA_R_16F
+#define HIP_R_32F CUDA_R_32F
+#define HIP_R_64F CUDA_R_64F
+#define HIP_C_16F CUDA_C_16F
+#define HIP_C_32F CUDA_C_32F
+#define HIP_C_64F CUDA_C_64F
+
 // hipTextureAddressMode
 #define hipTextureAddressMode cudaTextureAddressMode
 #define hipAddressModeWrap cudaAddressModeWrap

packaging/hip_base.postinst

@@ -17,9 +17,9 @@ mkdir -p $ROCMBINDIR
 pushd $ROCMBINDIR
 for f in $HIPBINFILES
 do
-    ln -s -f $f $(basename $f)
+    ln -r -s -f $f $(basename $f)
 done
-    ln -s -f $HIPDIR/bin/.hipVersion .hipVersion
+    ln -r -s -f $HIPDIR/bin/.hipVersion .hipVersion
 popd
 
 # Soft-link to headers
@@ -27,5 +27,5 @@ HIPINCDIR=$HIPDIR/include/hip
 ROCMINCDIR=$ROCMDIR/include
 mkdir -p $ROCMINCDIR
 pushd $ROCMINCDIR
-    ln -s -f $HIPINCDIR hip
+    ln -r -s -f $HIPINCDIR hip
 popd

packaging/hip_hcc.postinst

@@ -13,17 +13,17 @@ HIPDIR=$ROCMDIR/hip
 HIPLIBDIR=$ROCMDIR/hip/lib
 
 # Soft-link to library files
-HIPLIBFILES=$(ls -aF $HIPLIBDIR | grep -v [-/$])
+HIPLIBFILES=$(ls -A $HIPLIBDIR | grep -v [-/$])
 mkdir -p $ROCMLIBDIR
 mkdir -p $ROCMLIBDIR/cmake
 pushd $ROCMLIBDIR
 for f in $HIPLIBFILES
 do
-    ln -s -f $HIPLIBDIR/$f $(basename $f)
+    ln -s -r -f $HIPLIBDIR/$f $(basename $f)
 done
 # Make the hip cmake directory link.
 pushd cmake
-ln -s -f $HIPLIBDIR/cmake/hip hip
+ln -s -r -f $HIPLIBDIR/cmake/hip hip
 popd
 popd
 

packaging/hip_hcc.prerm

@@ -16,7 +16,7 @@ HIPLIBDIR=$ROCMDIR/hip/lib
 ([ ! -d $ROCMLIBDIR ] || [ ! -d $HIPLIBDIR ]) && exit 0
 
 # Remove soft-links to libraries
-HIPLIBFILES=$(ls -aF $HIPLIBDIR | grep -v [-/$])
+HIPLIBFILES=$(ls -A $HIPLIBDIR | grep -v [-/$])
 pushd $ROCMLIBDIR
 for f in $HIPLIBFILES; do
     [ -e $f ] || continue

src/hip_prof_api.h

@@ -38,9 +38,7 @@ class api_callbacks_table_templ {
   bool set_activity(uint32_t id, act_t fun, void* arg) {
     std::lock_guard<mutex_t> lock(mutex_);
     bool ret = true;
-    if (id == HIP_API_ID_ANY) {
+    if (id < HIP_API_ID_NUMBER) {
-      for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_activity(i, fun, arg);
-    } else if (id < HIP_API_ID_NUMBER) {
       cb_sync(id);
       callbacks_table_.arr[id].act = fun;
       callbacks_table_.arr[id].a_arg = arg;
@@ -54,9 +52,7 @@ class api_callbacks_table_templ {
   bool set_callback(uint32_t id, fun_t fun, void* arg) {
     std::lock_guard<mutex_t> lock(mutex_);
     bool ret = true;
-    if (id == HIP_API_ID_ANY) {
+    if (id < HIP_API_ID_NUMBER) {
-      for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_callback(i, fun, arg);
-    } else if (id < HIP_API_ID_NUMBER) {
       cb_sync(id);
       callbacks_table_.arr[id].fun = fun;
       callbacks_table_.arr[id].arg = arg;

tests/hipify-clang/unit_tests/samples/allocators.cu

@@ -1,9 +1,9 @@
 // RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
 
 #pragma once
-
+// CHECK: #include <hip/hip_runtime.h>
 #include <cuda_runtime.h>
-
+#include <math.h>
 
 /**
  * Allocate GPU memory for `count` elements of type `T`.
@@ -16,3 +16,38 @@ static T* gpuMalloc(size_t count) {
     return ret;
 }
 
+template<typename T>
+__global__ void add(int n, T* x, T* y) {
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = index; i < n; i += stride)
+    y[i] = x[i] + y[i];
+}
+
+int main(int argc, char* argv[]) {
+    size_t numElements = 50;
+    float *A = gpuMalloc<float>(numElements);
+    float* B = gpuMalloc<float>(numElements);
+    for (int i = 0; i < numElements; ++i) {
+        A[i] = 1.0f;
+        B[i] = 2.0f;
+    }
+    int blockSize = 512;
+    int numBlocks = (numElements + blockSize - 1) / blockSize;
+    dim3 dimGrid(numBlocks, 1, 1);
+    dim3 dimBlock(blockSize, 1, 1);
+    // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(add<float>), dim3(dimGrid), dim3(dimBlock), 0, 0, numElements, A, B);
+    add<float><<<dimGrid, dimBlock>>>(numElements, A, B);
+    // CHECK: hipDeviceSynchronize();
+    cudaDeviceSynchronize();
+    float maxError = 0.0f;
+    for (int i = 0; i < numElements; ++i)
+      maxError = fmax(maxError, fabs(B[i] - 3.0f));
+    // CHECK: hipFree(A);
+    cudaFree(A);
+    // CHECK: hipFree(B);
+    cudaFree(B);
+    if (maxError == 0.0f)
+      return 0;
+    return -1;
+  }

tests/src/hiprtc/hiprtcGetLoweredName.cpp

@@ -20,7 +20,7 @@ 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 LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -80,7 +80,16 @@ int main()
 
     for (auto&& x : variable_name_vec) hiprtcAddNameExpression(prog, x.c_str());
 
-    hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr);
+    hipDeviceProp_t props;
+    int device = 0;
+    hipGetDeviceProperties(&props, device);
+    std::string gfxName = "gfx" + std::to_string(props.gcnArch);
+    std::string sarg = "--gpu-architecture=" + gfxName;
+    const char* options[] = {
+            sarg.c_str()
+    };
+
+    hiprtcResult compileResult = hiprtcCompileProgram(prog, 1, options);
 
     // Obtain compilation log from the program.
     size_t logSize;

tests/src/hiprtc/hiprtcGetTypeName.cpp

@@ -20,7 +20,7 @@ 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 LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -76,7 +76,16 @@ int main()
 
     for (auto&& x : name_vec) hiprtcAddNameExpression(prog, x.c_str());
 
-    hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr);
+    hipDeviceProp_t props;
+    int device = 0;
+    hipGetDeviceProperties(&props, device);
+    std::string gfxName = "gfx" + std::to_string(props.gcnArch);
+    std::string sarg = "--gpu-architecture=" + gfxName;
+    const char* options[] = {
+            sarg.c_str()
+    };
+
+    hiprtcResult compileResult = hiprtcCompileProgram(prog, 1, options);
 
     size_t logSize;
     hiprtcGetProgramLogSize(prog, &logSize);

tests/src/hiprtc/saxpy.cpp

@@ -20,7 +20,7 @@ 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 LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -66,7 +66,16 @@ int main()
                         nullptr,    // headers
                         nullptr);   // includeNames
 
-    hiprtcResult compileResult{hiprtcCompileProgram(prog, 0, nullptr)};
+    hipDeviceProp_t props;
+    int device = 0;
+    hipGetDeviceProperties(&props, device);
+    std::string gfxName = "gfx" + std::to_string(props.gcnArch);
+    std::string sarg = "--gpu-architecture=" + gfxName;
+    const char* options[] = {
+        sarg.c_str()
+    };
+
+    hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
 
     size_t logSize;
     hiprtcGetProgramLogSize(prog, &logSize);

tests/src/runtimeApi/memory/hipArray.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
  * TEST: %t
  * HIT_END
  */

tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 // Simple test for hipLaunchCooperativeKernel API.
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
  * TEST: %t
  * HIT_END
  */