From c05e2e2261d606113e3990f7ca59eeea716a03c5 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 18 Apr 2018 15:27:56 -0400 Subject: [PATCH 1/2] initial gfx906 support [ROCm/hip commit: b898049412a1e101358118d75c04c110fdd340f4] --- projects/hip/CMakeLists.txt | 2 +- projects/hip/bin/hipcc | 21 +++++++++++++++++++ projects/hip/lpl_ca/ca.hpp | 4 ++-- projects/hip/lpl_ca/common.hpp | 4 ++-- projects/hip/lpl_ca/lpl.hpp | 4 ++-- .../hip/tests/src/deviceLib/hipTestHalf.cpp | 2 +- 6 files changed, 29 insertions(+), 8 deletions(-) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 075b916a4d..4f770b3429 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -201,7 +201,7 @@ if(HIP_PLATFORM STREQUAL "hcc") execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") - set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900") + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906") if(COMPILE_HIP_ATP_MARKER) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger") endif() diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index f796b0bf95..b574ed29bd 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -75,6 +75,7 @@ $target_gfx801 = 0; $target_gfx802 = 0; $target_gfx803 = 0; $target_gfx900 = 0; +$target_gfx906 = 0; $default_amdgpu_target = 1; if ($HIP_PLATFORM eq "hcc") { @@ -281,6 +282,12 @@ foreach $arg (@ARGV) $target_gfx900 = 1; $default_amdgpu_target = 0; } + if($arg eq '--amdgpu-target=gfx906') + { + $target_gfx906 = 1; + $default_amdgpu_target = 0; + } + if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0)) { @@ -373,6 +380,11 @@ if($HIP_PLATFORM eq "hcc"){ $target_gfx900 = 1; $default_amdgpu_target = 0; } + if($target eq 'gfx906') + { + $target_gfx906 = 1; + $default_amdgpu_target = 0; + } } } # Else try using rocm_agent_enumerator @@ -404,6 +416,10 @@ if($HIP_PLATFORM eq "hcc"){ $target_gfx900 = 1; $default_amdgpu_target = 0; } + if($val eq "gfx906") { + $target_gfx906 = 1; + $default_amdgpu_target = 0; + } } } # rocm_agent_enumerator failed! Throw an error and die if linking is required @@ -437,6 +453,11 @@ if($HIP_PLATFORM eq "hcc"){ $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX900__=1 "; $ENV{HCC_EXTRA_LIBRARIES_GFX900}="$HIP_PATH/lib/hip_hc_gfx803.ll\n"; } + if ($target_gfx906 eq 1) { + $HIPLDFLAGS .= " --amdgpu-target=gfx906"; + $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX906__=1 "; + $ENV{HCC_EXTRA_LIBRARIES_GFX906}="$HIP_PATH/lib/hip_hc_gfx803.ll\n"; + } } if ($hasC and $HIP_PLATFORM eq 'nvcc') { diff --git a/projects/hip/lpl_ca/ca.hpp b/projects/hip/lpl_ca/ca.hpp index bb1963bede..0ef8458c20 100644 --- a/projects/hip/lpl_ca/ca.hpp +++ b/projects/hip/lpl_ca/ca.hpp @@ -23,7 +23,7 @@ inline clara::Parser cmdline_parser(bool& help, std::vector& inputs "https://reviews.llvm.org/D13909; " "the code object format is documented at: " "https://www.llvm.org/docs/AMDGPUUsage.html#code-object.") | - clara::Opt{targets, "gfx803,gfx900 etc."}["-t"]["--targets"]( + clara::Opt{targets, "gfx803,gfx900,gfx906 etc."}["-t"]["--targets"]( "targets for which code objects are to be extracted from " "the fat binary; must be included in the set of processors " "with ROCm support from " @@ -76,4 +76,4 @@ inline void validate_inputs(const std::vector& inputs) { throw std::runtime_error{"Non existent file " + *it + " passed as input."}; } } -} // namespace hip_impl \ No newline at end of file +} // namespace hip_impl diff --git a/projects/hip/lpl_ca/common.hpp b/projects/hip/lpl_ca/common.hpp index 74f4db9221..7652f08c46 100644 --- a/projects/hip/lpl_ca/common.hpp +++ b/projects/hip/lpl_ca/common.hpp @@ -12,7 +12,7 @@ namespace hip_impl { inline const std::unordered_set& amdgpu_targets() { // The evolving list lives at: // https://www.llvm.org/docs/AMDGPUUsage.html#processors. static const std::unordered_set r{"gfx701", "gfx801", "gfx802", "gfx803", - "gfx900"}; + "gfx900", "gfx906"}; return r; } @@ -77,4 +77,4 @@ inline void validate_targets(const std::vector& x) { } } } -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. diff --git a/projects/hip/lpl_ca/lpl.hpp b/projects/hip/lpl_ca/lpl.hpp index c9346fd793..cbd7fe8386 100644 --- a/projects/hip/lpl_ca/lpl.hpp +++ b/projects/hip/lpl_ca/lpl.hpp @@ -132,9 +132,9 @@ inline clara::Parser cmdline_parser(bool& help, std::vector& source "file is documented at: https://reviews.llvm.org/D13909.") | clara::Arg{sources, "a.cpp b.cpp etc."}("inputs for compilation; must contain valid C++ code.") | - clara::Opt{targets, "gfx803,gfx900 etc."}["-t"]["--targets"]( + clara::Opt{targets, "gfx803,gfx900,gfx906 etc."}["-t"]["--targets"]( "targets for AMDGPU lowering; must be included in the set " "of processors with ROCm support from " "https://www.llvm.org/docs/AMDGPUUsage.html#processors."); } -} // namespace hip_impl \ No newline at end of file +} // namespace hip_impl diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp index 4a63260bf7..24a4d6c53e 100644 --- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #define HALF_SIZE 64 * sizeof(__half) #define HALF2_SIZE 64 * sizeof(__half2) -#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ +#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ __global__ void __halfMath(hipLaunchParm lp, __half* A, __half* B, __half* C) { int tx = threadIdx.x; From 9a591d6c2695d18b1e51af9627b2a7031408bb49 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 11 May 2018 10:53:07 -0400 Subject: [PATCH 2/2] Fix hipMathFunction for gfx906 [ROCm/hip commit: 848a24b5246f662f5b488fe428c8b414ec708dda] --- projects/hip/tests/src/deviceLib/hipMathFunctions.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp b/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp index 7fe0003672..78e85ba62a 100644 --- a/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp +++ b/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp @@ -29,7 +29,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ +#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ __global__ void kernel_abs_int64(hipLaunchParm lp, long long *input, long long *output) { int tx = threadIdx.x;