mirror of
https://github.com/saymrwulf/onnxruntime.git
synced 2026-07-30 20:18:08 +00:00
Hipify during build instead of before cmake config (#13333)
### Description Currently, hipify happens before cmake is configured and then cmake glob the directories. This get rids of thoes customized python threading logic and opt for build system itself to generate the files. This also supersede the half baked branch [sukha/hipify-with-cmake](https://github.com/microsoft/onnxruntime/tree/sukha/hipify-with-cmake)
This commit is contained in:
cloudhan
GitHub
parent
bb16ee712e
commit
928c9fc348
5 changed files with 253 additions and 321 deletions
|
|
@ -1871,6 +1871,10 @@ if (onnxruntime_USE_ROCM)
|
|||
set(CMAKE_HIP_ARCHITECTURES "gfx906;gfx908;gfx90a;gfx1030")
|
||||
endif()
|
||||
|
||||
if (NOT onnxruntime_HIPIFY_PERL)
|
||||
set(onnxruntime_HIPIFY_PERL ${onnxruntime_ROCM_HOME}/hip/bin/hipify-perl)
|
||||
endif()
|
||||
|
||||
# NOTE: HIP language is added in 3.21 and there are bugs before 3.23.1
|
||||
cmake_minimum_required(VERSION 3.23.1 FATAL_ERROR)
|
||||
find_package(HIP)
|
||||
|
|
|
|||
|
|
@ -93,16 +93,6 @@ file(GLOB_RECURSE onnxruntime_rocm_contrib_ops_cu_srcs CONFIGURE_DEPENDS
|
|||
"${ONNXRUNTIME_ROOT}/contrib_ops/rocm/*.cuh"
|
||||
)
|
||||
|
||||
file(GLOB_RECURSE onnxruntime_rocm_generated_contrib_ops_cc_srcs CONFIGURE_DEPENDS
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/contrib_ops/rocm/*.h"
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/contrib_ops/rocm/*.cc"
|
||||
)
|
||||
|
||||
file(GLOB_RECURSE onnxruntime_rocm_generated_contrib_ops_cu_srcs CONFIGURE_DEPENDS
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/contrib_ops/rocm/*.cu"
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/contrib_ops/rocm/*.cuh"
|
||||
)
|
||||
|
||||
file(GLOB onnxruntime_providers_common_srcs CONFIGURE_DEPENDS
|
||||
"${ONNXRUNTIME_ROOT}/core/providers/*.h"
|
||||
"${ONNXRUNTIME_ROOT}/core/providers/*.cc"
|
||||
|
|
@ -1276,6 +1266,7 @@ endif()
|
|||
|
||||
if (onnxruntime_USE_ROCM)
|
||||
add_definitions(-DUSE_ROCM=1)
|
||||
include(onnxruntime_rocm_hipify.cmake)
|
||||
|
||||
# Add search paths for default hip installation
|
||||
list(APPEND CMAKE_PREFIX_PATH ${onnxruntime_ROCM_HOME} ${onnxruntime_ROCM_HOME}/hip ${onnxruntime_ROCM_HOME}/hcc ${onnxruntime_ROCM_HOME}/miopen ${onnxruntime_ROCM_HOME}/hiprand ${onnxruntime_ROCM_HOME}/rocrand)
|
||||
|
|
@ -1306,15 +1297,7 @@ if (onnxruntime_USE_ROCM)
|
|||
"${ONNXRUNTIME_ROOT}/core/providers/rocm/*.cuh"
|
||||
)
|
||||
|
||||
file(GLOB_RECURSE onnxruntime_providers_rocm_generated_cc_srcs CONFIGURE_DEPENDS
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/core/providers/rocm/*.h"
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/core/providers/rocm/*.cc"
|
||||
)
|
||||
|
||||
file(GLOB_RECURSE onnxruntime_providers_rocm_generated_cu_srcs CONFIGURE_DEPENDS
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/core/providers/rocm/*.cu"
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/onnxruntime/core/providers/rocm/*.cuh"
|
||||
)
|
||||
hipify("onnxruntime/core/providers" provider_excluded_files onnxruntime_providers_rocm_generated_cc_srcs onnxruntime_providers_rocm_generated_cu_srcs)
|
||||
|
||||
source_group(TREE ${ONNXRUNTIME_ROOT}/core FILES ${onnxruntime_providers_rocm_cc_srcs} ${onnxruntime_providers_rocm_shared_srcs} ${onnxruntime_providers_rocm_cu_srcs})
|
||||
set(onnxruntime_providers_rocm_src ${onnxruntime_providers_rocm_cc_srcs} ${onnxruntime_providers_rocm_shared_srcs} ${onnxruntime_providers_rocm_cu_srcs})
|
||||
|
|
@ -1327,6 +1310,9 @@ if (onnxruntime_USE_ROCM)
|
|||
"${ONNXRUNTIME_ROOT}/contrib_ops/rocm/aten_ops/aten_op.cc"
|
||||
)
|
||||
endif()
|
||||
|
||||
hipify("onnxruntime/contrib_ops" contrib_ops_excluded_files onnxruntime_rocm_generated_contrib_ops_cc_srcs onnxruntime_rocm_generated_contrib_ops_cu_srcs)
|
||||
|
||||
# add using ONNXRUNTIME_ROOT so they show up under the 'contrib_ops' folder in Visual Studio
|
||||
source_group(TREE ${ONNXRUNTIME_ROOT} FILES ${onnxruntime_rocm_contrib_ops_cc_srcs} ${onnxruntime_rocm_contrib_ops_cu_srcs})
|
||||
list(APPEND onnxruntime_providers_rocm_src ${onnxruntime_rocm_contrib_ops_cc_srcs} ${onnxruntime_rocm_contrib_ops_cu_srcs})
|
||||
|
|
@ -1344,15 +1330,7 @@ if (onnxruntime_USE_ROCM)
|
|||
"${ORTTRAINING_SOURCE_DIR}/training_ops/rocm/*.cuh"
|
||||
)
|
||||
|
||||
file(GLOB_RECURSE onnxruntime_rocm_generated_training_ops_cc_srcs CONFIGURE_DEPENDS
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/orttraining/orttraining/training_ops/rocm/*.h"
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/orttraining/orttraining/training_ops/rocm/*.cc"
|
||||
)
|
||||
|
||||
file(GLOB_RECURSE onnxruntime_rocm_generated_training_ops_cu_srcs CONFIGURE_DEPENDS
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/orttraining/orttraining/training_ops/rocm/*.cu"
|
||||
"${CMAKE_CURRENT_BINARY_DIR}/amdgpu/orttraining/orttraining/training_ops/rocm/*.cuh"
|
||||
)
|
||||
hipify("orttraining/orttraining/training_ops" training_ops_excluded_files onnxruntime_rocm_generated_training_ops_cc_srcs onnxruntime_rocm_generated_training_ops_cu_srcs)
|
||||
|
||||
# NCCL is not support in Windows build
|
||||
if (WIN32 OR NOT onnxruntime_USE_NCCL)
|
||||
|
|
@ -1406,7 +1384,15 @@ if (onnxruntime_USE_ROCM)
|
|||
endif()
|
||||
endif()
|
||||
|
||||
add_dependencies(onnxruntime_providers_rocm onnxruntime_providers_shared ${onnxruntime_EXTERNAL_DEPENDENCIES})
|
||||
add_custom_target(generate_hipified_files DEPENDS
|
||||
${onnxruntime_providers_rocm_generated_cc_srcs}
|
||||
${onnxruntime_providers_rocm_generated_cu_srcs}
|
||||
${onnxruntime_rocm_generated_contrib_ops_cc_srcs}
|
||||
${onnxruntime_rocm_generated_contrib_ops_cu_srcs}
|
||||
${onnxruntime_rocm_generated_training_ops_cc_srcs}
|
||||
${onnxruntime_rocm_generated_training_ops_cu_srcs})
|
||||
|
||||
add_dependencies(onnxruntime_providers_rocm generate_hipified_files onnxruntime_providers_shared ${onnxruntime_EXTERNAL_DEPENDENCIES})
|
||||
target_link_libraries(onnxruntime_providers_rocm PRIVATE ${ONNXRUNTIME_ROCM_LIBS} ${ONNXRUNTIME_PROVIDERS_SHARED} ${ABSEIL_LIBS})
|
||||
target_include_directories(onnxruntime_providers_rocm SYSTEM
|
||||
PRIVATE
|
||||
|
|
|
|||
221
cmake/onnxruntime_rocm_hipify.cmake
Normal file
221
cmake/onnxruntime_rocm_hipify.cmake
Normal file
|
|
@ -0,0 +1,221 @@
|
|||
# Copyright (c) Microsoft Corporation. All rights reserved.
|
||||
# Licensed under the MIT License.
|
||||
|
||||
find_package(Python3 COMPONENTS Interpreter REQUIRED)
|
||||
|
||||
# GLOB pattern of file to be excluded
|
||||
set(contrib_ops_excluded_files
|
||||
"bert/attention.cc"
|
||||
"bert/attention.h"
|
||||
"bert/attention_impl.cu"
|
||||
"bert/attention_softmax.h"
|
||||
"bert/embed_layer_norm.cc"
|
||||
"bert/embed_layer_norm.h"
|
||||
"bert/embed_layer_norm_impl.cu"
|
||||
"bert/embed_layer_norm_impl.h"
|
||||
"bert/fast_gelu_impl.cu"
|
||||
"bert/fast_gelu_impl.h"
|
||||
"bert/fast_gelu.cc"
|
||||
"bert/fast_gelu.h"
|
||||
"bert/skip_layer_norm.cc"
|
||||
"bert/skip_layer_norm.h"
|
||||
"bert/skip_layer_norm_impl.cu"
|
||||
"bert/skip_layer_norm_impl.h"
|
||||
"bert/tensorrt_fused_multihead_attention/*"
|
||||
"bert/transformer_common.h"
|
||||
"bert/transformer_common.cc"
|
||||
"math/complex_mul.cc"
|
||||
"math/complex_mul.h"
|
||||
"math/complex_mul_impl.cu"
|
||||
"math/complex_mul_impl.h"
|
||||
"math/cufft_plan_cache.h"
|
||||
"math/fft_ops.cc"
|
||||
"math/fft_ops.h"
|
||||
"math/fft_ops_impl.cu"
|
||||
"math/fft_ops_impl.h"
|
||||
"quantization/attention_quantization.cc"
|
||||
"quantization/attention_quantization.h"
|
||||
"quantization/attention_quantization_impl.cu"
|
||||
"quantization/attention_quantization_impl.cuh"
|
||||
"quantization/quantize_dequantize_linear.cc"
|
||||
"quantization/qordered_ops/qordered_attention_impl.cu"
|
||||
"quantization/qordered_ops/qordered_attention_impl.h"
|
||||
"quantization/qordered_ops/qordered_attention_input_enum.h"
|
||||
"quantization/qordered_ops/qordered_attention.cc"
|
||||
"quantization/qordered_ops/qordered_attention.h"
|
||||
"quantization/qordered_ops/qordered_common.cuh"
|
||||
"quantization/qordered_ops/qordered_layer_norm.h"
|
||||
"quantization/qordered_ops/qordered_layer_norm.cc"
|
||||
"quantization/qordered_ops/qordered_layer_norm_impl.h"
|
||||
"quantization/qordered_ops/qordered_layer_norm_impl.cu"
|
||||
"quantization/qordered_ops/qordered_longformer_attention.cc"
|
||||
"quantization/qordered_ops/qordered_longformer_attention.h"
|
||||
"quantization/qordered_ops/qordered_matmul.h"
|
||||
"quantization/qordered_ops/qordered_matmul.cc"
|
||||
"quantization/qordered_ops/qordered_matmul_utils.h"
|
||||
"quantization/qordered_ops/qordered_matmul_utils.cc"
|
||||
"quantization/qordered_ops/qordered_qdq_impl.cu"
|
||||
"quantization/qordered_ops/qordered_qdq_impl.h"
|
||||
"quantization/qordered_ops/qordered_qdq.cc"
|
||||
"quantization/qordered_ops/qordered_qdq.h"
|
||||
"quantization/qordered_ops/qordered_unary_ops.h"
|
||||
"quantization/qordered_ops/qordered_unary_ops.cc"
|
||||
"quantization/qordered_ops/qordered_unary_ops_impl.h"
|
||||
"quantization/qordered_ops/qordered_unary_ops_impl.cu"
|
||||
"tensor/crop.cc"
|
||||
"tensor/crop.h"
|
||||
"tensor/crop_impl.cu"
|
||||
"tensor/crop_impl.h"
|
||||
"tensor/dynamicslice.cc"
|
||||
"tensor/image_scaler.cc"
|
||||
"tensor/image_scaler.h"
|
||||
"tensor/image_scaler_impl.cu"
|
||||
"tensor/image_scaler_impl.h"
|
||||
"transformers/beam_search.cc"
|
||||
"transformers/beam_search.h"
|
||||
"transformers/generation_device_helper.cc"
|
||||
"transformers/generation_device_helper.h"
|
||||
"transformers/beam_search_impl.cu"
|
||||
"transformers/beam_search_impl.h"
|
||||
"transformers/greedy_search.cc"
|
||||
"transformers/greedy_search.h"
|
||||
"transformers/dump_cuda_tensor.cc"
|
||||
"transformers/dump_cuda_tensor.h"
|
||||
"conv_transpose_with_dynamic_pads.cc"
|
||||
"conv_transpose_with_dynamic_pads.h"
|
||||
"cuda_contrib_kernels.cc"
|
||||
"cuda_contrib_kernels.h"
|
||||
"inverse.cc"
|
||||
"fused_conv.cc"
|
||||
)
|
||||
|
||||
set(provider_excluded_files
|
||||
"atomic/common.cuh"
|
||||
"controlflow/if.cc"
|
||||
"controlflow/if.h"
|
||||
"controlflow/loop.cc"
|
||||
"controlflow/loop.h"
|
||||
"controlflow/scan.cc"
|
||||
"controlflow/scan.h"
|
||||
"cu_inc/common.cuh"
|
||||
"math/einsum_utils/einsum_auxiliary_ops.cc"
|
||||
"math/einsum_utils/einsum_auxiliary_ops.h"
|
||||
"math/einsum_utils/einsum_auxiliary_ops_diagonal.cu"
|
||||
"math/einsum_utils/einsum_auxiliary_ops_diagonal.h"
|
||||
"math/einsum.cc"
|
||||
"math/einsum.h"
|
||||
"math/gemm.cc"
|
||||
"math/matmul.cc"
|
||||
"math/softmax_impl.cu"
|
||||
"math/softmax_warpwise_impl.cuh"
|
||||
"math/softmax_common.cc"
|
||||
"math/softmax.cc"
|
||||
"nn/conv.cc"
|
||||
"nn/conv.h"
|
||||
"nn/conv_transpose.cc"
|
||||
"nn/conv_transpose.h"
|
||||
"reduction/reduction_ops.cc"
|
||||
"rnn/cudnn_rnn_base.cc"
|
||||
"rnn/cudnn_rnn_base.h"
|
||||
"rnn/gru.cc"
|
||||
"rnn/gru.h"
|
||||
"rnn/lstm.cc"
|
||||
"rnn/lstm.h"
|
||||
"rnn/rnn.cc"
|
||||
"rnn/rnn.h"
|
||||
"rnn/rnn_impl.cu"
|
||||
"rnn/rnn_impl.h"
|
||||
"shared_inc/cuda_call.h"
|
||||
"shared_inc/fpgeneric.h"
|
||||
"cuda_allocator.cc"
|
||||
"cuda_allocator.h"
|
||||
"cuda_call.cc"
|
||||
"cuda_common.cc"
|
||||
"cuda_common.h"
|
||||
"cuda_execution_provider_info.cc"
|
||||
"cuda_execution_provider_info.h"
|
||||
"cuda_execution_provider.cc"
|
||||
"cuda_execution_provider.h"
|
||||
"cuda_memory_check.cc"
|
||||
"cuda_memory_check.h"
|
||||
"cuda_fence.cc"
|
||||
"cuda_fence.h"
|
||||
"cuda_fwd.h"
|
||||
"cuda_kernel.h"
|
||||
"cuda_pch.cc"
|
||||
"cuda_pch.h"
|
||||
"cuda_profiler.cc"
|
||||
"cuda_profiler.h"
|
||||
"cuda_provider_factory.cc"
|
||||
"cuda_provider_factory.h"
|
||||
"cuda_utils.cu"
|
||||
"cudnn_common.cc"
|
||||
"cudnn_common.h"
|
||||
"fpgeneric.cu"
|
||||
"gpu_data_transfer.cc"
|
||||
"gpu_data_transfer.h"
|
||||
"integer_gemm.cc"
|
||||
)
|
||||
|
||||
set(training_ops_excluded_files
|
||||
"activation/gelu_grad_impl_common.cuh" # uses custom tanh
|
||||
"collective/adasum_kernels.cc"
|
||||
"collective/adasum_kernels.h"
|
||||
"math/div_grad.cc" # miopen API differs from cudnn, no double type support
|
||||
"nn/batch_norm_grad.cc" # no double type support
|
||||
"nn/batch_norm_grad.h" # miopen API differs from cudnn
|
||||
"nn/batch_norm_internal.cc" # miopen API differs from cudnn, no double type support
|
||||
"nn/batch_norm_internal.h" # miopen API differs from cudnn, no double type support
|
||||
"nn/conv_grad.cc"
|
||||
"nn/conv_grad.h"
|
||||
"reduction/reduction_all.cc" # deterministic = true, ignore ctx setting
|
||||
"reduction/reduction_ops.cc" # no double type support
|
||||
"cuda_training_kernels.cc"
|
||||
"cuda_training_kernels.h"
|
||||
)
|
||||
|
||||
|
||||
# cuda_dir must be relative to REPO_ROOT
|
||||
function(hipify cuda_dir in_excluded_file_patterns out_generated_cc_files out_generated_cu_files)
|
||||
set(hipify_tool ${REPO_ROOT}/tools/ci_build/amd_hipify.py)
|
||||
|
||||
file(GLOB_RECURSE srcs CONFIGURE_DEPENDS
|
||||
"${REPO_ROOT}/${cuda_dir}/cuda/*.h"
|
||||
"${REPO_ROOT}/${cuda_dir}/cuda/*.cc"
|
||||
"${REPO_ROOT}/${cuda_dir}/cuda/*.cuh"
|
||||
"${REPO_ROOT}/${cuda_dir}/cuda/*.cu"
|
||||
)
|
||||
|
||||
# do exclusion
|
||||
set(excluded_file_patterns ${${in_excluded_file_patterns}})
|
||||
list(TRANSFORM excluded_file_patterns PREPEND "${REPO_ROOT}/${cuda_dir}/cuda/")
|
||||
file(GLOB_RECURSE excluded_srcs CONFIGURE_DEPENDS ${excluded_file_patterns})
|
||||
foreach(f ${excluded_srcs})
|
||||
message(STATUS "Excluded from hipify: ${f}")
|
||||
endforeach()
|
||||
list(REMOVE_ITEM srcs ${excluded_srcs})
|
||||
|
||||
foreach(f ${srcs})
|
||||
file(RELATIVE_PATH cuda_f_rel "${REPO_ROOT}" ${f})
|
||||
string(REPLACE "cuda" "rocm" rocm_f_rel ${cuda_f_rel})
|
||||
set(f_out "${CMAKE_CURRENT_BINARY_DIR}/amdgpu/${rocm_f_rel}")
|
||||
add_custom_command(
|
||||
OUTPUT ${f_out}
|
||||
COMMAND Python3::Interpreter ${hipify_tool}
|
||||
--hipify_perl ${onnxruntime_HIPIFY_PERL}
|
||||
${f} -o ${f_out}
|
||||
DEPENDS ${hipify_tool} ${f}
|
||||
COMMENT "Hipify: ${cuda_f_rel} -> amdgpu/${rocm_f_rel}"
|
||||
)
|
||||
if(f MATCHES "\\..*cuh?")
|
||||
list(APPEND generated_cu_files ${f_out})
|
||||
else()
|
||||
list(APPEND generated_cc_files ${f_out})
|
||||
endif()
|
||||
endforeach()
|
||||
|
||||
set_source_files_properties(generated_cc_files PROPERTIES GENERATED TRUE)
|
||||
set_source_files_properties(generated_cu_files PROPERTIES GENERATED TRUE)
|
||||
set(${out_generated_cc_files} ${generated_cc_files} PARENT_SCOPE)
|
||||
set(${out_generated_cu_files} ${generated_cu_files} PARENT_SCOPE)
|
||||
endfunction()
|
||||
|
|
@ -1,232 +1,18 @@
|
|||
# Copyright (c) Microsoft Corporation. All rights reserved.
|
||||
# Licensed under the MIT License.
|
||||
|
||||
import concurrent.futures
|
||||
import fnmatch
|
||||
import functools
|
||||
import argparse
|
||||
import os
|
||||
import shutil
|
||||
import subprocess
|
||||
|
||||
from logger import get_logger
|
||||
|
||||
log = get_logger("amd_hipify")
|
||||
|
||||
|
||||
def path_in_repo(path):
|
||||
repo_root = os.path.relpath(os.path.join(os.path.dirname(__file__), "../.."))
|
||||
return os.path.join(repo_root, path)
|
||||
|
||||
|
||||
contrib_ops_path = path_in_repo("onnxruntime/contrib_ops")
|
||||
providers_path = path_in_repo("onnxruntime/core/providers")
|
||||
training_ops_path = path_in_repo("orttraining/orttraining/training_ops")
|
||||
|
||||
|
||||
def is_excluded(f, excluded_patterns):
|
||||
return any([fnmatch.fnmatch(f, pat) for pat in excluded_patterns])
|
||||
|
||||
|
||||
contrib_ops_excluded_files = [
|
||||
"bert/attention.cc",
|
||||
"bert/attention.h",
|
||||
"bert/attention_impl.cu",
|
||||
"bert/attention_softmax.h",
|
||||
"bert/embed_layer_norm.cc",
|
||||
"bert/embed_layer_norm.h",
|
||||
"bert/embed_layer_norm_impl.cu",
|
||||
"bert/embed_layer_norm_impl.h",
|
||||
"bert/fast_gelu_impl.cu",
|
||||
"bert/fast_gelu_impl.h",
|
||||
"bert/fast_gelu.cc",
|
||||
"bert/fast_gelu.h",
|
||||
"bert/skip_layer_norm.cc",
|
||||
"bert/skip_layer_norm.h",
|
||||
"bert/skip_layer_norm_impl.cu",
|
||||
"bert/skip_layer_norm_impl.h",
|
||||
"bert/tensorrt_fused_multihead_attention/*",
|
||||
"bert/transformer_common.h",
|
||||
"bert/transformer_common.cc",
|
||||
"math/complex_mul.cc",
|
||||
"math/complex_mul.h",
|
||||
"math/complex_mul_impl.cu",
|
||||
"math/complex_mul_impl.h",
|
||||
"math/cufft_plan_cache.h",
|
||||
"math/fft_ops.cc",
|
||||
"math/fft_ops.h",
|
||||
"math/fft_ops_impl.cu",
|
||||
"math/fft_ops_impl.h",
|
||||
"quantization/attention_quantization.cc",
|
||||
"quantization/attention_quantization.h",
|
||||
"quantization/attention_quantization_impl.cu",
|
||||
"quantization/attention_quantization_impl.cuh",
|
||||
"quantization/quantize_dequantize_linear.cc",
|
||||
"quantization/qordered_ops/qordered_attention_impl.cu",
|
||||
"quantization/qordered_ops/qordered_attention_impl.h",
|
||||
"quantization/qordered_ops/qordered_attention_input_enum.h",
|
||||
"quantization/qordered_ops/qordered_attention.cc",
|
||||
"quantization/qordered_ops/qordered_attention.h",
|
||||
"quantization/qordered_ops/qordered_common.cuh",
|
||||
"quantization/qordered_ops/qordered_layer_norm.h",
|
||||
"quantization/qordered_ops/qordered_layer_norm.cc",
|
||||
"quantization/qordered_ops/qordered_layer_norm_impl.h",
|
||||
"quantization/qordered_ops/qordered_layer_norm_impl.cu",
|
||||
"quantization/qordered_ops/qordered_longformer_attention.cc",
|
||||
"quantization/qordered_ops/qordered_longformer_attention.h",
|
||||
"quantization/qordered_ops/qordered_matmul.h",
|
||||
"quantization/qordered_ops/qordered_matmul.cc",
|
||||
"quantization/qordered_ops/qordered_matmul_utils.h",
|
||||
"quantization/qordered_ops/qordered_matmul_utils.cc",
|
||||
"quantization/qordered_ops/qordered_qdq_impl.cu",
|
||||
"quantization/qordered_ops/qordered_qdq_impl.h",
|
||||
"quantization/qordered_ops/qordered_qdq.cc",
|
||||
"quantization/qordered_ops/qordered_qdq.h",
|
||||
"quantization/qordered_ops/qordered_unary_ops.h",
|
||||
"quantization/qordered_ops/qordered_unary_ops.cc",
|
||||
"quantization/qordered_ops/qordered_unary_ops_impl.h",
|
||||
"quantization/qordered_ops/qordered_unary_ops_impl.cu",
|
||||
"tensor/crop.cc",
|
||||
"tensor/crop.h",
|
||||
"tensor/crop_impl.cu",
|
||||
"tensor/crop_impl.h",
|
||||
"tensor/dynamicslice.cc",
|
||||
"tensor/image_scaler.cc",
|
||||
"tensor/image_scaler.h",
|
||||
"tensor/image_scaler_impl.cu",
|
||||
"tensor/image_scaler_impl.h",
|
||||
"transformers/beam_search.cc",
|
||||
"transformers/beam_search.h",
|
||||
"transformers/generation_device_helper.cc",
|
||||
"transformers/generation_device_helper.h",
|
||||
"transformers/beam_search_impl.cu",
|
||||
"transformers/beam_search_impl.h",
|
||||
"transformers/greedy_search.cc",
|
||||
"transformers/greedy_search.h",
|
||||
"transformers/dump_cuda_tensor.cc",
|
||||
"transformers/dump_cuda_tensor.h",
|
||||
"conv_transpose_with_dynamic_pads.cc",
|
||||
"conv_transpose_with_dynamic_pads.h",
|
||||
"cuda_contrib_kernels.cc",
|
||||
"cuda_contrib_kernels.h",
|
||||
"inverse.cc",
|
||||
"fused_conv.cc",
|
||||
]
|
||||
|
||||
provider_excluded_files = [
|
||||
"atomic/common.cuh",
|
||||
"controlflow/if.cc",
|
||||
"controlflow/if.h",
|
||||
"controlflow/loop.cc",
|
||||
"controlflow/loop.h",
|
||||
"controlflow/scan.cc",
|
||||
"controlflow/scan.h",
|
||||
"cu_inc/common.cuh",
|
||||
"math/einsum_utils/einsum_auxiliary_ops.cc",
|
||||
"math/einsum_utils/einsum_auxiliary_ops.h",
|
||||
"math/einsum_utils/einsum_auxiliary_ops_diagonal.cu",
|
||||
"math/einsum_utils/einsum_auxiliary_ops_diagonal.h",
|
||||
"math/einsum.cc",
|
||||
"math/einsum.h",
|
||||
"math/gemm.cc",
|
||||
"math/matmul.cc",
|
||||
"math/softmax_impl.cu",
|
||||
"math/softmax_warpwise_impl.cuh",
|
||||
"math/softmax_common.cc",
|
||||
"math/softmax.cc",
|
||||
"nn/conv.cc",
|
||||
"nn/conv.h",
|
||||
"nn/conv_transpose.cc",
|
||||
"nn/conv_transpose.h",
|
||||
"reduction/reduction_ops.cc",
|
||||
"rnn/cudnn_rnn_base.cc",
|
||||
"rnn/cudnn_rnn_base.h",
|
||||
"rnn/gru.cc",
|
||||
"rnn/gru.h",
|
||||
"rnn/lstm.cc",
|
||||
"rnn/lstm.h",
|
||||
"rnn/rnn.cc",
|
||||
"rnn/rnn.h",
|
||||
"rnn/rnn_impl.cu",
|
||||
"rnn/rnn_impl.h",
|
||||
"shared_inc/cuda_call.h",
|
||||
"shared_inc/fpgeneric.h",
|
||||
"cuda_allocator.cc",
|
||||
"cuda_allocator.h",
|
||||
"cuda_call.cc",
|
||||
"cuda_common.cc",
|
||||
"cuda_common.h",
|
||||
"cuda_execution_provider_info.cc",
|
||||
"cuda_execution_provider_info.h",
|
||||
"cuda_execution_provider.cc",
|
||||
"cuda_execution_provider.h",
|
||||
"cuda_memory_check.cc",
|
||||
"cuda_memory_check.h",
|
||||
"cuda_fence.cc",
|
||||
"cuda_fence.h",
|
||||
"cuda_fwd.h",
|
||||
"cuda_kernel.h",
|
||||
"cuda_pch.cc",
|
||||
"cuda_pch.h",
|
||||
"cuda_profiler.cc",
|
||||
"cuda_profiler.h",
|
||||
"cuda_provider_factory.cc",
|
||||
"cuda_provider_factory.h",
|
||||
"cuda_utils.cu",
|
||||
"cudnn_common.cc",
|
||||
"cudnn_common.h",
|
||||
"fpgeneric.cu",
|
||||
"gpu_data_transfer.cc",
|
||||
"gpu_data_transfer.h",
|
||||
"integer_gemm.cc",
|
||||
"symbols.txt",
|
||||
]
|
||||
|
||||
training_ops_excluded_files = [
|
||||
"activation/gelu_grad_impl_common.cuh", # uses custom tanh
|
||||
"collective/adasum_kernels.cc",
|
||||
"collective/adasum_kernels.h",
|
||||
"math/div_grad.cc", # miopen API differs from cudnn, no double type support
|
||||
"nn/batch_norm_grad.cc", # no double type support
|
||||
"nn/batch_norm_grad.h", # miopen API differs from cudnn
|
||||
"nn/batch_norm_internal.cc", # miopen API differs from cudnn, no double type support
|
||||
"nn/batch_norm_internal.h", # miopen API differs from cudnn, no double type support
|
||||
"nn/conv_grad.cc",
|
||||
"nn/conv_grad.h",
|
||||
"reduction/reduction_all.cc", # deterministic = true, ignore ctx setting
|
||||
"reduction/reduction_ops.cc", # no double type support
|
||||
"cuda_training_kernels.cc",
|
||||
"cuda_training_kernels.h",
|
||||
]
|
||||
|
||||
|
||||
@functools.lru_cache(maxsize=1)
|
||||
def get_hipify_path():
|
||||
# prefer the hipify-perl in PATH
|
||||
HIPIFY_PERL = shutil.which("hipify-perl")
|
||||
# if not found, attempt hard-coded location 1
|
||||
if HIPIFY_PERL is None:
|
||||
print("hipify-perl not found, trying default location 1")
|
||||
hipify_path = "/opt/rocm/hip/bin/hipify-perl"
|
||||
HIPIFY_PERL = hipify_path if os.access(hipify_path, os.X_OK) else None
|
||||
# if not found, attempt hard-coded location 2
|
||||
if HIPIFY_PERL is None:
|
||||
print("hipify-perl not found, trying default location 2")
|
||||
hipify_path = "/opt/rocm/bin/hipify-perl"
|
||||
HIPIFY_PERL = hipify_path if os.access(hipify_path, os.X_OK) else None
|
||||
# fail
|
||||
if HIPIFY_PERL is None:
|
||||
raise RuntimeError("Could not locate hipify-perl script")
|
||||
return HIPIFY_PERL
|
||||
|
||||
|
||||
def hipify(src_file_path, dst_file_path):
|
||||
dst_file_path = dst_file_path.replace("cuda", "rocm")
|
||||
def hipify(hipify_perl_path, src_file_path, dst_file_path):
|
||||
dir_name = os.path.dirname(dst_file_path)
|
||||
if not os.path.exists(dir_name):
|
||||
os.makedirs(dir_name, exist_ok=True)
|
||||
# Run hipify-perl first, capture output
|
||||
s = subprocess.run(
|
||||
[get_hipify_path(), "-roc", src_file_path], stdout=subprocess.PIPE, universal_newlines=True, check=False
|
||||
[hipify_perl_path, "-roc", src_file_path], stdout=subprocess.PIPE, universal_newlines=True, check=False
|
||||
).stdout
|
||||
|
||||
# Additional exact-match replacements.
|
||||
|
|
@ -374,71 +160,15 @@ def hipify(src_file_path, dst_file_path):
|
|||
# Deletions
|
||||
s = s.replace('#include "device_atomic_functions.h"', "") # HIP atomics in main hip header already
|
||||
|
||||
do_write = True
|
||||
if os.path.exists(dst_file_path):
|
||||
with open(dst_file_path, "r", encoding="utf-8") as fout_old:
|
||||
do_write = fout_old.read() != s
|
||||
if do_write:
|
||||
with open(dst_file_path, "w") as f:
|
||||
f.write(s)
|
||||
return 'Hipified: "{}" -> "{}"'.format(src_file_path, dst_file_path)
|
||||
else:
|
||||
return 'Repeated: "{}" -> "{}"'.format(src_file_path, dst_file_path)
|
||||
|
||||
|
||||
def list_files(prefix, path):
|
||||
all_files = []
|
||||
curr_path = os.path.join(prefix, path)
|
||||
for root, dirs, files in os.walk(curr_path):
|
||||
for file in files:
|
||||
full_path = os.path.join(root, file)
|
||||
all_files.append(os.path.relpath(full_path, curr_path))
|
||||
return all_files
|
||||
|
||||
|
||||
def amd_hipify(config_build_dir):
|
||||
# determine hipify script path now to avoid doing so concurrently in the thread pool
|
||||
print("Using %s" % get_hipify_path())
|
||||
with concurrent.futures.ThreadPoolExecutor() as executor:
|
||||
cuda_path = os.path.join(contrib_ops_path, "cuda")
|
||||
rocm_path = os.path.join(config_build_dir, "amdgpu", contrib_ops_path, "rocm")
|
||||
contrib_files = list_files(cuda_path, "")
|
||||
contrib_results = [
|
||||
executor.submit(hipify, os.path.join(cuda_path, f), os.path.join(rocm_path, f))
|
||||
for f in contrib_files
|
||||
if not is_excluded(f, contrib_ops_excluded_files)
|
||||
]
|
||||
|
||||
cuda_path = os.path.join(providers_path, "cuda")
|
||||
rocm_path = os.path.join(config_build_dir, "amdgpu", providers_path, "rocm")
|
||||
provider_files = list_files(cuda_path, "")
|
||||
provider_results = [
|
||||
executor.submit(hipify, os.path.join(cuda_path, f), os.path.join(rocm_path, f))
|
||||
for f in provider_files
|
||||
if not is_excluded(f, provider_excluded_files)
|
||||
]
|
||||
|
||||
cuda_path = os.path.join(training_ops_path, "cuda")
|
||||
rocm_path = os.path.join(config_build_dir, "amdgpu", training_ops_path, "rocm")
|
||||
training_files = list_files(cuda_path, "")
|
||||
training_results = [
|
||||
executor.submit(hipify, os.path.join(cuda_path, f), os.path.join(rocm_path, f))
|
||||
for f in training_files
|
||||
if not is_excluded(f, training_ops_excluded_files)
|
||||
]
|
||||
# explicitly wait so that hipify warnings finish printing before logging the hipify statements
|
||||
concurrent.futures.wait(contrib_results)
|
||||
concurrent.futures.wait(provider_results)
|
||||
concurrent.futures.wait(training_results)
|
||||
for result in contrib_results:
|
||||
log.debug(result.result())
|
||||
for result in provider_results:
|
||||
log.debug(result.result())
|
||||
for result in training_results:
|
||||
log.debug(result.result())
|
||||
with open(dst_file_path, "w") as f:
|
||||
f.write(s)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
import sys
|
||||
parser = argparse.ArgumentParser()
|
||||
parser.add_argument("--hipify_perl", required=True)
|
||||
parser.add_argument("--output", "-o", help="output file")
|
||||
parser.add_argument("src", help="src")
|
||||
args = parser.parse_args()
|
||||
|
||||
amd_hipify(sys.argv[1])
|
||||
hipify(args.hipify_perl, args.src, args.output)
|
||||
|
|
|
|||
|
|
@ -14,8 +14,6 @@ import sys
|
|||
from distutils.version import LooseVersion
|
||||
from pathlib import Path
|
||||
|
||||
from amd_hipify import amd_hipify
|
||||
|
||||
SCRIPT_DIR = os.path.dirname(os.path.realpath(__file__))
|
||||
REPO_DIR = os.path.normpath(os.path.join(SCRIPT_DIR, "..", ".."))
|
||||
|
||||
|
|
@ -1467,24 +1465,17 @@ def setup_dml_build(args, cmake_path, build_dir, configs):
|
|||
run_subprocess(cmd_args)
|
||||
|
||||
|
||||
def setup_rocm_build(args, configs):
|
||||
|
||||
def setup_rocm_build(args):
|
||||
rocm_home = None
|
||||
|
||||
if args.use_rocm:
|
||||
print("rocm_home = {}".format(args.rocm_home))
|
||||
rocm_home = args.rocm_home or None
|
||||
|
||||
rocm_home_not_valid = rocm_home and not os.path.exists(rocm_home)
|
||||
|
||||
if rocm_home_not_valid:
|
||||
raise BuildError(
|
||||
"rocm_home paths must be specified and valid.",
|
||||
"rocm_home='{}' valid={}.".format(rocm_home, rocm_home_not_valid),
|
||||
)
|
||||
|
||||
for config in configs:
|
||||
amd_hipify(get_config_build_dir(args.build_dir, config))
|
||||
return rocm_home or ""
|
||||
|
||||
|
||||
|
|
@ -2506,7 +2497,7 @@ def main():
|
|||
migraphx_home = setup_migraphx_vars(args)
|
||||
|
||||
# if using rocm, setup rocm paths
|
||||
rocm_home = setup_rocm_build(args, configs)
|
||||
rocm_home = setup_rocm_build(args)
|
||||
|
||||
# if using cann, setup cann paths
|
||||
cann_home = setup_cann_vars(args)
|
||||
|
|
|
|||
Loading…
Reference in a new issue