diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 81f394102a..35fbb54397 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/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. -# 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 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"; diff --git a/projects/hip/bin/hipconfig b/projects/hip/bin/hipconfig index 87da86f3b1..e4bf118c61 100755 --- a/projects/hip/bin/hipconfig +++ b/projects/hip/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'; -$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"; diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 8116cfbe7c..25e02352ad 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/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; diff --git a/projects/hip/hip_prof_gen.py b/projects/hip/hip_prof_gen.py index 6a1564cbb9..d1203a64d6 100755 --- a/projects/hip/hip_prof_gen.py +++ b/projects/hip/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: diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index 9bbb52a2fa..df83bf8d57 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/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& 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; diff --git a/projects/hip/include/hip/hcc_detail/library_types.h b/projects/hip/include/hip/hcc_detail/library_types.h new file mode 100644 index 0000000000..74bae81321 --- /dev/null +++ b/projects/hip/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 + diff --git a/projects/hip/include/hip/hip_runtime.h b/projects/hip/include/hip/hip_runtime.h index 937ba61ecf..b02eb0eef7 100644 --- a/projects/hip/include/hip/hip_runtime.h +++ b/projects/hip/include/hip/hip_runtime.h @@ -63,5 +63,6 @@ THE SOFTWARE. #include #include +#include -#endif \ No newline at end of file +#endif diff --git a/projects/hip/include/hip/library_types.h b/projects/hip/include/hip/library_types.h new file mode 100644 index 0000000000..4a988df52b --- /dev/null +++ b/projects/hip/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 + +#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__) +#include +#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 diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 97551d15d4..02f83b329e 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/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 diff --git a/projects/hip/packaging/hip_base.postinst b/projects/hip/packaging/hip_base.postinst index 7dfb0369d7..a8b3eed9fc 100755 --- a/projects/hip/packaging/hip_base.postinst +++ b/projects/hip/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 diff --git a/projects/hip/packaging/hip_hcc.postinst b/projects/hip/packaging/hip_hcc.postinst index b7c3e628ec..1cf21dd755 100755 --- a/projects/hip/packaging/hip_hcc.postinst +++ b/projects/hip/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 diff --git a/projects/hip/packaging/hip_hcc.prerm b/projects/hip/packaging/hip_hcc.prerm index 8d7d8b93b2..aab670eb4f 100755 --- a/projects/hip/packaging/hip_hcc.prerm +++ b/projects/hip/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 diff --git a/projects/hip/src/hip_prof_api.h b/projects/hip/src/hip_prof_api.h index eb3112bdb4..8a69746f93 100644 --- a/projects/hip/src/hip_prof_api.h +++ b/projects/hip/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 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 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; diff --git a/projects/hip/tests/hipify-clang/unit_tests/samples/allocators.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/allocators.cu index a1bce9d4cc..4d33e315bf 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/samples/allocators.cu +++ b/projects/hip/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 #include - +#include /** * Allocate GPU memory for `count` elements of type `T`. @@ -16,3 +16,38 @@ static T* gpuMalloc(size_t count) { return ret; } +template +__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(numElements); + float* B = gpuMalloc(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), dim3(dimGrid), dim3(dimBlock), 0, 0, numElements, A, B); + add<<>>(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; + } diff --git a/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp index 8041aa73a3..e7b88d26d2 100644 --- a/projects/hip/tests/src/hiprtc/hiprtcGetLoweredName.cpp +++ b/projects/hip/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; diff --git a/projects/hip/tests/src/hiprtc/hiprtcGetTypeName.cpp b/projects/hip/tests/src/hiprtc/hiprtcGetTypeName.cpp index 0263dca28b..812229f81f 100644 --- a/projects/hip/tests/src/hiprtc/hiprtcGetTypeName.cpp +++ b/projects/hip/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); diff --git a/projects/hip/tests/src/hiprtc/saxpy.cpp b/projects/hip/tests/src/hiprtc/saxpy.cpp index 268d718b27..191e48b846 100644 --- a/projects/hip/tests/src/hiprtc/saxpy.cpp +++ b/projects/hip/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); diff --git a/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp b/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp index 9b3b18521e..d99ba6aee3 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp +++ b/projects/hip/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 */ diff --git a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp index 7680229855..89d003ea94 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp +++ b/projects/hip/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 */