From 11ad6077ae065c0493f5e9d98a6c66cb3855d17b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sat, 2 Jun 2018 18:01:09 -0400 Subject: [PATCH 01/20] Add more function declarations for hip-clang [ROCm/hip commit: d726ba1d9ac5ea19559ec481f171d5e1e434d6a6] --- projects/hip/include/hip/hcc_detail/hip_runtime.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 1a6b0f7dda..7c1f2a4d16 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -98,6 +98,9 @@ struct Empty_launch_parm {}; #include "grid_launch_GGL.hpp" #endif // GENERIC_GRID_LAUNCH +#endif // HCC + +#if __HCC_OR_HIP_CLANG__ extern int HIP_TRACE_API; #ifdef __cplusplus @@ -106,15 +109,16 @@ extern int HIP_TRACE_API; #include #include #include +#if __HCC___ #include #include +#endif // __HCC__ // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define. #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__) #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__ #endif - // TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call. #if __HIP_DEVICE_COMPILE__ == 1 #undef assert @@ -179,10 +183,6 @@ extern int HIP_TRACE_API; #define __HCC_C__ #endif -#endif // defined __HCC__ - -#if __HCC_OR_HIP_CLANG__ - // TODO - hipify-clang - change to use the function call. //#define warpSize hc::__wavesize() static constexpr int warpSize = 64; From 0573ee79de3781d9a08bd3a6fd24724d9f8fbe4b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 13 Jun 2018 10:01:14 -0400 Subject: [PATCH 02/20] Let hipcc handle static library for hip-clang only if it contains bundles [ROCm/hip commit: f0935e6d4f424c82f065c0c22ca5b48e6ed2099a] --- projects/hip/bin/hipcc | 21 +++++++++++++++++++-- 1 file changed, 19 insertions(+), 2 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 03f35b27fc..ca504fae4e 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -377,14 +377,31 @@ foreach $arg (@ARGV) while (my $line = <$in>) { chomp $line; if ($line =~ m/\.a$/) { + my $libFile = $line; my $path = abs_path($line); my @objs = split ('\n', `cd $tmpdir; ar xv $path`); + ## Check if all files in .a are object files. + my $allIsObj = 1; + my $realObjs = ""; foreach my $obj (@objs) { chomp $obj; $obj =~ s/^x - //; $obj = "$tmpdir/$obj"; - push (@inputs, $obj); - $new_arg = "$new_arg $obj"; + my $fileType = `file $obj`; + my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/); + $allIsObj = ($allIsObj and $isObj); + if ($isObj) { + $realObjs = $realObjs . " " . $obj; + } else { + push (@inputs, $obj); + $new_arg = "$new_arg $obj"; + } + } + if ($allIsObj) { + print $out "$line\n"; + } else { + system("cd $tmpdir; ar c $libFile $realObjs"); + print $out "$tmpdir/$libFile\n"; } } else { print $out "$line\n"; From acaf783f5269bb435150abfd8bfbf6e448aa1614 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sun, 17 Jun 2018 12:18:37 -0400 Subject: [PATCH 03/20] Fix handling of static library in hipcc for hip-clang [ROCm/hip commit: 2b32dbd414226adcac8c8b31211031128fffd6a8] --- projects/hip/bin/hipcc | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index ca504fae4e..4c526bd2c7 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -391,17 +391,20 @@ foreach $arg (@ARGV) my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/); $allIsObj = ($allIsObj and $isObj); if ($isObj) { - $realObjs = $realObjs . " " . $obj; + $realObjs = ($realObjs . " " . $obj); } else { push (@inputs, $obj); $new_arg = "$new_arg $obj"; } } + chomp $realObjs; if ($allIsObj) { print $out "$line\n"; - } else { - system("cd $tmpdir; ar c $libFile $realObjs"); - print $out "$tmpdir/$libFile\n"; + } elsif ($realObjs) { + my $libBaseName = basename($libFile, ".a"); + $libBaseName = mktemp($libBaseName . "XXXX") . ".a"; + system("cd $tmpdir; ar c $libBaseName $realObjs"); + print $out "$tmpdir/$libBaseName\n"; } } else { print $out "$line\n"; From 0b10ab2b7e4b7671d923bbdc98abc6e01deb2902 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 5 Jun 2018 08:46:20 -0400 Subject: [PATCH 04/20] Add missing macro MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT to hip_complex.h [ROCm/hip commit: 5eeb57b0a6b2113bf5ec5fdcb333f6834522698c] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 973d5f564b..22808ed576 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -94,6 +94,9 @@ THE SOFTWARE. ret.y = lhs.y * rhs; \ return ret; \ } +#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ + ComplexT(T val) : x(val), y(val) {} \ + ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif From cf6cdab0294a95f91448accf1407589a5b61cf0f Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 31 May 2018 13:58:44 -0400 Subject: [PATCH 05/20] Include cmath instead of math.h in hip_complex.h [ROCm/hip commit: 325cf3ccf0ec626bed2dccfacef6176e3e40cc35] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 22808ed576..3dc7a2ed56 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -24,7 +24,7 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "hip/hcc_detail/hip_vector_types.h" -#include +#include #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ From 0e29f327e2e2507986a4dca9231b87ce1b0654c7 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 6 Jun 2018 11:15:54 -0400 Subject: [PATCH 06/20] Includes or by __cplusplus in hip_complex.h [ROCm/hip commit: 2523c39a3706fc8d95d196a1eb4179497dda4a7e] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 3dc7a2ed56..72f139ba84 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -24,7 +24,15 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "hip/hcc_detail/hip_vector_types.h" + +// TODO: Clang has a bug which allows device functions to call std functions +// when std functions are introduced into default namespace by using statement. +// math.h may be included after this bug is fixed. +#if __cplusplus #include +#else +#include "math.h" +#endif #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ From fa9e73ccdc3be16d5c2564a7c238967647515317 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sun, 17 Jun 2018 12:18:05 -0400 Subject: [PATCH 07/20] Add missing __device__ __host__ to complex constructor Also add missing typedef value_type [ROCm/hip commit: 7a5605d006862c423710171f33641be7f54432ac] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 72f139ba84..783833dc7e 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -103,14 +103,15 @@ THE SOFTWARE. return ret; \ } #define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ - ComplexT(T val) : x(val), y(val) {} \ - ComplexT(T val1, T val2) : x(val1), y(val2) {} + __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ + __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif struct hipFloatComplex { #ifdef __cplusplus public: + typedef float value_type; __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {} __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {} __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {} @@ -130,6 +131,7 @@ struct hipFloatComplex { struct hipDoubleComplex { #ifdef __cplusplus public: + typedef double value_type; __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {} From a46f62a5c05aae9d2509d1aaa376711e5312ca7b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 11:57:57 -0400 Subject: [PATCH 08/20] Add abs/real/imag functions for hipFloatComplex/hipDoubleComplex [ROCm/hip commit: 9181fbb0b7e5f7db4df486b8600ea78092c4ffb1] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 783833dc7e..7f4dd840b7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -234,7 +234,6 @@ __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hi __device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); } - __device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; } __device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; } @@ -314,4 +313,12 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } +#define __DEFINE_HIP_COMPLEX_FUN(func) \ +__device__ __host__ inline float func(const hipFloatComplex& z) { return hipC##func##f(z); } \ +__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipC##func(z); } + +__DEFINE_HIP_COMPLEX_FUN(abs) +__DEFINE_HIP_COMPLEX_FUN(real) +__DEFINE_HIP_COMPLEX_FUN(imag) + #endif From 1eb91f1bab0007836c1fe663956442d02be23c29 Mon Sep 17 00:00:00 2001 From: pradeepisro Date: Tue, 19 Jun 2018 12:26:18 +0530 Subject: [PATCH 09/20] added changes to FindHIP.cmake which would allow us to build using ninja #467 [ROCm/hip commit: 4f0aefbd6320881804feadf38c2c00f09c021c95] --- projects/hip/cmake/FindHIP.cmake | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/projects/hip/cmake/FindHIP.cmake b/projects/hip/cmake/FindHIP.cmake index a45fb1b44d..e540ac35a6 100644 --- a/projects/hip/cmake/FindHIP.cmake +++ b/projects/hip/cmake/FindHIP.cmake @@ -471,7 +471,13 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files INPUT "${custom_target_script_pregen}" ) set(main_dep DEPENDS ${source_file}) - set(verbose_output "$(VERBOSE)") + if(CMAKE_GENERATOR MATCHES "Makefiles") + set(verbose_output "$(VERBOSE)") + elseif(HIP_VERBOSE_BUILD) + set(verbose_output ON) + else() + set(verbose_output OFF) + endif() # Create up the comment string file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}") From 4c9fad371f0a61480f46cd5f4bb01715019054b3 Mon Sep 17 00:00:00 2001 From: pradeepisro Date: Tue, 19 Jun 2018 13:50:24 +0530 Subject: [PATCH 10/20] added option HIP_BUILD_VERBOSE to enable verbosity in HIP build [ROCm/hip commit: f6938961b9011dd42a393fed34a868fd9bf5260f] --- projects/hip/cmake/FindHIP.cmake | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip/cmake/FindHIP.cmake b/projects/hip/cmake/FindHIP.cmake index e540ac35a6..b761356c9a 100644 --- a/projects/hip/cmake/FindHIP.cmake +++ b/projects/hip/cmake/FindHIP.cmake @@ -20,6 +20,7 @@ foreach(config ${_hip_configuration_types}) mark_as_advanced(HIP_HIPCC_FLAGS_${config_upper} HIP_HCC_FLAGS_${config_upper} HIP_NVCC_FLAGS_${config_upper}) endforeach() option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON) +option(HIP_VERBOSE_BUILD "Print out the commands run while compiling the HIP source file. With the Makefile generator this defaults to VERBOSE variable specified on the command line, but can be forced on with this option." OFF) mark_as_advanced(HIP_HOST_COMPILATION_CPP) ############################################################################### From b8ae2784e955b798f241a5b01c52a24d3044a253 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 14:35:48 -0400 Subject: [PATCH 11/20] Add conj, operator-,==,!= for hipFloatComplex/hipDoubleComplex [ROCm/hip commit: 84da72dae8b204cd4c89e78b2d99219e33011290] --- .../hip/include/hip/hcc_detail/hip_complex.h | 44 ++++++++++++++++--- 1 file changed, 38 insertions(+), 6 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 7f4dd840b7..48a2852a5a 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -35,6 +35,24 @@ THE SOFTWARE. #endif #if __cplusplus +#define COMPLEX_NEG_OP_OVERLOAD(type) \ + __device__ __host__ static inline type operator-(const type& op) { \ + type ret; \ + ret.x = -op.x; \ + ret.y = -op.y; \ + return ret; \ + } + +#define COMPLEX_EQ_OP_OVERLOAD(type) \ + __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \ + return lhs.x == rhs.x && lhs.y == rhs.y; \ + } + +#define COMPLEX_NE_OP_OVERLOAD(type) \ + __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \ + return !(lhs == rhs); \ + } + #define COMPLEX_ADD_OP_OVERLOAD(type) \ __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \ type ret; \ @@ -150,6 +168,9 @@ struct hipDoubleComplex { #if __cplusplus +COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex) +COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex) +COMPLEX_NE_OP_OVERLOAD(hipFloatComplex) COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex) COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex) COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex) @@ -169,6 +190,9 @@ COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double) COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long) COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long) +COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex) +COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex) +COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex) COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex) COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex) COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex) @@ -313,12 +337,20 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } -#define __DEFINE_HIP_COMPLEX_FUN(func) \ -__device__ __host__ inline float func(const hipFloatComplex& z) { return hipC##func##f(z); } \ -__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipC##func(z); } +// Complex functions returning real numbers. +#define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \ +__device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \ +__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); } -__DEFINE_HIP_COMPLEX_FUN(abs) -__DEFINE_HIP_COMPLEX_FUN(real) -__DEFINE_HIP_COMPLEX_FUN(imag) +__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs) +__DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal) +__DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag) + +// Complex functions returning complex numbers. +#define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \ +__device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \ +__device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); } + +__DEFINE_HIP_COMPLEX_FUN(conj, hipConj) #endif From 8696038043a5d32a143dd884d0b8ed38f54add13 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 18 Jun 2018 21:43:24 -0400 Subject: [PATCH 12/20] Let hipcc handle library with extension lo for hip-clang [ROCm/hip commit: 46d3c1d51e26f6e464e32872504d4219c3c5dec8] --- projects/hip/bin/hipcc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 4c526bd2c7..f5fe4249ca 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -376,7 +376,7 @@ foreach $arg (@ARGV) open my $out, ">", $new_file or die "$new_file: $!"; while (my $line = <$in>) { chomp $line; - if ($line =~ m/\.a$/) { + if ($line =~ m/\.a$/ || $line =~ m/\.lo$/) { my $libFile = $line; my $path = abs_path($line); my @objs = split ('\n', `cd $tmpdir; ar xv $path`); @@ -401,8 +401,8 @@ foreach $arg (@ARGV) if ($allIsObj) { print $out "$line\n"; } elsif ($realObjs) { - my $libBaseName = basename($libFile, ".a"); - $libBaseName = mktemp($libBaseName . "XXXX") . ".a"; + my($libBaseName, $libDir, $libExt) = fileparse($libFile); + $libBaseName = mktemp($libBaseName . "XXXX") . $libExt; system("cd $tmpdir; ar c $libBaseName $realObjs"); print $out "$tmpdir/$libBaseName\n"; } From a7baff62e7f97ee3d86113cb9056457f5c005405 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 21 Jun 2018 18:12:55 +0000 Subject: [PATCH 13/20] HIPCC - Use clang if clang++ is not compiled [ROCm/hip commit: 73023e0a54127e00d986508250fe05b1a5bbbd82] --- projects/hip/bin/hipcc | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 03f35b27fc..1be33d55e1 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -115,8 +115,13 @@ if ($HIP_PLATFORM eq "clang") { $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm"; $HIPCC="$HIP_CLANG_PATH/clang++"; $HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include"; - $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc"; + + # If $HIPCC clang++ is not compiled, use clang instead + if ( ! -e $HIPCC ) { + $HIPCC="$HIP_CLANG_PATH/clang"; + $HIPLDFLAGS = "--driver-mode=g++ " . $HIPLDFLAGS; + } } elsif ($HIP_PLATFORM eq "hcc") { $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa"; From 260e50769bc1d68260c74d51c513b5681fc4d9ae Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 22 Jun 2018 14:37:19 +0000 Subject: [PATCH 14/20] Add HIP Compute Mode [ROCm/hip commit: ff924ecb3d2a0135b03f653e266f63b01e9f22a4] --- projects/hip/include/hip/hip_runtime_api.h | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index e18c17e07b..2ff562cc01 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -293,6 +293,12 @@ typedef enum hipDeviceAttribute_t { hipDeviceAttributeIntegrated, ///< iGPU } hipDeviceAttribute_t; +enum hipComputeMode { + hipComputeModeDefault = 0, + hipComputeModeExclusive = 1, + hipComputeModeProhibited = 2, + hipComputeModeExcusiveProcess = 3 +}; /** * @} From 9fed6fed2bd929f724b9ecff9a3c9a1895fdb614 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 25 Jun 2018 19:16:27 +0100 Subject: [PATCH 15/20] Removes use of unimplemented OCML functionality. [ROCm/hip commit: 6c7a64efa2e4169a6425009f0d7354e23cb2b486] --- .../include/hip/hcc_detail/math_functions.h | 18 ++++++------- .../hip/include/hip/hcc_detail/math_fwd.h | 26 ------------------- 2 files changed, 9 insertions(+), 35 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index a9288dc204..8f7c885fc7 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -508,19 +508,19 @@ inline float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); } __device__ inline -float __fdiv_rd(float x, float y) { return __ocml_div_rtp_f32(x, y); } +float __fdiv_rd(float x, float y) { return x / y; } __device__ inline -float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); } +float __fdiv_rn(float x, float y) { return x / y; } __device__ inline -float __fdiv_ru(float x, float y) { return __ocml_div_rtn_f32(x, y); } +float __fdiv_ru(float x, float y) { return x / y; } __device__ inline -float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); } +float __fdiv_rz(float x, float y) { return x / y; } __device__ inline -float __fdividef(float x, float y) { return __ocml_div_rte_f32(x, y); } +float __fdividef(float x, float y) { return x / y; } __device__ inline float __fmaf_rd(float x, float y, float z) @@ -1028,16 +1028,16 @@ inline double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); } __device__ inline -double __ddiv_rd(double x, double y) { return __ocml_div_rtp_f64(x, y); } +double __ddiv_rd(double x, double y) { return x / y; } __device__ inline -double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); } +double __ddiv_rn(double x, double y) { return x / y; } __device__ inline -double __ddiv_ru(double x, double y) { return __ocml_div_rtn_f64(x, y); } +double __ddiv_ru(double x, double y) { return x / y; } __device__ inline -double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); } +double __ddiv_rz(double x, double y) { return x / y; } __device__ inline double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); } diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h index ac5d6e7dc6..102714e117 100644 --- a/projects/hip/include/hip/hcc_detail/math_fwd.h +++ b/projects/hip/include/hip/hcc_detail/math_fwd.h @@ -69,8 +69,6 @@ float __ocml_cosh_f32(float); __device__ float __ocml_cospi_f32(float); __device__ -float __ocml_div_rtz_f32(float, float); -__device__ float __ocml_i0_f32(float); __device__ float __ocml_i1_f32(float); @@ -290,18 +288,6 @@ __attribute__((const)) float __ocml_mul_rtz_f32(float, float); __device__ __attribute__((const)) -float __ocml_div_rte_f32(float, float); -__device__ -__attribute__((const)) -float __ocml_div_rtn_f32(float, float); -__device__ -__attribute__((const)) -float __ocml_div_rtp_f32(float, float); -__device__ -__attribute__((const)) -float __ocml_div_rtz_f32(float, float); -__device__ -__attribute__((const)) float __ocml_sqrt_rte_f32(float); __device__ __attribute__((const)) @@ -598,18 +584,6 @@ __attribute__((const)) double __ocml_mul_rtz_f64(double, double); __device__ __attribute__((const)) -double __ocml_div_rte_f64(double, double); -__device__ -__attribute__((const)) -double __ocml_div_rtn_f64(double, double); -__device__ -__attribute__((const)) -double __ocml_div_rtp_f64(double, double); -__device__ -__attribute__((const)) -double __ocml_div_rtz_f64(double, double); -__device__ -__attribute__((const)) double __ocml_sqrt_rte_f64(double); __device__ __attribute__((const)) From c205c090f4beb0a421329942731578cb2b08f564 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 25 Jun 2018 15:33:41 -0400 Subject: [PATCH 16/20] Include host_defines.h in hip_fp16.h since it uses __host__ __device__ attributes [ROCm/hip commit: 02b160491d3526dfbe1036880bf5a0356cbe62a5] --- projects/hip/include/hip/hcc_detail/hip_fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 8657bc30a3..777113a9de 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -21,7 +21,7 @@ THE SOFTWARE. */ #pragma once - +#include "hip/hcc_detail/host_defines.h" #include #if defined(__cplusplus) #include From af0e9b54de8373fc0c0a499bf3c300850f48b4ed Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 25 Jun 2018 18:12:36 -0400 Subject: [PATCH 17/20] Add workaround to hipcc for build failure in tensorflow due to missing symbol __cpu_model https://github.com/tensorflow/tensorflow/issues/9593 [ROCm/hip commit: 021728cd863f212fa5ebd18cba31af1af9a3f789] --- projects/hip/bin/hipcc | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 03f35b27fc..731c9bcc96 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -601,6 +601,9 @@ if ($needLDFLAGS and not $compileOnly) { $CMD .= " $HIPLDFLAGS"; } $CMD .= " $toolArgs"; +if ($needLDFLAGS and not $compileOnly and $HIP_PLATFORM eq "clang") { + $CMD .= " -lgcc_s -lgcc"; +} if ($verbose & 0x1) { print "hipcc-cmd: ", $CMD, "\n"; From fe5afd44f45d9f2a6df8f728f5e2fee70705a178 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 28 Jun 2018 11:19:22 +0530 Subject: [PATCH 18/20] Fix typo [ROCm/hip commit: dde875f23ba37bc4247942654e69d576fbd89b29] --- projects/hip/include/hip/hcc_detail/hip_runtime.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 7c1f2a4d16..ca4e73eee0 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -109,7 +109,7 @@ extern int HIP_TRACE_API; #include #include #include -#if __HCC___ +#if __HCC__ #include #include #endif // __HCC__ From b7367558ef58563868c58ed51504e8cebb492136 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 2 Jul 2018 10:37:20 +0530 Subject: [PATCH 19/20] Update hip-targets.cmake for hip::host & hip::device [ROCm/hip commit: 4b660c8382a99cddfc21b52e084cbe82919305da] --- projects/hip/packaging/hip-targets.cmake | 20 +++++++++++++++++++- 1 file changed, 19 insertions(+), 1 deletion(-) diff --git a/projects/hip/packaging/hip-targets.cmake b/projects/hip/packaging/hip-targets.cmake index 65370eec9e..5aff2ca1ed 100644 --- a/projects/hip/packaging/hip-targets.cmake +++ b/projects/hip/packaging/hip-targets.cmake @@ -16,7 +16,7 @@ set(CMAKE_IMPORT_FILE_VERSION 1) set(_targetsDefined) set(_targetsNotDefined) set(_expectedTargets) -foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device) +foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device hip::host hip::device) list(APPEND _expectedTargets ${_expectedTarget}) if(NOT TARGET ${_expectedTarget}) list(APPEND _targetsNotDefined ${_expectedTarget}) @@ -65,6 +65,24 @@ set_target_properties(hip::hip_device PROPERTIES INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include" ) +# Create imported target hip::host +add_library(hip::host INTERFACE IMPORTED) + +set_target_properties(hip::host PROPERTIES + INTERFACE_LINK_LIBRARIES "hip::hip_hcc" +) + +# Create imported target hip::device +add_library(hip::device INTERFACE IMPORTED) + +set_target_properties(hip::device PROPERTIES + INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am" +) + +if(CMAKE_VERSION VERSION_LESS 3.0.0) + message(FATAL_ERROR "This file relies on consumers using CMake 3.0.0 or greater.") +endif() + # Load information for each installed configuration. get_filename_component(_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH) file(GLOB CONFIG_FILES "${_DIR}/hip-targets-*.cmake") From 19cb14d54cd10975c8e9ec4a6b678bf67dc1fd97 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 4 Jul 2018 09:33:51 +0530 Subject: [PATCH 20/20] Update hip_hcc_internal.h Adding missing include for hip_hcc_internal in order to build with HCC [ROCm/hip commit: 0c2f985553d6223d1219a3868ecb8d474cd6de20] --- projects/hip/src/hip_hcc_internal.h | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 4008df1574..be257aff4f 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -26,6 +26,7 @@ THE SOFTWARE. #include #include #include +#include #include "hsa/hsa_ext_amd.h" #include "hip/hip_runtime.h"