saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-29 20:14:01 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Merge remote-tracking branch 'upstream/master' into DmlDev

Browse source
This commit is contained in:
ISS Build Account 2020-11-10 00:46:04 +00:00
commit de85638543
14 changed files with 1172 additions and 980 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

31
include/onnxruntime/core/framework/op_kernel.h
View file

@ -226,6 +226,13 @@ class OpKernelContext {
*/
_Ret_maybenull_ onnxruntime::concurrency::ThreadPool* GetOperatorThreadPool() const { return threadpool_; }
/**
Returns whether deterministic computation is preferred.
*/
virtual bool GetUseDeterministicCompute() const {
return true;
}
protected:
onnxruntime::NodeIndex GetNodeIndex() const;
@ -442,18 +449,18 @@ using BuildKernelCreateInfoFn = KernelCreateInfo (*)();
#define ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_CLASS_NAME(provider, domain, startver, endver, type1, type2, name) \
provider##_##name##_##domain##_ver##startver##_##endver##_##type1##_##type2
#define ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_EX(name, domain, startver, endver, type1, type2, provider, builder, ...) \
class ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_CLASS_NAME(provider, domain, startver, endver, type1, type2, name); \
template <> \
KernelCreateInfo \
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_CLASS_NAME(provider, domain, startver, endver, type1, type2, name)>() { \
return KernelCreateInfo( \
builder.SetName(#name) \
.SetDomain(domain) \
.SinceVersion(startver, endver) \
.Provider(provider) \
.Build(), \
static_cast<KernelCreatePtrFn>([](const OpKernelInfo& info) -> OpKernel* { return new __VA_ARGS__(info); })); \
#define ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_EX(name, domain, startver, endver, type1, type2, provider, builder, ...) \
class ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_CLASS_NAME(provider, domain, startver, endver, type1, type2, name); \
template <> \
KernelCreateInfo \
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TWO_TYPED_KERNEL_CLASS_NAME(provider, domain, startver, endver, type1, type2, name)>() { \
return KernelCreateInfo( \
builder.SetName(#name) \
.SetDomain(domain) \
.SinceVersion(startver, endver) \
.Provider(provider) \
.Build(), \
static_cast<KernelCreatePtrFn>([](const OpKernelInfo& info) -> OpKernel* { return new __VA_ARGS__(info); })); \
}
// Use within macro definitions to create a custom vector of constraints.

4
onnxruntime/core/framework/op_kernel_context_internal.h
View file

@ -37,10 +37,10 @@ class OpKernelContextInternal : public OpKernelContext {
}
}
bool GetUseDeterministicCompute() const {
bool GetUseDeterministicCompute() const override {
return session_state_.GetUseDeterministicCompute();
}
const SessionState* SubgraphSessionState(const std::string& attribute_name) {
return session_state_.GetSubgraphSessionState(GetNodeIndex(), attribute_name);
}

1805
onnxruntime/core/providers/cpu/cpu_execution_provider.cc
View file

File diff suppressed because it is too large Load diff

27
onnxruntime/core/providers/cpu/math/element_wise_ops.cc
View file

@ -191,8 +191,10 @@ REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(Min, 12, 12, Min_8, float, double, MLFloat
// Supposed to add BFloat16 but we are not supporting now, however, separate registration
REG_ELEMENTWISE_KERNEL_NONT(Min, 13, Min_8, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 12, float, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 12, double, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 8, float, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 8, double, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 9, 12, float, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 9, 12, double, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 9, 12, int32_t, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 9, 12, int64_t, Less);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Less, 13, float, Less);
@ -200,8 +202,10 @@ REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Less, 13, double, Less);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Less, 13, int32_t, Less);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Less, 13, int64_t, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 7, 12, float, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 7, 12, double, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 7, 8, float, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 7, 8, double, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 9, 12, float, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 9, 12, double, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 9, 12, int32_t, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Greater, 9, 12, int64_t, Greater);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Greater, 13, float, Greater);
@ -209,11 +213,16 @@ REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Greater, 13, double, Greater);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Greater, 13, int32_t, Greater);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Greater, 13, int64_t, Greater);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 12, bool, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 12, int32_t, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 12, int64_t, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 12, float, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 12, double, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 10, bool, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 10, int32_t, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 10, int64_t, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 10, float, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 7, 10, double, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 11, 12, bool, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 11, 12, int32_t, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 11, 12, int64_t, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 11, 12, float, Equal);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Equal, 11, 12, double, Equal);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Equal, 13, bool, Equal);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Equal, 13, int32_t, Equal);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Equal, 13, int64_t, Equal);

7
onnxruntime/core/providers/cuda/reduction/reduction_ops.cc
View file

@ -661,12 +661,7 @@ Status ReduceKernel<allow_multi_axes>::ComputeImpl(OpKernelContext* ctx, cudnnRe
axes,
prepare_reduce_metadata));
Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims);
bool fast_reduction = fast_reduction_;
if (fast_reduction) {
auto ctx_internal = dynamic_cast<OpKernelContextInternal*>(ctx);
if (ctx_internal && ctx_internal->GetUseDeterministicCompute())
fast_reduction = false;
}
const bool fast_reduction = fast_reduction_ && !ctx->GetUseDeterministicCompute();
return ReduceComputeCore<T, ReduceTensorIndices>(*cuda_ep_, *X, prepare_reduce_metadata, *Y, cudnn_reduce_op, axes_,
calculate_log_, calculate_sqt_, log_sum_exp_, fast_reduction);

27
onnxruntime/core/providers/rocm/reduction/reduction_ops.cc
View file

@ -27,7 +27,7 @@ namespace rocm {
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
name, \
kOnnxDomain, \
11, 12, \
11, 12, \
T, \
kRocmExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
@ -133,20 +133,19 @@ Status ReduceKernel<allow_multi_axes>::ReduceKernelShared(
typedef typename ToHipType<T>::MappedType HipT;
const auto rank = input_shape.NumDimensions();
// Block of fast matrix row reduction.
// It relies on new atomicAdd for half type, so old hip can't use it.
// Block of fast matrix row reduction.
// It relies on new atomicAdd for half type, so old hip can't use it.
const auto stride = input_shape[input_shape.NumDimensions() - 1];
const auto reduction_size = input_shape.Size() / stride;
if (fast_reduction_ && reduction_size <= std::numeric_limits<int>::max() && stride <= std::numeric_limits<int>::max() &&
is_matrix_row_reduction(miopen_reduce_op,
static_cast<int>(reduction_size),
static_cast<int>(stride), rank, axes_)) {
static_cast<int>(reduction_size),
static_cast<int>(stride), rank, axes_)) {
reduce_matrix_rows(
reinterpret_cast<const HipT*>(X),
reinterpret_cast<HipT*>(Y),
static_cast<int>(reduction_size),
static_cast<int>(stride));
reinterpret_cast<const HipT*>(X),
reinterpret_cast<HipT*>(Y),
static_cast<int>(reduction_size),
static_cast<int>(stride));
return Status::OK();
}
@ -330,7 +329,6 @@ Status ReduceComputeCore(const Tensor& input, PrepareReduceMetadata& prepare_red
return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "reduction2 is not supported");
}
template <bool allow_multi_axes>
template <typename T>
Status ReduceKernel<allow_multi_axes>::ComputeImpl(OpKernelContext* ctx, miopenReduceTensorOp_t miopen_reduce_op) const {
@ -342,12 +340,7 @@ Status ReduceKernel<allow_multi_axes>::ComputeImpl(OpKernelContext* ctx, miopenR
axes_,
prepare_reduce_metadata));
Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims);
bool fast_reduction = fast_reduction_;
if (fast_reduction) {
auto ctx_internal = static_cast<OpKernelContextInternal*>(ctx);
if (ctx_internal && ctx_internal->GetUseDeterministicCompute())
fast_reduction = false;
}
const bool fast_reduction = fast_reduction_ && !ctx->GetUseDeterministicCompute();
return ReduceComputeCore<T>(*X, prepare_reduce_metadata, *Y, miopen_reduce_op, axes_,
calculate_log_, calculate_sqt_, log_sum_exp_, fast_reduction);

68
onnxruntime/core/providers/rocm/reduction/reduction_utils.cuh Normal file
View file

@ -0,0 +1,68 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#pragma once
#include "core/providers/rocm/cu_inc/common.cuh"
namespace onnxruntime {
namespace rocm {
__forceinline__ __host__ __device__ int least_pow2_bound(int value) {
unsigned int value_ = static_cast<unsigned int>(value);
--value_;
value_ |= value_ >> 1;
value_ |= value_ >> 2;
value_ |= value_ >> 4;
value_ |= value_ >> 8;
value_ |= value_ >> 16;
return static_cast<int>(++value_);
}
template <typename TAccumulated, typename TValue>
struct Cast {
__forceinline__ __device__ TAccumulated operator()(const TValue& value) {
return TAccumulated(value);
}
};
template <typename TAccumulated, typename TValue>
struct Square {
__forceinline__ __device__ TAccumulated operator()(const TValue& value) {
return TAccumulated(value) * TAccumulated(value);
}
};
template <typename TAccumulated, typename TValue>
struct Abs {
__forceinline__ __device__ TAccumulated operator()(const TValue& value) {
TAccumulated value_ = TAccumulated(value);
return value_ > TAccumulated(0) ? value_ : -value_;
}
};
template <typename T>
struct Sqrt {
__forceinline__ __device__ T operator()(const T& value) {
return _Sqrt(value);
}
};
template <typename T>
struct Identity {
__forceinline__ __device__ T operator()(const T& value) {
return value;
}
};
template <typename T>
struct ToBuffer {
typedef T Type;
};
template <>
struct ToBuffer<half> {
typedef float Type;
};
} // namespace rocm
} // namespace onnxruntime

46
onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc vendored
View file

@ -56,16 +56,6 @@
"^test_logsoftmax_large_number_expanded_cpu",
"^test_logsoftmax_negative_axis_cpu",
"^test_logsoftmax_negative_axis_expanded_cpu",
"^test_reduce_sum_default_axes_keepdims_example_cpu", // NOT_IMPLEMENTED : Could not find an implementation for the node ReduceSum(13)
"^test_reduce_sum_default_axes_keepdims_random_cpu",
"^test_reduce_sum_do_not_keepdims_example_cpu",
"^test_reduce_sum_do_not_keepdims_random_cpu",
"^test_reduce_sum_empty_axes_input_noop_example_cpu",
"^test_reduce_sum_empty_axes_input_noop_random_cpu",
"^test_reduce_sum_keepdims_example_cpu",
"^test_reduce_sum_keepdims_random_cpu",
"^test_reduce_sum_negative_axes_keepdims_example_cpu",
"^test_reduce_sum_negative_axes_keepdims_random_cpu",
"^test_resize_downsample_scales_cubic_A_n0p5_exclude_outside_cpu", // NOT_IMPLEMENTED : Could not find an implementation for the node Resize(13)
"^test_resize_downsample_scales_cubic_cpu",
"^test_resize_downsample_scales_linear_cpu",
@ -103,23 +93,7 @@
"^test_softmax_large_number_cpu",
"^test_softmax_large_number_expanded_cpu",
"^test_softmax_negative_axis_cpu",
"^test_softmax_negative_axis_expanded_cpu",
"^test_split_equal_parts_1d_cpu", // NOT_IMPLEMENTED : Could not find an implementation for the node Split(13)
"^test_split_equal_parts_2d_cpu",
"^test_split_equal_parts_default_axis_cpu",
"^test_split_variable_parts_1d_cpu",
"^test_split_variable_parts_2d_cpu",
"^test_split_variable_parts_default_axis_cpu",
"^test_split_zero_size_splits_cpu",
"^test_squeeze_cpu", // NOT_IMPLEMENTED : Could not find an implementation for the node Squeeze(13)
"^test_squeeze_negative_axes_cpu",
"^test_unsqueeze_axis_0_cpu", // NOT_IMPLEMENTED : Could not find an implementation for the node Unsqueeze(13)
"^test_unsqueeze_axis_1_cpu",
"^test_unsqueeze_axis_2_cpu",
"^test_unsqueeze_negative_axes_cpu",
"^test_unsqueeze_three_axes_cpu",
"^test_unsqueeze_two_axes_cpu",
"^test_unsqueeze_unsorted_axes_cpu"
"^test_softmax_negative_axis_expanded_cpu"
],
"current_failing_tests_x86": [
"^test_vgg19",
@ -203,7 +177,14 @@
"^test_softmax_cross_entropy.*", // Does not support 5-D or above tensors for SUB op.
"^test_sce.*",
"^test_nllloss.*",
"^test_gather_negative_indices.*"
"^test_gather_negative_indices.*",
"^test_reduce_sum_do_not_keepdims*", // Does not support axes as input
"^test_reduce_sum_keepdims*",
"^test_reduce_sum_default_axes_keepdims*",
"^test_reduce_sum_negative_axes_keepdims*",
"^test_reduce_sum_empty_axes_input_noop*",
"^test_unsqueeze_*", // Does not support axes as input
"^test_squeeze_*" // Does not support axes as input
],
"current_failing_tests_OPENVINO_CPU_FP32": [
"^test_operator_permute2",
@ -215,7 +196,14 @@
"^test_gather_negative_indices.*",
"^test_sce.*",
"^test_nllloss.*",
"^test_upsample_nearest.*"
"^test_upsample_nearest.*",
"^test_reduce_sum_do_not_keepdims*", // Does not support axes as input
"^test_reduce_sum_keepdims*",
"^test_reduce_sum_default_axes_keepdims*",
"^test_reduce_sum_negative_axes_keepdims*",
"^test_reduce_sum_empty_axes_input_noop*",
"^test_unsqueeze_*", // Does not support axes as input
"^test_squeeze_*" // Does not support axes as input
],
// ORT first supported opset 7, so models with nodes that require versions prior to opset 7 are not supported
"tests_with_pre_opset7_dependencies": [

3
orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cc
View file

@ -46,8 +46,7 @@ Status ReduceAllL2<TIn, TOut>::ComputeInternal(OpKernelContext* ctx) const {
CudaTOut* p_output = reinterpret_cast<CudaTOut*>(output->template MutableData<TOut>());
CUDA_RETURN_IF_ERROR(cudaMemsetAsync(p_output, 0, sizeof(CudaTOut)));
auto ctx_internal = dynamic_cast<OpKernelContextInternal*>(ctx);
bool deterministic = ctx_internal && ctx_internal->GetUseDeterministicCompute();
const bool deterministic = ctx->GetUseDeterministicCompute();
if (!deterministic) {
typedef MultiTensorReduceL2<CudaTIn, CudaTOut> TFunctor;

7
orttraining/orttraining/training_ops/cuda/reduction/reduction_ops.cc
View file

@ -57,12 +57,7 @@ Status ReduceKernel<allow_multi_axes>::ComputeImplEx(OpKernelContext* ctx, cudnn
axes,
prepare_reduce_metadata));
Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims);
bool fast_reduction = fast_reduction_;
if (fast_reduction) {
auto ctx_internal = dynamic_cast<OpKernelContextInternal*>(ctx);
if (ctx_internal && ctx_internal->GetUseDeterministicCompute())
fast_reduction = false;
}
const bool fast_reduction = fast_reduction_ && !ctx->GetUseDeterministicCompute();
return ReduceComputeCore<T, ReduceTensorIndices>(*cuda_ep_, *X, prepare_reduce_metadata, *Y, cudnn_reduce_op, axes,
calculate_log_, calculate_sqt_, log_sum_exp_, fast_reduction);

3
orttraining/orttraining/training_ops/rocm/reduction/reduction_all.cc
View file

@ -55,8 +55,7 @@ Status ReduceAllL2<TIn, TOut>::ComputeInternal(OpKernelContext* ctx) const {
HipTOut* p_output = reinterpret_cast<HipTOut*>(output->template MutableData<TOut>());
HIP_RETURN_IF_ERROR(hipMemsetAsync(p_output, 0, sizeof(HipTOut)));
// auto ctx_internal = static_cast<OpKernelContextInternal*>(ctx);
// bool deterministic = ctx_internal && ctx_internal->GetUseDeterministicCompute();
// bool deterministic = ctx->GetUseDeterministicCompute();
bool deterministic = true;
if (!deterministic) {

114
orttraining/orttraining/training_ops/rocm/reduction/reduction_all.cu Normal file
View file

@ -0,0 +1,114 @@
#include "hip/hip_runtime.h"
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "orttraining/training_ops/rocm/reduction/reduction_all.h"
#include "core/providers/rocm/cu_inc/common.cuh"
#include "core/providers/rocm/rocm_common.h"
#include "core/providers/rocm/atomic/common.cuh"
#include "core/providers/rocm/reduction/reduction_utils.cuh"
namespace onnxruntime {
namespace rocm {
template<typename Tin, typename Tout>
__global__ void _ScalarSqrtImpl(Tin* input, Tout* output) {
*output = (Tout)_Sqrt(*input);
};
template<typename Tin, typename Tout>
void ScalarSqrt(Tin* input, Tout* output) {
hipLaunchKernelGGL(_ScalarSqrtImpl, dim3(1), dim3(1), 0, 0, input, output);
};
template void ScalarSqrt(float* input, float* output);
template void ScalarSqrt(half* input, half* output);
template void ScalarSqrt(float* input, half* output);
template <typename TIn, typename TOut, typename TBuf, typename TInOp, typename TOutOp>
__global__ void _MultiTensorReduceImpl(ChunkGroup<1> chunk_group, TOut* output) {
const int group_index = chunk_group.block_index_to_tensor_group_index[blockIdx.x];
const int tensor_size = chunk_group.tensor_sizes[group_index];
const int chunk_size = chunk_group.chunk_size;
const int chunk_start = chunk_group.block_index_to_chunk_start_index[blockIdx.x];
const TIn* w = reinterpret_cast<const TIn*>(chunk_group.tensor_ptrs[0][group_index]) + chunk_start;
TOut* w_norm = output;
TBuf w_sum = TBuf(0.f);
constexpr int load_count_per_thread = 4;
for (int i = threadIdx.x; i < chunk_size && i + chunk_start < tensor_size; i += blockDim.x * load_count_per_thread) {
#pragma unroll
for (int j = 0; j < load_count_per_thread; ++j) {
const int index_in_chunk = i + j * blockDim.x;
const int index_in_tensor = chunk_start + index_in_chunk;
if (index_in_chunk < chunk_size && index_in_tensor < tensor_size) {
const TBuf w_element = TBuf(w[index_in_chunk]);
w_sum += TInOp()(w_element);
}
}
}
// Thread count in a block must be a multiple of GPU_WARP_SIZE.
#pragma unroll
for (int stride = GPU_WARP_SIZE / 2; stride > 0; stride /= 2) {
w_sum += WARP_SHFL_DOWN(w_sum, stride);
}
const int warp_count_in_block = blockDim.x / GPU_WARP_SIZE;
const int lid = threadIdx.x % GPU_WARP_SIZE;
const int wid = threadIdx.x / GPU_WARP_SIZE;
// Shape is 2 x warp_count_in_block.
HIP_DYNAMIC_SHARED( unsigned char, shared_memory_)
TBuf* shared_memory = reinterpret_cast<TBuf*>(shared_memory_);
if (lid == 0) {
shared_memory[wid] = w_sum;
}
__syncthreads();
#pragma unroll
for (int stride = warp_count_in_block / 2; stride > 0; stride /= 2) {
if (threadIdx.x < stride) {
shared_memory[threadIdx.x] += shared_memory[threadIdx.x + stride];
}
__syncthreads();
}
if (threadIdx.x == 0) {
atomic_add(w_norm, TOutOp()(shared_memory[0]));
}
};
template <typename TIn, typename TOut, typename TBuf, typename TInOp, typename TOutOp>
void MultiTensorReduce(ChunkGroup<1> chunk_group, TOut* output) {
// thread count per block.
constexpr int thread_count = ChunkGroup<1>::thread_count_per_block;
// shared memory's size per block.
const int shared_memory_size = thread_count / GPU_WARP_SIZE * sizeof(TBuf);
// Enforce assumptions used inside this reduction ROCM kernel.
ORT_ENFORCE(thread_count % GPU_WARP_SIZE == 0);
ORT_ENFORCE((thread_count & (thread_count - 1)) == 0);
hipLaunchKernelGGL(HIP_KERNEL_NAME(_MultiTensorReduceImpl<TIn, TOut, TBuf, TInOp, TOutOp>), dim3(chunk_group.chunk_count), dim3(thread_count), shared_memory_size, 0, chunk_group, output);
}
template <typename TIn, typename TOut>
void MultiTensorReduceL2<TIn, TOut>::operator()(ChunkGroup<1> chunk_group, TOut* output) {
typedef typename ToBuffer<TIn>::Type TBuf;
MultiTensorReduce<TIn, TOut, TBuf, Square<TBuf, TIn>, Cast<TOut, TBuf>>(chunk_group, output);
}
#define INSTANTIATE_MULTI_TENSOR_REDUCTION_L2_FUNCTOR(TIn, TOut) \
template void MultiTensorReduceL2<TIn, TOut>::operator()(ChunkGroup<1> chunk_group, TOut* output);
INSTANTIATE_MULTI_TENSOR_REDUCTION_L2_FUNCTOR(double, float)
INSTANTIATE_MULTI_TENSOR_REDUCTION_L2_FUNCTOR(float, float)
INSTANTIATE_MULTI_TENSOR_REDUCTION_L2_FUNCTOR(half, float)
INSTANTIATE_MULTI_TENSOR_REDUCTION_L2_FUNCTOR(float, half)
INSTANTIATE_MULTI_TENSOR_REDUCTION_L2_FUNCTOR(half, half)
} // namespace rocm
} // namespace onnxruntime

7
orttraining/orttraining/training_ops/rocm/reduction/reduction_ops.cc
View file

@ -57,12 +57,7 @@ Status ReduceKernel<allow_multi_axes>::ComputeImplEx(OpKernelContext* ctx, miope
axes,
prepare_reduce_metadata));
Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims);
bool fast_reduction = fast_reduction_;
if (fast_reduction) {
auto ctx_internal = static_cast<OpKernelContextInternal*>(ctx);
if (ctx_internal && ctx_internal->GetUseDeterministicCompute())
fast_reduction = false;
}
const bool fast_reduction = fast_reduction_ && !ctx->GetUseDeterministicCompute();
return ReduceComputeCore<T>(*X, prepare_reduce_metadata, *Y, miopen_reduce_op, axes,
calculate_log_, calculate_sqt_, log_sum_exp_, fast_reduction);

3
tools/ci_build/amd_hipify.py
View file

@ -133,10 +133,12 @@ core_ops_files = [
'object_detection/roialign.h',
'object_detection/roialign_impl.cu',
'object_detection/roialign_impl.h',
'reduction/reduction_functions.cc',
'reduction/reduction_functions.cu',
'reduction/reduction_functions.h',
'reduction/reduction_ops.cc',
'reduction/reduction_ops.h',
'reduction/reduction_utils.cuh',
'rnn/cudnn_rnn_base.cc',
'rnn/cudnn_rnn_base.h',
'rnn/gru.cc',
@ -258,6 +260,7 @@ training_ops_files = [
'optimizer/adam.cu',
'optimizer/lamb.cc',
'reduction/reduction_all.cc',
'reduction/reduction_all.cu',
'reduction/reduction_ops.cc',
'tensor/gather_elements_grad.cc',
'tensor/gather_elements_grad.h',