Remove HCC (#1546)

* deprecate HCC

.clang-tidy

 CheckOptions: 
   - key:             bugprone-reserved-identifier.AllowedIdentifiers
-    value:           '__HIP_PLATFORM_HCC__;__HIP_ROCclr__'
+    value:           '__HIP_PLATFORM_AMD__;__HIP_ROCclr__'
   - key:             bugprone-unused-return-value.CheckedFunctions
     value:           '::std::async;::std::launder;::std::remove;::std::remove_if;::std::unique;::std::unique_ptr::release;::std::basic_string::empty;::std::vector::empty;::std::find;::std::find_if;::std::find_if_not;::std::all_of;::std::any_of;::std::none_of;::std::count;::std::count_if;::std::mismatch;::std::find_end;::std::find_first_of;::std::adjacent_find;::std::search;::std::search_n;::std::nth_element;::std::lower_bound;::std::upper_bound;::std::binary_search;::std::equal_range;::std::max;::std::max_element;::std::min;::std::min_element;::std::minmax;::std::minmax_element;::std::equal;::std::lexicographical_compare;::std::accumulate;::std::inner_product'
   - key:             cppcoreguidelines-macro-usage.AllowedRegexp

CMakeLists.txt

@@ -95,7 +95,7 @@ endif()
 list(APPEND CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake)
 include(EnableCompilerWarnings)
 include(ROCMClangTidy)
-if(CMAKE_CXX_COMPILER MATCHES ".*hcc" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+")
+if(CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+")
     set(MIGRAPHX_TIDY_ERRORS ERRORS * -readability-inconsistent-declaration-parameter-name)
 # Enable tidy on hip
 elseif(MIGRAPHX_ENABLE_GPU)

Jenkinsfile

@@ -33,11 +33,9 @@ def rocmtestnode(Map conf) {
         """
         echo cmd
         sh cmd
-        if (compiler != "hcc") {
-            // Only archive from master or develop
-            if (env.BRANCH_NAME == "develop" || env.BRANCH_NAME == "master") {
-                archiveArtifacts artifacts: "build/*.deb", allowEmptyArchive: true, fingerprint: true
-            }
+        // Only archive from master or develop
+        if (env.BRANCH_NAME == "develop" || env.BRANCH_NAME == "master") {
+            archiveArtifacts artifacts: "build/*.deb", allowEmptyArchive: true, fingerprint: true
         }
     }
     node(name) {

install_deps.cmake

 #!/usr/bin/cmake -P
 
-#####################################################################################
+# ####################################################################################
 # The MIT License (MIT)
 #
 # Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
@@ -22,9 +22,10 @@
 # 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.
-#####################################################################################
+# ####################################################################################
 
 set(ARGS)
+
 foreach(i RANGE 3 ${CMAKE_ARGC})
     list(APPEND ARGS ${CMAKE_ARGV${i}})
 endforeach()
@@ -38,18 +39,19 @@ set(multiValueArgs)
 cmake_parse_arguments(PARSE "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGS})
 
 if(PARSE_help)
-message("Usage: install_deps.cmake [options] [cmake-args]")
-message("")
-message("Options:")
-message("  --prefix               Set the prefix to install the dependencies.")
-message("")
-message("Commands:")
-message("  help                   Show this message and exit.")
-message("")
-return()
+    message("Usage: install_deps.cmake [options] [cmake-args]")
+    message("")
+    message("Options:")
+    message("  --prefix               Set the prefix to install the dependencies.")
+    message("")
+    message("Commands:")
+    message("  help                   Show this message and exit.")
+    message("")
+    return()
 endif()
 
 set(_PREFIX /usr/local)
+
 if(PARSE_--prefix)
     set(_PREFIX ${PARSE_--prefix})
 endif()
@@ -63,21 +65,23 @@ if(NOT CMakeGet_FOUND)
     file(DOWNLOAD https://raw.githubusercontent.com/pfultz2/cmake-get/master/install.cmake ${FILENAME} STATUS RESULT_LIST)
     list(GET RESULT_LIST 0 RESULT)
     list(GET RESULT_LIST 1 RESULT_MESSAGE)
+
     if(NOT RESULT EQUAL 0)
         message(FATAL_ERROR "Download for install.cmake failed: ${RESULT_MESSAGE}")
     endif()
+
     execute_process(COMMAND ${CMAKE_COMMAND} -P ${FILENAME} ${PREFIX})
     file(REMOVE ${FILENAME})
     find_package(CMakeGet REQUIRED PATHS ${PREFIX})
 endif()
 
-# Set compiler to hcc if not set
+# Set compiler to clang++ if not set
 if(NOT DEFINED ENV{CXX} AND NOT DEFINED CMAKE_CXX_COMPILER AND NOT DEFINED CMAKE_TOOLCHAIN_FILE)
-    find_program(HCC hcc PATHS /opt/rocm PATH_SUFFIXES bin)
-    if(HCC)
-        set(ENV{CXX} ${HCC})
+    find_program(CLANG clang++ PATHS /opt/rocm /opt/rocm/llvm PATH_SUFFIXES bin)
+    if(CLANG)
+        set(ENV{CXX} ${CLANG})
     else()
-        message(FATAL_ERROR "Cannot find hcc")
+        message(FATAL_ERROR "Cannot find clang++")
     endif()
 endif()
 

src/compile_src.cpp

@@ -70,9 +70,6 @@ std::vector<char> src_compiler::compile(const std::vector<src_file>& srcs) const
     if(not fs::exists(out_path))
         MIGRAPHX_THROW("Output file missing: " + out);
 
-    if(process)
-        out_path = process(out_path);
-
     return read_buffer(out_path.string());
 }
 

src/targets/gpu/CMakeLists.txt

@@ -22,7 +22,7 @@
 # THE SOFTWARE.
 # ####################################################################################
 
-list(APPEND CMAKE_PREFIX_PATH /opt/rocm /opt/rocm/hip /opt/rocm/hcc)
+list(APPEND CMAKE_PREFIX_PATH /opt/rocm /opt/rocm/hip)
 find_package(miopen)
 
 # rocblas
@@ -170,21 +170,6 @@ register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp
 rocm_set_soversion(migraphx_gpu ${MIGRAPHX_SO_VERSION})
 rocm_clang_tidy_check(migraphx_gpu)
 
-get_filename_component(CMAKE_CXX_COMPILER_PATH "${CMAKE_CXX_COMPILER}" PATH)
-
-if(NOT CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+$")
-    find_program(MIGRAPHX_EXTRACT_KERNEL extractkernel
-        PATH_SUFFIXES bin
-        HINTS ${CMAKE_CXX_COMPILER_PATH}
-        PATHS
-        /opt/rocm/hip
-        /opt/rocm/hcc
-        /opt/rocm
-    )
-endif()
-
-message(STATUS "extractkernel: ${MIGRAPHX_EXTRACT_KERNEL}")
-
 set(MIGRAPHX_ENABLE_MLIR OFF CACHE BOOL "")
 
 if(MIGRAPHX_ENABLE_MLIR)
@@ -220,7 +205,6 @@ else()
     target_compile_definitions(migraphx_gpu PRIVATE
         "-DMIGRAPHX_HIP_COMPILER=${CMAKE_CXX_COMPILER}"
         "-DMIGRAPHX_HIP_COMPILER_FLAGS=${HIP_COMPILER_FLAGS}"
-        "-DMIGRAPHX_EXTRACT_KERNEL=${MIGRAPHX_EXTRACT_KERNEL}"
     )
 
     if(DEFINED CMAKE_CXX_COMPILER_LAUNCHER)
@@ -252,8 +236,6 @@ else()
     message(STATUS "MIOpen does not have find mode api")
 endif()
 
-# Workaround broken rocblas headers
-target_compile_definitions(migraphx_gpu PUBLIC -D__HIP_PLATFORM_HCC__=1)
 target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
 target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
 

src/targets/gpu/compile_hip.cpp

@@ -207,12 +207,6 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
 
 #else // MIGRAPHX_USE_HIPRTC
 
-bool is_hcc_compiler()
-{
-    static const auto result = ends_with(MIGRAPHX_STRINGIZE(MIGRAPHX_HIP_COMPILER), "hcc");
-    return result;
-}
-
 bool is_hip_clang_compiler()
 {
     static const auto result = ends_with(MIGRAPHX_STRINGIZE(MIGRAPHX_HIP_COMPILER), "clang++");
@@ -236,7 +230,7 @@ std::vector<std::vector<char>>
 compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std::string& arch)
 {
     assert(not srcs.empty());
-    if(not is_hcc_compiler() and not is_hip_clang_compiler())
+    if(not is_hip_clang_compiler())
         MIGRAPHX_THROW("Unknown hip compiler: " +
                        std::string(MIGRAPHX_STRINGIZE(MIGRAPHX_HIP_COMPILER)));
 
@@ -246,16 +240,9 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
     if(enabled(MIGRAPHX_GPU_DEBUG_SYM{}))
         params += " -g";
     params += " -c";
-    if(is_hcc_compiler())
-    {
-        params += " -amdgpu-target=" + arch;
-    }
-    else if(is_hip_clang_compiler())
-    {
-        params += " --offload-arch=" + arch;
-        params += " --cuda-device-only";
-        params += " -O" + string_value_of(MIGRAPHX_GPU_OPTIMIZE{}, "3") + " ";
-    }
+    params += " --offload-arch=" + arch;
+    params += " --cuda-device-only";
+    params += " -O" + string_value_of(MIGRAPHX_GPU_OPTIMIZE{}, "3") + " ";
 
     if(enabled(MIGRAPHX_GPU_DEBUG{}))
         params += " -DMIGRAPHX_DEBUG";
@@ -270,24 +257,6 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
     if(has_compiler_launcher())
         compiler.launcher = MIGRAPHX_STRINGIZE(MIGRAPHX_HIP_COMPILER_LAUNCHER);
 #endif
-
-    if(is_hcc_compiler())
-        compiler.process = [&](const fs::path& obj_path) -> fs::path {
-            process{MIGRAPHX_STRINGIZE(MIGRAPHX_EXTRACT_KERNEL) + std::string{" -i "} +
-                    obj_path.string()}
-                .cwd(obj_path.parent_path());
-            for(const auto& entry : fs::directory_iterator{obj_path.parent_path()})
-            {
-                const auto& hsaco_path = entry.path();
-                if(not fs::is_regular_file(hsaco_path))
-                    continue;
-                if(hsaco_path.extension() != ".hsaco")
-                    continue;
-                return hsaco_path;
-            }
-            MIGRAPHX_THROW("Missing hsaco");
-        };
-
     if(enabled(MIGRAPHX_GPU_DUMP_SRC{}))
     {
         for(const auto& src : srcs)

src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp

@@ -112,14 +112,8 @@ inline auto gs_launch(hipStream_t stream, index_int n, index_int local = 1024)
 #ifdef MIGRAPHX_USE_CLANG_TIDY
 #define MIGRAPHX_DEVICE_SHARED
 #else
-// Workaround hcc's broken tile_static macro
-#ifdef tile_static
-#undef tile_static
-#define MIGRAPHX_DEVICE_SHARED __attribute__((tile_static))
-#else
 #define MIGRAPHX_DEVICE_SHARED __shared__
 #endif
-#endif
 
 } // namespace device
 } // namespace gpu

src/targets/gpu/kernels/include/migraphx/kernels/hip.hpp

@@ -25,10 +25,6 @@
 #define MIGRAPHX_GUARD_KERNELS_HIP_HPP
 
 #ifndef MIGRAPHX_USE_HIPRTC
-// Workaround macro redefinition issue with clang tidy
-#if defined(__HIP_PLATFORM_HCC__) && defined(MIGRAPHX_USE_CLANG_TIDY)
-#undef __HIP_PLATFORM_HCC__ // NOLINT
-#endif
 #include <hip/hip_runtime.h>
 #include <hip/hip_fp16.h>
 #include <hip/math_functions.h>

tools/build_and_test_onnxrt.sh

@@ -25,7 +25,7 @@ cd /onnxruntime
 pip3 install -r requirements-dev.txt
 # Add newer cmake to the path
 export PATH="/opt/cmake/bin:$PATH"
-export CXXFLAGS="-D__HIP_PLATFORM_HCC__=1 -w"
+export CXXFLAGS="-D__HIP_PLATFORM_AMD__=1 -w"
 ./build.sh --config Release  --cmake_extra_defines CMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ --update --build --parallel --cmake_extra_defines ONNXRUNTIME_VERSION=$(cat ./VERSION_NUMBER) --skip_tests --rocm_home /opt/rocm --use_migraphx --migraphx_home /opt/rocm --rocm_version=`cat /opt/rocm/.info/version-dev`
 
 cd build/Linux/Release