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 diff --git a/bin/hipcc b/bin/hipcc index 29eb51d5ab..b649bb6e7f 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -53,11 +53,14 @@ 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'}; $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 +123,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") { @@ -188,13 +193,20 @@ 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"; } } 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"; @@ -268,7 +280,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"); @@ -306,6 +320,7 @@ my $runCmd = 1; my $buildDeps = 0; my $linkType = 1; my $setLinkType = 0; +my $useCodeObjectV3 = 0; my @options = (); my @inputs = (); @@ -409,6 +424,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"; } @@ -742,6 +761,12 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') { $HIPCXXFLAGS .= " --cuda-host-only"; } +if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= " -mcode-object-v3"; +} elsif ($HIP_PLATFORM eq 'clang') { + $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"; @@ -754,9 +779,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_LIB_PATH/libhip_hcc.so "; } } @@ -766,13 +791,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/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/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); } 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. 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 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))) 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); 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 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)); }