Merge 'master' into 'amd-master'

Change-Id: I9e11a66df0398a3794eaeaf0d3994da225649be0 [ROCm/hip commit: 0402ffdfd3]
2019-10-14 05:22:07 -04:00
commit 94ebeed6da
@@ -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.
-# HSA_PATH  : Path to HSA dir (default /opt/rocm/hsa).  Used on AMD platforms only.
+# HCC_HOME  : Path to HCC SDK (defaults to ../../hcc relative to this
+#             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";

@@ -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';
-$HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa';
+$HCC_HOME=$ENV{'HCC_HOME'} // "$ROCM_PATH/hcc";
+$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";
@@ -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;
@@ -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:
@@ -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;
@@ -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
+
@@ -63,5 +63,6 @@ THE SOFTWARE.

 #include <hip/hip_runtime_api.h>
 #include <hip/hip_vector_types.h>
+#include <hip/library_types.h>

-#endif
+#endif
@@ -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
@@ -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
@@ -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
@@ -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

@@ -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
@@ -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) {
-      for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_activity(i, fun, arg);
-    } else if (id < HIP_API_ID_NUMBER) {
+    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) {
-      for (unsigned i = 0; i < HIP_API_ID_NUMBER; ++i) set_callback(i, fun, arg);
-    } else if (id < HIP_API_ID_NUMBER) {
+    if (id < HIP_API_ID_NUMBER) {
      cb_sync(id);
      callbacks_table_.arr[id].fun = fun;
      callbacks_table_.arr[id].arg = arg;
@@ -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;
+  }
@@ -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;
@@ -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);
@@ -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);
@@ -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
 */
@@ -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
 */