diff --git a/onnxruntime/core/providers/cuda/math/topk_impl.cu b/onnxruntime/core/providers/cuda/math/topk_impl.cu index 9f81f428dd..f24a0f9ffb 100644 --- a/onnxruntime/core/providers/cuda/math/topk_impl.cu +++ b/onnxruntime/core/providers/cuda/math/topk_impl.cu @@ -56,14 +56,15 @@ struct NumericLimits { template __global__ void BitonicTopK(const T* X, T* V, int64_t* I, const TArray elem_nums, size_t size, int32_t axis, int64_t K, int64_t aligned_K, int64_t largest, int64_t sorted, int64_t dimension, int64_t aligned_dimension, T type_min, T type_max) { - auto tid = threadIdx.x; - auto bid = blockIdx.x; + int64_t tid = threadIdx.x; + int64_t bid = blockIdx.x; + int64_t bdim = blockDim.x; extern __shared__ char shared_mem[]; auto S = (KV*)(shared_mem); auto mid_dim = axis == size - 1 ? 1 : elem_nums[axis + 1]; auto left_dim = bid / mid_dim * elem_nums[axis]; auto right_dim = axis == size - 1 ? 0 : bid % elem_nums[axis + 1]; - for (auto i = tid; i < aligned_dimension; i += blockDim.x) { + for (auto i = tid; i < aligned_dimension; i += bdim) { S[i].key = i < dimension ? X[FROM(i)] : TRIVIAL; S[i].val = i; } diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc index 0d60177c3b..5ba380f0e4 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc @@ -1551,7 +1551,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, // BuildKernelCreateInfo, // BuildKernelCreateInfo, // BuildKernelCreateInfo, @@ -1581,7 +1581,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { // BuildKernelCreateInfo, // BuildKernelCreateInfo, // BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, // BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1659,7 +1659,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/tools/ci_build/amd_hipify.py b/tools/ci_build/amd_hipify.py index 99de6b213c..bd88a243ed 100644 --- a/tools/ci_build/amd_hipify.py +++ b/tools/ci_build/amd_hipify.py @@ -90,10 +90,6 @@ provider_excluded_files = [ 'math/matmul_integer.h', 'math/softmax_impl.cu', 'math/softmax.cc', - 'math/topk.cc', - 'math/topk.h', - 'math/topk_impl.cu', - 'math/topk_impl.h', 'nn/batch_norm.cc', 'nn/batch_norm.h', 'nn/conv.cc', @@ -244,6 +240,8 @@ def hipify(src_file_path, dst_file_path): s = s.replace('std::log', 'logf') s = s.replace('#include ', '#include \n#include ') + s = s.replace('#include "cub/device/device_radix_sort.cuh"', + '#include \n#include ') s = s.replace('#include ', '#include ') s = s.replace('#include ', @@ -254,6 +252,14 @@ def hipify(src_file_path, dst_file_path): '#include ') s = s.replace('#include ', '#include ') + s = s.replace('#include ', + '#include ') + s = s.replace('#include "cub/util_allocator.cuh"', + '#include ') + s = s.replace('#include ', + '#include ') + s = s.replace('#include "cub/util_type.cuh"', + '#include ') s = s.replace('typedef half MappedType', 'typedef __half MappedType') # CUBLAS -> HIPBLAS @@ -301,6 +307,9 @@ def hipify(src_file_path, dst_file_path): s = s.replace('ROCM_VERSION', 'CUDA_VERSION') # semantically different meanings, cannot hipify s = s.replace('__ROCM_ARCH__', '__CUDA_ARCH__') # semantically different meanings, cannot hipify + # 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: