From a474dc7579103de358dd2e60492a7609a1895ef1 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 7 Feb 2019 11:07:50 -0500 Subject: [PATCH 01/12] Fix hipcc and hip_vector_types.h for windows Disable linux specific compilation and linking options for windows. In hip_vector_types.h, hip-clang needs HIP's own implementation for vector types even on windows. MSVC specific implementation is for building runtime itself, which is compiled by MSVC. --- bin/hipcc | 22 +++++++++++++++------- include/hip/hcc_detail/hip_vector_types.h | 2 +- 2 files changed, 16 insertions(+), 8 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index d422174ad8..80106f06aa 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -53,6 +53,8 @@ sub parse_config_file { $verbose = $ENV{'HIPCC_VERBOSE'} // 0; # Verbose: 0x1=commands, 0x2=paths, 0x4=hipcc args +$isWindows = $^O eq 'MSWin32'; + $HIPCC_COMPILE_FLAGS_APPEND=$ENV{'HIPCC_COMPILE_FLAGS_APPEND'}; $HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'}; @@ -182,7 +184,11 @@ if ($HIP_PLATFORM eq "clang") { # is required. $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH"; - $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS -L$HIP_LIB_PATH -Wl,--rpath=$HIP_LIB_PATH -lhip_hcc"; + $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS -L$HIP_LIB_PATH"; + if (not $isWindows) { + $HIPLDFLAGS .= " -Wl,--rpath=$HIP_LIB_PATH"; + } + $HIPLDFLAGS .= " -lhip_hcc"; if ($HIP_CLANG_HCC_COMPAT_MODE) { ## Allow __fp16 as function parameter and return type. $HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1"; @@ -262,7 +268,9 @@ if ($HIP_PLATFORM eq "clang") { $HIPLDFLAGS .= " -L$marker_lib_path -lCXLActivityLogger -Wl,--rpath=$marker_lib_path"; } - $HIPLDFLAGS .= " -lm"; + if (not $isWindows) { + $HIPLDFLAGS .= " -lm"; + } if ($verbose & 0x2) { print ("HSA_PATH=$HSA_PATH\n"); @@ -760,13 +768,13 @@ if ($needHipHcc) { # pass-through CPP mode. if ($HIP_PLATFORM eq "clang") { - $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm"; + if ($isWindows) { + $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility"; + } else { + $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm"; + } } -# Windows specific options -if ($^O eq 'MSWin32' and $HIP_PLATFORM eq "clang") { - $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility" -} if ($HIPCC_COMPILE_FLAGS_APPEND) { $HIPCXXFLAGS .= " $HIPCC_COMPILE_FLAGS_APPEND"; diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index acfb16423f..1df6385fa6 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -34,7 +34,7 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -#if !defined(_MSC_VER) || __HIP_DEVICE_COMPILE__ +#if !defined(_MSC_VER) || __clang__ #if defined(__clang__) #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) #elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax. From ee4ce7ae496df6ab57e3632b5b93e5e5360e0b58 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 7 Feb 2019 17:29:36 -0500 Subject: [PATCH 02/12] Fix nan() for windows --- include/hip/hcc_detail/math_functions.h | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 6374375491..8ac87425b9 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -899,6 +899,7 @@ __DEVICE__ inline double nan(const char* tagp) { +#if !_WIN32 union { double val; struct ieee_double { @@ -906,8 +907,7 @@ double nan(const char* tagp) uint32_t quiet : 1; uint32_t exponent : 11; uint32_t sign : 1; - } bits; - + } bits; static_assert(sizeof(double) == sizeof(ieee_double), ""); } tmp; @@ -917,6 +917,11 @@ double nan(const char* tagp) tmp.bits.mantissa = __make_mantissa(tagp); return tmp.val; +#else + uint64_t val = __make_mantissa(tagp); + val |= 0xFFF << 51; + return reinterpret_cast(val); +#endif } __DEVICE__ inline From 31dd6bc83032386d34732aa7ee56d765d635bb41 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 5 Feb 2019 16:06:42 +0530 Subject: [PATCH 03/12] [packaging] Add checkpoints in prerm scripts - Add failsafe in case of non existant files or directories. Change-Id: I85db04381c3e16f04638086152945463a8004c19 --- packaging/hip_base.prerm | 12 ++++++++---- packaging/hip_hcc.prerm | 8 +++++--- 2 files changed, 13 insertions(+), 7 deletions(-) diff --git a/packaging/hip_base.prerm b/packaging/hip_base.prerm index 5985cc30ec..837008f3e8 100755 --- a/packaging/hip_base.prerm +++ b/packaging/hip_base.prerm @@ -10,14 +10,17 @@ popd () { ROCMDIR=/opt/rocm HIPDIR=$ROCMDIR/hip +([ ! -d $ROCMDIR ] || [ ! -d $HIPDIR ]) && exit 0 # Remove soft-link to bin files -HIPBINFILES=$HIPDIR/bin/* +HIPBINDIR=$HIPDIR/bin ROCMBINDIR=$ROCMDIR/bin +([ ! -d $ROCMBINDIR ] || [ ! -d $HIPBINDIR ]) && exit 0 +HIPBINFILES=$HIPBINDIR/* pushd $ROCMBINDIR -for f in $HIPBINFILES -do -rm $(basename $f) +for f in $HIPBINFILES; do + [ -e $f ] || continue + rm $(basename $f) done rm .hipVersion popd @@ -26,6 +29,7 @@ rmdir --ignore-fail-on-non-empty $ROCMBINDIR # Remove soft-link to headers HIPINCDIR=$HIPDIR/include ROCMINCDIR=$ROCMDIR/include +([ ! -d $ROCMINCDIR ] || [ ! -d $HIPINCDIR ]) && exit 0 pushd $ROCMINCDIR rm hip popd diff --git a/packaging/hip_hcc.prerm b/packaging/hip_hcc.prerm index ee64aea632..8d7d8b93b2 100755 --- a/packaging/hip_hcc.prerm +++ b/packaging/hip_hcc.prerm @@ -12,13 +12,15 @@ ROCMDIR=/opt/rocm ROCMLIBDIR=$ROCMDIR/lib HIPDIR=$ROCMDIR/hip HIPLIBDIR=$ROCMDIR/hip/lib +([ ! -d $ROCMDIR ] || [ ! -d $HIPDIR ]) && exit 0 +([ ! -d $ROCMLIBDIR ] || [ ! -d $HIPLIBDIR ]) && exit 0 # Remove soft-links to libraries HIPLIBFILES=$(ls -aF $HIPLIBDIR | grep -v [-/$]) pushd $ROCMLIBDIR -for f in $HIPLIBFILES -do -rm $(basename $f) +for f in $HIPLIBFILES; do + [ -e $f ] || continue + rm $(basename $f) done pushd cmake unlink hip From 7ba8c440e6837d7c7c9e73a87810ee1dadaf0001 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 8 Feb 2019 17:54:28 +0300 Subject: [PATCH 04/12] [HIPIFY][RAND] revise + add missing functions + add minimum rocRAND support + updated CURAND_API_supported_by_HIP.md --- docs/markdown/CURAND_API_supported_by_HIP.md | 2 ++ hipify-clang/src/CUDA2HIP.cpp | 4 ++-- hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp | 2 ++ hipify-clang/src/HipifyAction.cpp | 6 ++++-- .../libraries/cuRAND/benchmark_curand_kernel.cpp | 10 ++++------ 5 files changed, 14 insertions(+), 10 deletions(-) diff --git a/docs/markdown/CURAND_API_supported_by_HIP.md b/docs/markdown/CURAND_API_supported_by_HIP.md index c35eeb26b6..0561dd6144 100644 --- a/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/docs/markdown/CURAND_API_supported_by_HIP.md @@ -134,6 +134,8 @@ | `curandSetPseudoRandomGeneratorSeed` | `hiprandSetPseudoRandomGeneratorSeed` | | `curandSetQuasiRandomGeneratorDimensions` | `hiprandSetQuasiRandomGeneratorDimensions` | | `curandSetStream` | `hiprandSetStream` | +| `curandMakeMTGP32Constants` | `hiprandMakeMTGP32Constants` | +| `curandMakeMTGP32KernelState` | `hiprandMakeMTGP32KernelState` | ## **3. Device API Functions** diff --git a/hipify-clang/src/CUDA2HIP.cpp b/hipify-clang/src/CUDA2HIP.cpp index 920b310da8..7536c5c0f7 100644 --- a/hipify-clang/src/CUDA2HIP.cpp +++ b/hipify-clang/src/CUDA2HIP.cpp @@ -48,9 +48,9 @@ const std::map CUDA_INCLUDE_MAP{ {"curand_lognormal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mrg32k3a.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, - {"curand_mtgp32_host.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32_host.h", {"hiprand_mtgp32_host.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, - {"curand_mtgp32dc_p_11213.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32dc_p_11213.h", {"rocrand_mtgp32_11213.h", "", CONV_INCLUDE, API_RAND}}, {"curand_normal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_normal_static.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_philox4x32_x.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, diff --git a/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp b/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp index 65985d6f5c..525aa1e5f3 100644 --- a/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp +++ b/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp @@ -51,6 +51,8 @@ const std::map CUDA_RAND_FUNCTION_MAP{ {"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", "", CONV_LIB_FUNC, API_RAND}}, {"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", "", CONV_LIB_FUNC, API_RAND}}, {"curandSetStream", {"hiprandSetStream", "", CONV_LIB_FUNC, API_RAND}}, + {"curandMakeMTGP32Constants", {"hiprandMakeMTGP32Constants", "", CONV_LIB_FUNC, API_RAND}}, + {"curandMakeMTGP32KernelState", {"hiprandMakeMTGP32KernelState", "", CONV_LIB_FUNC, API_RAND}}, // RAND Device functions {"curand", {"hiprand", "", CONV_LIB_DEVICE_FUNC, API_RAND}}, diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index f4d68cd820..baa66e064c 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -191,8 +191,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) { case CONV_INCLUDE: switch (hipToken.apiType) { case API_RAND: - if (insertedRAND_kernelHeader) { return true; } - insertedRAND_kernelHeader = true; + if (hipToken.hipName == "hiprand_kernel.h") { + if (insertedRAND_kernelHeader) { return true; } + insertedRAND_kernelHeader = true; + } return false; default: return false; diff --git a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp index 9cb1a72420..b1cdec702a 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp +++ b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp @@ -37,9 +37,9 @@ #include // CHECK: #include #include -// CHECK-NOT: #include -// CHECK-NOT: #include +// CHECK: #include #include +// CHECK: #include #include // CHECK: if ((x) != hipSuccess) { @@ -201,11 +201,9 @@ struct runner CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateMtgp32_t))); // CHECK: CUDA_CALL(hipMalloc((void **)&d_param, sizeof(mtgp32_kernel_params))); CUDA_CALL(cudaMalloc((void **)&d_param, sizeof(mtgp32_kernel_params))); - // curandMakeMTGP32Constants is yet unsupported by HIP - // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); + // CHECK: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); CURAND_CALL(curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); - // curandMakeMTGP32KernelState is yet unsupported by HIP - // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); + // CHECK: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); } From eb33507774444a984e6f8bed9671a0b0e78c7a07 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 8 Feb 2019 18:29:14 +0000 Subject: [PATCH 05/12] hipcc should consume -mcode-object-v3 flag --- bin/hipcc | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/bin/hipcc b/bin/hipcc index d422174ad8..d173d761dd 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -300,6 +300,7 @@ my $runCmd = 1; my $buildDeps = 0; my $linkType = 1; my $setLinkType = 0; +my $useCodeObjectV3 = 0; my @options = (); my @inputs = (); @@ -403,6 +404,10 @@ foreach $arg (@ARGV) $swallowArg = 1; } + if (($trimarg eq '-mcode-object-v3') and ($HIP_PLATFORM eq 'clang') ) { + $useCodeObjectV3 = 1; + } + if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) { $arg = "--cuda-device-only"; } @@ -736,6 +741,12 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') { $HIPCXXFLAGS .= " --cuda-host-only"; } +if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= " -mcode-object-v3"; +} else { + $HIPCXXFLAGS .= " -mno-code-object-v3"; +} + # Add --hip-link only if there are no source files. if (!$needCXXFLAGS and $HIP_PLATFORM eq 'clang') { $HIPLDFLAGS .= " --hip-link"; From 4eaf4b7c0a91600ffbd59b0c1192cb44b051243e Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Fri, 8 Feb 2019 13:50:13 -0800 Subject: [PATCH 06/12] Allow HIP_LIB_PATH to be overriden at the command line. Use HIP_LIB_PATH for rpath instead of HIP_PATH/lib --- bin/hipcc | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 29eb51d5ab..c6e9baaec6 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -58,6 +58,7 @@ $HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'}; $HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc $HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'}; +$HIP_LIB_PATH=$ENV{'HIP_LIB_PATH'}; $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'}; $DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'}; $HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibility mode @@ -120,7 +121,9 @@ if (defined $HIP_VDI_HOME) { } $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang"; $HIP_INCLUDE_PATH = "$HIP_VDI_HOME/include"; - $HIP_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits; + if (!defined $HIP_LIB_PATH) { + $HIP_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits; + } } if ($HIP_COMPILER eq "clang") { @@ -754,9 +757,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc') if ($needHipHcc) { if ($linkType eq 0) { - substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a " ; + substr($HIPLDFLAGS,0,0) = " $HIP_LIB_PATH/libhip_hcc_static.a " ; } else { - substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so "; + substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_LIB_PATH $HIP_PATH/lib/libhip_hcc.so "; } } From eba3bd30823c1178c15b72653b76197f800b5b79 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Mon, 11 Feb 2019 09:31:30 -0600 Subject: [PATCH 07/12] Introduce __HIP_DISABLE_CPP_FUNCTIONS__ macro to disable C++ wrappers hipMalloc and hipHostMalloc are C++ function templates by default. In some applications they need to discover the decltype of underlying C functions. In this commit, a new macro __HIP_DISABLE_CPP_FUNCTIONS__ is introduced to suppress the C++ function templates. Existing HIP applications don't need to be changed. hipMalloc and hipHostMalloc will behave just the same. For new applications which depends on identifying decltypes of underlying C functions, __HIP_DISABLE_CPP_FUNCTIONS__ shall be defined prior to including HIP headers. --- include/hip/hip_runtime_api.h | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index cd7af65265..34363689e4 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -318,9 +318,13 @@ enum hipComputeMode { * * Perform automatic type conversion to eliminate need for excessive typecasting (ie void**) * + * __HIP_DISABLE_CPP_FUNCTIONS__ macro can be defined to suppress these + * wrappers. It is useful for applications which need to obtain decltypes of + * HIP runtime APIs. + * * @see hipMalloc */ -#ifdef __cplusplus +#if defined(__cplusplus) && !defined(__HIP_DISABLE_CPP_FUNCTIONS__) template static inline hipError_t hipMalloc(T** devPtr, size_t size) { return hipMalloc((void**)devPtr, size); From 1c6365fbe77c4d15b83f08f3a32b7bd3f1ee2640 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 11 Feb 2019 17:18:12 +0000 Subject: [PATCH 08/12] Change code-object flag to only HIP-Clang --- bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index d173d761dd..792d1da59a 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -743,7 +743,7 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') { if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') { $HIPCXXFLAGS .= " -mcode-object-v3"; -} else { +} elsif ($HIP_PLATFORM eq 'clang') { $HIPCXXFLAGS .= " -mno-code-object-v3"; } From 63b400beea2a2d273b1d253167bcf1cc039b3b38 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 12 Feb 2019 16:56:22 +0000 Subject: [PATCH 09/12] Fix texture using constant addrspace 4 --- include/hip/hcc_detail/texture_functions.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index c293d65558..1f22e01177 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/include/hip/hcc_detail/texture_functions.h @@ -50,7 +50,7 @@ union TData { #define __TEXTURE_FUNCTIONS_DECL__ static inline __device__ -#if (__hcc_workweek__ >= 18114) +#if (__hcc_workweek__ >= 18114) || __clang__ #define ADDRESS_SPACE_CONSTANT __attribute__((address_space(4))) #else #define ADDRESS_SPACE_CONSTANT __attribute__((address_space(2))) From f0618df8f40c3bc855e7b7982ef2a8f6e2ea9a05 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 13 Feb 2019 09:45:07 +0530 Subject: [PATCH 10/12] Fix signature for unsigned long long variant of atomicExch This fixes #888 Change-Id: I2b2dc5ed1ff2a26083ad9d102eccac90f67801ae --- include/hip/hcc_detail/hip_atomic.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_atomic.h b/include/hip/hcc_detail/hip_atomic.h index ce4b7c9008..a5ac94a74b 100644 --- a/include/hip/hcc_detail/hip_atomic.h +++ b/include/hip/hcc_detail/hip_atomic.h @@ -112,7 +112,7 @@ unsigned int atomicExch(unsigned int* address, unsigned int val) } __device__ inline -unsigned int atomicExch(unsigned long long* address, unsigned long long val) +unsigned long long atomicExch(unsigned long long* address, unsigned long long val) { return __atomic_exchange_n(address, val, __ATOMIC_RELAXED); } From c8579ead673ad77527a2d2420eb6a7d92334c52d Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 13 Feb 2019 09:57:21 +0530 Subject: [PATCH 11/12] Initialize HIP_LIB_PATH on hcc path Change-Id: I8f392fffcd6a3714c20cb17e876dbaad999c33a0 --- bin/hipcc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index ac4cde2611..b649bb6e7f 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -204,6 +204,9 @@ if ($HIP_PLATFORM eq "clang") { } } elsif ($HIP_PLATFORM eq "hcc") { $HIP_INCLUDE_PATH = "$HIP_PATH/include"; + if (! defined $HIP_LIB_PATH) { + $HIP_LIB_PATH = "$HIP_PATH/lib"; + } $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa"; $HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc"; @@ -778,7 +781,7 @@ if ($needHipHcc) { if ($linkType eq 0) { substr($HIPLDFLAGS,0,0) = " $HIP_LIB_PATH/libhip_hcc_static.a " ; } else { - substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_LIB_PATH $HIP_PATH/lib/libhip_hcc.so "; + substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_LIB_PATH $HIP_LIB_PATH/libhip_hcc.so "; } } From 336a781a2d80ab40c8d36ba1b73a2da2c435e9e2 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 13 Feb 2019 13:22:08 +0530 Subject: [PATCH 12/12] [ci] Switch CI tests to cuda-10.x Change-Id: I99e99bcf528d0c38924bd3a86a3067628f2bb3ea --- Jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 0be0bf29bd..71e8d6e537 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -295,7 +295,7 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// The following launches 3 builds in parallel: rocm-head, rocm-2.1.x and cuda-9.x +// The following launches 3 builds in parallel: rocm-head, rocm-2.1.x and cuda-10.x parallel rocm_2_1: { node('hip-rocm') @@ -386,14 +386,14 @@ rocm_head: */ } }, -cuda_9_x: +cuda_10_x: { node('hip-cuda') { //////////////////////////////////////////////////////////////////////// // Block of string constants customizing behavior for cuda - String nvcc_ver = 'nvcc-9.x' - String from_image = 'ci_test_nodes/cuda-9.x/ubuntu-16.04:latest' + String nvcc_ver = 'cuda-10.x' + String from_image = 'ci_test_nodes/cuda-10.x/ubuntu-16.04:latest' String inside_args = '--runtime=nvidia'; // Checkout source code, dependencies and version files