mirror of
https://github.com/saymrwulf/onnxruntime.git
synced 2026-07-11 17:48:34 +00:00
Cuda Resize Operator for opset 11. (#2484)
* Cuda Resize Operator for opset 11.
This commit is contained in:
Zhang Lei
GitHub
parent
c42148a0c3
commit
50eb140119
10 changed files with 699 additions and 309 deletions
|
|
@ -29,6 +29,7 @@ ONNX_CPU_OPERATOR_VERSIONED_TYPED_KERNEL(
|
|||
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<uint8_t>()),
|
||||
Upsample<uint8_t>);
|
||||
|
||||
|
||||
template <typename T>
|
||||
void UpsampleNearest2x(int64_t batch_size,
|
||||
int64_t num_channels,
|
||||
|
|
@ -594,7 +595,7 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context,
|
|||
|
||||
if (dims.size() != scales.size())
|
||||
return Status(ONNXRUNTIME, INVALID_ARGUMENT,
|
||||
is_resize ? "Resize: input tensor's dimension does not match the scales."
|
||||
is_resize_ ? "Resize: input tensor's dimension does not match the scales."
|
||||
: "Upsample: input tensor's dimension does not match the scales.");
|
||||
|
||||
if (roi.size() != 2 * X->Shape().GetDims().size())
|
||||
|
|
@ -614,7 +615,7 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context,
|
|||
switch (mode_) {
|
||||
case UpsampleMode::NN:
|
||||
return UpsampleNearest<T>(X->template Data<T>(), Y->template MutableData<T>(), X->Shape(), Y->Shape(), scales, roi,
|
||||
is_resize, use_extrapolation_, extrapolation_value_, use_nearest2x_optimization,
|
||||
is_resize_, use_extrapolation_, extrapolation_value_, use_nearest2x_optimization_,
|
||||
get_original_coordinate_, get_nearest_pixel_);
|
||||
case UpsampleMode::LINEAR: {
|
||||
//The correct behavior of 'linear' mode for an N-D input is not clear right now,
|
||||
|
|
@ -623,7 +624,7 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context,
|
|||
std::ostringstream oss;
|
||||
oss << "'Linear' mode only support 2-D inputs ('Bilinear') or 4-D inputs "
|
||||
"with the corresponding outermost 2 scale values being 1 in the ";
|
||||
oss << (is_resize ? "Resize operator" : "Upsample operator");
|
||||
oss << (is_resize_ ? "Resize operator" : "Upsample operator");
|
||||
return Status(ONNXRUNTIME, FAIL, oss.str());
|
||||
}
|
||||
|
||||
|
|
@ -659,7 +660,7 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context,
|
|||
return Status::OK();
|
||||
}
|
||||
default:
|
||||
return Status(ONNXRUNTIME, FAIL, is_resize ? "Resize: unexpected mode" : "Upsample: unexpected mode");
|
||||
return Status(ONNXRUNTIME, FAIL, is_resize_ ? "Resize: unexpected mode" : "Upsample: unexpected mode");
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -25,13 +25,32 @@ enum UpsampleMode {
|
|||
CUBIC = 2, // cubic interpolation
|
||||
};
|
||||
|
||||
enum ResizeCoordinateTransformationMode {
|
||||
HALF_PIXEL = 0,
|
||||
ASYMMETRIC = 1,
|
||||
PYTORCH_HALF_PIXEL = 2,
|
||||
TF_HALF_PIXEL_FOR_NN = 3,
|
||||
ALIGN_CORNERS = 4,
|
||||
TF_CROP_AND_RESIZE = 5,
|
||||
CoordinateTransformationModeCount = 6,
|
||||
};
|
||||
|
||||
enum ResizeNearestMode {
|
||||
SIMPLE = 0, // For resize op 10
|
||||
ROUND_PREFER_FLOOR = 1,
|
||||
ROUND_PREFER_CEIL = 2,
|
||||
FLOOR = 3,
|
||||
CEIL = 4,
|
||||
NearestModeCount = 5,
|
||||
};
|
||||
|
||||
class UpsampleBase {
|
||||
protected:
|
||||
UpsampleBase(OpKernelInfo info) : scales_cached_(false), roi_cached_(false), use_extrapolation_(false) {
|
||||
int start;
|
||||
int end;
|
||||
info.GetKernelDef().SinceVersion(&start, &end);
|
||||
is_resize = (start >= 10);
|
||||
is_resize_ = (start >= 10);
|
||||
|
||||
std::string mode;
|
||||
ORT_ENFORCE(info.GetAttr<std::string>("mode", &mode).IsOK());
|
||||
|
|
@ -47,14 +66,18 @@ class UpsampleBase {
|
|||
extrapolation_value_ = info.GetAttrOrDefault<float>("extrapolation_value", 0.0f);
|
||||
|
||||
// Coordinate transformation mode attr was introduced in version 11, before that asymmetric mode was the only available transformation mode
|
||||
std::string coordinate_transform_mode = start > 10
|
||||
? info.GetAttrOrDefault<std::string>("coordinate_transformation_mode", "half_pixel")
|
||||
: "asymmetric";
|
||||
get_original_coordinate_ = GetOriginalCoordinateFromResizedCoordinate(coordinate_transform_mode);
|
||||
use_extrapolation_ = need_roi_input_ = coordinate_transform_mode == "tf_crop_and_resize" ? true : false;
|
||||
std::string coordinate_transform_mode_name = start > 10
|
||||
? info.GetAttrOrDefault<std::string>("coordinate_transformation_mode", "half_pixel")
|
||||
: "asymmetric";
|
||||
coordinate_transform_mode_ = StringToCoordinateTransformationMode(coordinate_transform_mode_name);
|
||||
get_original_coordinate_ = GetOriginalCoordinateFromResizedCoordinate(coordinate_transform_mode_);
|
||||
use_extrapolation_ = need_roi_input_ = (coordinate_transform_mode_ == TF_CROP_AND_RESIZE);
|
||||
|
||||
std::string nearest_mode = info.GetAttrOrDefault<std::string>("nearest_mode", "round_prefer_floor");
|
||||
get_nearest_pixel_ = GetNearestPixelFromOriginal(nearest_mode, start);
|
||||
std::string nearest_mode_name = (mode_ == NN && start >= 11)
|
||||
? info.GetAttrOrDefault<std::string>("nearest_mode", "round_prefer_floor")
|
||||
: "";
|
||||
nearest_mode_ = StringToNearstMode(nearest_mode_name);
|
||||
get_nearest_pixel_ = GetNearestPixelFromOriginal(nearest_mode_);
|
||||
|
||||
cubic_coeff_a_ = info.GetAttrOrDefault<float>("cubic_coeff_a", -0.75f);
|
||||
exclude_outside_ = info.GetAttrOrDefault<int64_t>("exclude_outside", 0) == 0 ? false : true;
|
||||
|
|
@ -65,7 +88,7 @@ class UpsampleBase {
|
|||
|
||||
// after version 11 update, this optimization is no longer applicable for all the available modes...
|
||||
// TODO : needs more testing to enable this for version 11
|
||||
use_nearest2x_optimization = start > 10 ? false : true;
|
||||
use_nearest2x_optimization_ = start > 10 ? false : true;
|
||||
|
||||
if (start > 10) {
|
||||
roi_input_idx_ = 1;
|
||||
|
|
@ -98,13 +121,15 @@ class UpsampleBase {
|
|||
}
|
||||
}
|
||||
|
||||
UpsampleMode mode_;
|
||||
ResizeCoordinateTransformationMode coordinate_transform_mode_;
|
||||
GetOriginalCoordinateFunc get_original_coordinate_;
|
||||
ResizeNearestMode nearest_mode_;
|
||||
GetNearestPixelFunc get_nearest_pixel_;
|
||||
float cubic_coeff_a_;
|
||||
bool exclude_outside_;
|
||||
float extrapolation_value_;
|
||||
UpsampleMode mode_;
|
||||
bool use_nearest2x_optimization = false;
|
||||
bool use_nearest2x_optimization_ = false;
|
||||
|
||||
std::vector<float> scales_;
|
||||
std::vector<float> roi_;
|
||||
|
|
@ -112,7 +137,7 @@ class UpsampleBase {
|
|||
bool roi_cached_;
|
||||
bool need_roi_input_;
|
||||
bool use_extrapolation_;
|
||||
bool is_resize = false;
|
||||
bool is_resize_ = false;
|
||||
|
||||
int roi_input_idx_ = -1;
|
||||
int scales_input_idx_ = -1;
|
||||
|
|
@ -125,7 +150,6 @@ class UpsampleBase {
|
|||
if (mode == UpsampleModeLinear) {
|
||||
return UpsampleMode::LINEAR;
|
||||
}
|
||||
|
||||
if (mode == UpsampleModeCubic) {
|
||||
return UpsampleMode::CUBIC;
|
||||
}
|
||||
|
|
@ -133,78 +157,114 @@ class UpsampleBase {
|
|||
UpsampleModeNN + "(default) or " + UpsampleModeLinear + " or " + UpsampleModeCubic + ".");
|
||||
}
|
||||
|
||||
ResizeCoordinateTransformationMode StringToCoordinateTransformationMode(
|
||||
const std::string& coordinate_transform_mode_name) {
|
||||
if (coordinate_transform_mode_name == "asymmetric") {
|
||||
return ASYMMETRIC;
|
||||
}
|
||||
if (coordinate_transform_mode_name == "pytorch_half_pixel") {
|
||||
return PYTORCH_HALF_PIXEL;
|
||||
}
|
||||
if (coordinate_transform_mode_name == "tf_half_pixel_for_nn") {
|
||||
return TF_HALF_PIXEL_FOR_NN;
|
||||
}
|
||||
if (coordinate_transform_mode_name == "align_corners") {
|
||||
return ALIGN_CORNERS;
|
||||
}
|
||||
if (coordinate_transform_mode_name == "tf_crop_and_resize") {
|
||||
return TF_CROP_AND_RESIZE;
|
||||
}
|
||||
if (coordinate_transform_mode_name == "half_pixel") {
|
||||
return HALF_PIXEL;
|
||||
}
|
||||
ORT_THROW("coordinate_transform_mode:[" + coordinate_transform_mode_name + "] is not supportted!");
|
||||
}
|
||||
|
||||
GetOriginalCoordinateFunc GetOriginalCoordinateFromResizedCoordinate(
|
||||
const std::string& coordinate_transform_mode) {
|
||||
if (coordinate_transform_mode == "asymmetric") {
|
||||
return [](float x_resized, float x_scale, float, float, float, float) {
|
||||
return x_resized / x_scale;
|
||||
};
|
||||
} else if (coordinate_transform_mode == "pytorch_half_pixel") {
|
||||
return [](float x_resized, float x_scale, float length_resized, float, float, float) {
|
||||
return length_resized > 1 ? (x_resized + 0.5f) / x_scale - 0.5f : 0.0f;
|
||||
};
|
||||
} else if (coordinate_transform_mode == "tf_half_pixel_for_nn") {
|
||||
return [](float x_resized, float x_scale, float, float, float, float) {
|
||||
return (x_resized + 0.5f) / x_scale;
|
||||
};
|
||||
} else if (coordinate_transform_mode == "align_corners") {
|
||||
return [](float x_resized, float, float length_resized, float length_original, float, float) {
|
||||
return length_resized == 1 ? 0 : x_resized * (length_original - 1) / (length_resized - 1);
|
||||
};
|
||||
} else if (coordinate_transform_mode == "tf_crop_and_resize") {
|
||||
return [](float x_resized, float, float length_resized, float length_original, float roi_start, float roi_end) {
|
||||
auto orig = length_resized > 1
|
||||
? roi_start * (length_original - 1) + (x_resized * (roi_end - roi_start) * (length_original - 1)) / (length_resized - 1)
|
||||
: 0.5 * (roi_start + roi_end) * (length_original - 1);
|
||||
return static_cast<float>(orig);
|
||||
};
|
||||
} else { // "half_pixel"
|
||||
return [](float x_resized, float x_scale, float, float, float, float) {
|
||||
return ((x_resized + 0.5f) / x_scale) - 0.5f;
|
||||
};
|
||||
ResizeCoordinateTransformationMode coordinate_transform_mode) {
|
||||
switch (coordinate_transform_mode) {
|
||||
case ASYMMETRIC:
|
||||
return [](float x_resized, float x_scale, float, float, float, float) {
|
||||
return x_resized / x_scale;
|
||||
};
|
||||
case PYTORCH_HALF_PIXEL:
|
||||
return [](float x_resized, float x_scale, float length_resized, float, float, float) {
|
||||
return length_resized > 1 ? (x_resized + 0.5f) / x_scale - 0.5f : 0.0f;
|
||||
};
|
||||
case TF_HALF_PIXEL_FOR_NN:
|
||||
return [](float x_resized, float x_scale, float, float, float, float) {
|
||||
return (x_resized + 0.5f) / x_scale;
|
||||
};
|
||||
case ALIGN_CORNERS:
|
||||
return [](float x_resized, float, float length_resized, float length_original, float, float) {
|
||||
return length_resized == 1 ? 0 : x_resized * (length_original - 1) / (length_resized - 1);
|
||||
};
|
||||
case TF_CROP_AND_RESIZE:
|
||||
return [](float x_resized, float, float length_resized, float length_original, float roi_start, float roi_end) {
|
||||
auto orig = length_resized > 1
|
||||
? roi_start * (length_original - 1) + (x_resized * (roi_end - roi_start) * (length_original - 1)) / (length_resized - 1)
|
||||
: 0.5 * (roi_start + roi_end) * (length_original - 1);
|
||||
return static_cast<float>(orig);
|
||||
};
|
||||
default: // "half_pixel"
|
||||
return [](float x_resized, float x_scale, float, float, float, float) {
|
||||
return ((x_resized + 0.5f) / x_scale) - 0.5f;
|
||||
};
|
||||
}
|
||||
}
|
||||
|
||||
GetNearestPixelFunc GetNearestPixelFromOriginal(
|
||||
const std::string& nearest_mode, int opset_version) {
|
||||
// versions older than 11 did not have nearest_mode attr. Use the original logic in this case
|
||||
// to maintain backward compatibility
|
||||
if (opset_version < 11) {
|
||||
return [](float x_original, bool isDownSample) {
|
||||
if (isDownSample) {
|
||||
return static_cast<int64_t>(std::ceil(x_original));
|
||||
} else {
|
||||
return static_cast<int64_t>(x_original);
|
||||
}
|
||||
};
|
||||
ResizeNearestMode StringToNearstMode(const std::string& nearst_mode_name) {
|
||||
if (nearst_mode_name == "round_prefer_floor") {
|
||||
return ROUND_PREFER_FLOOR;
|
||||
} else if (nearst_mode_name == "round_prefer_ceil") {
|
||||
return ROUND_PREFER_CEIL;
|
||||
} else if (nearst_mode_name == "floor") {
|
||||
return FLOOR;
|
||||
} else if (nearst_mode_name == "ceil") {
|
||||
return CEIL;
|
||||
} else if (nearst_mode_name == "") {
|
||||
return SIMPLE;
|
||||
}
|
||||
ORT_THROW("nearst_mode:[" + nearst_mode_name + "] is not supportted!");
|
||||
}
|
||||
|
||||
// if opset version >=11 choose the rounding mode based on nearest_mode attr
|
||||
if (nearest_mode == "round_prefer_ceil") {
|
||||
return [](float x_original, bool) {
|
||||
return static_cast<int64_t>(std::round(x_original));
|
||||
};
|
||||
} else if (nearest_mode == "floor") {
|
||||
return [](float x_original, bool) {
|
||||
return static_cast<int64_t>(std::floor(x_original));
|
||||
};
|
||||
} else if (nearest_mode == "ceil") {
|
||||
return [](float x_original, bool) {
|
||||
return static_cast<int64_t>(std::ceil(x_original));
|
||||
};
|
||||
} else { // default is round_prefer_floor
|
||||
return [](float x_original, bool) {
|
||||
// for half way cases prefer floor
|
||||
if (x_original == static_cast<int64_t>(x_original) + 0.5f) {
|
||||
GetNearestPixelFunc GetNearestPixelFromOriginal(ResizeNearestMode nearest_mode) {
|
||||
switch (nearest_mode) {
|
||||
case SIMPLE:
|
||||
// versions older than 11 did not have nearest_mode attr. Use the original logic in this case
|
||||
// to maintain backward compatibility
|
||||
return [](float x_original, bool isDownSample) {
|
||||
if (isDownSample) {
|
||||
return static_cast<int64_t>(std::ceil(x_original));
|
||||
} else {
|
||||
return static_cast<int64_t>(x_original);
|
||||
}
|
||||
};
|
||||
case ROUND_PREFER_CEIL:
|
||||
return [](float x_original, bool) {
|
||||
return static_cast<int64_t>(std::round(x_original));
|
||||
};
|
||||
case FLOOR:
|
||||
return [](float x_original, bool) {
|
||||
return static_cast<int64_t>(std::floor(x_original));
|
||||
}
|
||||
return static_cast<int64_t>(std::round(x_original));
|
||||
};
|
||||
};
|
||||
case CEIL:
|
||||
return [](float x_original, bool) {
|
||||
return static_cast<int64_t>(std::ceil(x_original));
|
||||
};
|
||||
default: // default is round_prefer_floor
|
||||
return [](float x_original, bool) {
|
||||
// for half way cases prefer floor
|
||||
if (x_original == static_cast<int64_t>(x_original) + 0.5f) {
|
||||
return static_cast<int64_t>(std::floor(x_original));
|
||||
}
|
||||
return static_cast<int64_t>(std::round(x_original));
|
||||
};
|
||||
}
|
||||
}
|
||||
|
||||
void ScalesValidation(const std::vector<float>& scales, const UpsampleMode mode) const {
|
||||
if (!is_resize) {
|
||||
void ScalesValidation(const std::vector<float>& scales, const UpsampleMode mode) const {
|
||||
if (!is_resize_) {
|
||||
for (auto& scale : scales) {
|
||||
ORT_ENFORCE(scale >= 1, "Scale value should be greater than or equal to 1.");
|
||||
}
|
||||
|
|
@ -218,7 +278,7 @@ class UpsampleBase {
|
|||
ORT_ENFORCE(scales.size() == 2 || (scales.size() == 4 && scales[0] == 1 && scales[1] == 1),
|
||||
"'Linear' mode and 'Cubic' mode only support 2-D inputs ('Bilinear', 'Bicubic') or 4-D inputs "
|
||||
"with the corresponding outermost 2 scale values being 1 in the ",
|
||||
is_resize ? "Resize operator" : "Upsample operator");
|
||||
is_resize_ ? "Resize operator" : "Upsample operator");
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -257,7 +317,7 @@ class UpsampleBase {
|
|||
output_dims[i] = static_cast<int64_t>(scales[i] * input_dims[i]);
|
||||
}
|
||||
}
|
||||
};
|
||||
}; // UpsampleBase
|
||||
|
||||
template <typename T>
|
||||
class Upsample : public UpsampleBase, public OpKernel {
|
||||
|
|
|
|||
|
|
@ -557,11 +557,11 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kO
|
|||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, double, MaxPool);
|
||||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, MaxPool);
|
||||
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, NonMaxSuppression);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int32_t, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, Resize);
|
||||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, float, Resize);
|
||||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, double, Resize);
|
||||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, Resize);
|
||||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, int32_t, Resize);
|
||||
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, uint8_t, Resize);
|
||||
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, ReverseSequence);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, RoiAlign);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, RoiAlign);
|
||||
|
|
@ -652,6 +652,11 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
|
|||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, MaxPool);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, MaxPool);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, MaxPool);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, bool, Equal);
|
||||
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Equal);
|
||||
|
|
@ -1007,11 +1012,11 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
|
|||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, double, MaxPool)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, MaxPool)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, NonMaxSuppression)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int32_t, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, float, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, double, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, MLFloat16, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, int32_t, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, uint8_t, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, ReverseSequence)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, RoiAlign)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, RoiAlign)>,
|
||||
|
|
@ -1103,6 +1108,11 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
|
|||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, MaxPool)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, MaxPool)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, MaxPool)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, bool, Equal)>,
|
||||
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Equal)>,
|
||||
|
|
|
|||
|
|
@ -6,15 +6,27 @@
|
|||
namespace onnxruntime {
|
||||
namespace cuda {
|
||||
#define REGISTER_KERNEL_TYPED(T) \
|
||||
ONNX_OPERATOR_TYPED_KERNEL_EX( \
|
||||
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
|
||||
Resize, \
|
||||
kOnnxDomain, \
|
||||
10, \
|
||||
10, 10, \
|
||||
T, \
|
||||
kCudaExecutionProvider, \
|
||||
KernelDefBuilder() \
|
||||
.InputMemoryType<OrtMemTypeCPUInput>(1) \
|
||||
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
|
||||
Resize<T>); \
|
||||
ONNX_OPERATOR_TYPED_KERNEL_EX( \
|
||||
Resize, \
|
||||
kOnnxDomain, \
|
||||
11, \
|
||||
T, \
|
||||
kCudaExecutionProvider, \
|
||||
KernelDefBuilder() \
|
||||
.InputMemoryType<OrtMemTypeCPUInput>(1) \
|
||||
.InputMemoryType<OrtMemTypeCPUInput>(2) \
|
||||
.InputMemoryType<OrtMemTypeCPUInput>(3) \
|
||||
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
|
||||
Resize<T>);
|
||||
|
||||
REGISTER_KERNEL_TYPED(float)
|
||||
|
|
|
|||
|
|
@ -3,147 +3,245 @@
|
|||
|
||||
namespace onnxruntime {
|
||||
namespace cuda {
|
||||
|
||||
using onnxruntime::ResizeCoordinateTransformationMode;
|
||||
using onnxruntime::ResizeNearestMode;
|
||||
using onnxruntime::UpsampleMode;
|
||||
|
||||
__device__ int NearestPixel_SIMPLE(float x_original, bool is_down_sampling) {
|
||||
if (is_down_sampling) {
|
||||
return static_cast<int>(ceil(x_original));
|
||||
} else {
|
||||
return static_cast<int>(x_original);
|
||||
}
|
||||
}
|
||||
|
||||
__device__ int NearestPixel_ROUND_PREFER_FLOOR(float x_original, bool) {
|
||||
if (x_original == static_cast<int>(x_original) + 0.5f) {
|
||||
return static_cast<int>(floor(x_original));
|
||||
}
|
||||
return static_cast<int>(round(x_original));
|
||||
}
|
||||
|
||||
__device__ int NearestPixel_ROUND_PREFER_CEIL(float x_original, bool) {
|
||||
return static_cast<int>(round(x_original));
|
||||
}
|
||||
|
||||
__device__ int NearestPixel_FLOOR(float x_original, bool) {
|
||||
return static_cast<int>(floor(x_original));
|
||||
}
|
||||
|
||||
__device__ int NearestPixel_CEIL(float x_original, bool) {
|
||||
return static_cast<int>(ceil(x_original));
|
||||
}
|
||||
|
||||
using CudaFunctionNearestPixel = int (*)(float, bool);
|
||||
__device__ CudaFunctionNearestPixel func_NearestPixel_SIMPLE = NearestPixel_SIMPLE;
|
||||
__device__ CudaFunctionNearestPixel func_NearestPixel_ROUND_PREFER_FLOOR = NearestPixel_ROUND_PREFER_FLOOR;
|
||||
__device__ CudaFunctionNearestPixel func_NearestPixel_ROUND_PREFER_CEIL = NearestPixel_ROUND_PREFER_CEIL;
|
||||
__device__ CudaFunctionNearestPixel func_NearestPixel_FLOOR = NearestPixel_FLOOR;
|
||||
__device__ CudaFunctionNearestPixel func_NearestPixel_CEIL = NearestPixel_CEIL;
|
||||
|
||||
CudaFunctionNearestPixel GetDeviceNearstPixelFunction(ResizeNearestMode nearest_mode) {
|
||||
static bool already_copied = false;
|
||||
static std::mutex s_mutext;
|
||||
static CudaFunctionNearestPixel s_nearest_pixel[ResizeNearestMode::NearestModeCount];
|
||||
if (!already_copied) {
|
||||
std::lock_guard<std::mutex> lock(s_mutext);
|
||||
if (!already_copied) {
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::SIMPLE],
|
||||
func_NearestPixel_SIMPLE, sizeof(CudaFunctionNearestPixel)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::ROUND_PREFER_FLOOR],
|
||||
func_NearestPixel_ROUND_PREFER_FLOOR, sizeof(CudaFunctionNearestPixel)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::ROUND_PREFER_CEIL],
|
||||
func_NearestPixel_ROUND_PREFER_CEIL, sizeof(CudaFunctionNearestPixel)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::FLOOR],
|
||||
func_NearestPixel_FLOOR, sizeof(CudaFunctionNearestPixel)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_nearest_pixel[ResizeNearestMode::CEIL],
|
||||
func_NearestPixel_CEIL, sizeof(CudaFunctionNearestPixel)));
|
||||
already_copied = true;
|
||||
}
|
||||
}
|
||||
return s_nearest_pixel[nearest_mode];
|
||||
}
|
||||
|
||||
__device__ float TransformCoordinate_ASYMMETRIC(float x_resized, float x_scale, float, float, float, float) {
|
||||
return x_resized / x_scale;
|
||||
}
|
||||
|
||||
__device__ float TransformCoordinate_HALF_PIXEL(float x_resized, float x_scale, float, float, float, float) {
|
||||
return ((x_resized + 0.5f) / x_scale) - 0.5f;
|
||||
}
|
||||
|
||||
__device__ float TransformCoordinate_PYTORCH_HALF_PIXEL(
|
||||
float x_resized, float x_scale, float length_resized, float, float, float) {
|
||||
return length_resized > 1 ? (x_resized + 0.5f) / x_scale - 0.5f : 0.0f;
|
||||
}
|
||||
|
||||
__device__ float TransformCoordinate_TF_HALF_PIXEL_FOR_NN(
|
||||
float x_resized, float x_scale, float, float, float, float) {
|
||||
return (x_resized + 0.5f) / x_scale;
|
||||
}
|
||||
|
||||
__device__ float TransformCoordinate_ALIGN_CORNERS(
|
||||
float x_resized, float, float length_resized, float length_original, float, float) {
|
||||
return length_resized == 1 ? 0 : x_resized * (length_original - 1) / (length_resized - 1);
|
||||
}
|
||||
|
||||
__device__ float TransformCoordinate_TF_CROP_AND_RESIZE(
|
||||
float x_resized, float, float length_resized, float length_original, float roi_start, float roi_end) {
|
||||
auto orig = length_resized > 1
|
||||
? roi_start * (length_original - 1) + (x_resized * (roi_end - roi_start) * (length_original - 1)) / (length_resized - 1)
|
||||
: 0.5 * (roi_start + roi_end) * (length_original - 1);
|
||||
return static_cast<float>(orig);
|
||||
}
|
||||
|
||||
using CudaFunctionOriginalCoordinate = float (*)(float, float, float, float, float, float);
|
||||
|
||||
__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_ASYMMETRIC = TransformCoordinate_ASYMMETRIC;
|
||||
__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_HALF_PIXEL = TransformCoordinate_HALF_PIXEL;
|
||||
__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_PYTORCH_HALF_PIXEL = TransformCoordinate_PYTORCH_HALF_PIXEL;
|
||||
__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_ALIGN_CORNERS = TransformCoordinate_ALIGN_CORNERS;
|
||||
__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_TF_HALF_PIXEL_FOR_NN = TransformCoordinate_TF_HALF_PIXEL_FOR_NN;
|
||||
__device__ CudaFunctionOriginalCoordinate func_TransformCoordinate_TF_CROP_AND_RESIZE = TransformCoordinate_TF_CROP_AND_RESIZE;
|
||||
|
||||
CudaFunctionOriginalCoordinate GetDeviceOriginalCoordinateFunc(ResizeCoordinateTransformationMode coordinate_transform_mode) {
|
||||
static bool already_copied = false;
|
||||
static std::mutex s_mutext;
|
||||
static CudaFunctionOriginalCoordinate s_coordinate_tranforms[ResizeCoordinateTransformationMode::CoordinateTransformationModeCount];
|
||||
if (!already_copied) {
|
||||
std::lock_guard<std::mutex> lock(s_mutext);
|
||||
if (!already_copied) {
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::HALF_PIXEL],
|
||||
func_TransformCoordinate_HALF_PIXEL, sizeof(CudaFunctionOriginalCoordinate)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::ASYMMETRIC],
|
||||
func_TransformCoordinate_ASYMMETRIC, sizeof(CudaFunctionOriginalCoordinate)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::PYTORCH_HALF_PIXEL],
|
||||
func_TransformCoordinate_PYTORCH_HALF_PIXEL, sizeof(CudaFunctionOriginalCoordinate)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::ALIGN_CORNERS],
|
||||
func_TransformCoordinate_ALIGN_CORNERS, sizeof(CudaFunctionOriginalCoordinate)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::TF_HALF_PIXEL_FOR_NN],
|
||||
func_TransformCoordinate_TF_HALF_PIXEL_FOR_NN, sizeof(CudaFunctionOriginalCoordinate)));
|
||||
CUDA_CALL(cudaMemcpyFromSymbol(&s_coordinate_tranforms[ResizeCoordinateTransformationMode::TF_CROP_AND_RESIZE],
|
||||
func_TransformCoordinate_TF_CROP_AND_RESIZE, sizeof(CudaFunctionOriginalCoordinate)));
|
||||
already_copied = true;
|
||||
}
|
||||
}
|
||||
return s_coordinate_tranforms[coordinate_transform_mode];
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void _ResizeNearestKernel(const size_t rank,
|
||||
const int64_t* input_pitches,
|
||||
const fast_divmod* output_div_pitches,
|
||||
const float* scales,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N) {
|
||||
__global__ void _ResizeNearestKernel(
|
||||
const size_t rank,
|
||||
const int64_t* input_shape,
|
||||
const int64_t* output_shape,
|
||||
const int64_t* input_pitches,
|
||||
const fast_divmod* output_div_pitches,
|
||||
const float* scales,
|
||||
const float* roi,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N,
|
||||
bool extrapolation_enabled,
|
||||
float extrapolation_value,
|
||||
CudaFunctionOriginalCoordinate transform_coordinate,
|
||||
CudaFunctionNearestPixel calc_nearest_pixel) {
|
||||
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
|
||||
CUDA_LONG input_index = 0;
|
||||
CUDA_LONG output_index = id;
|
||||
|
||||
int div, mod;
|
||||
bool extrapolation_occured = false;
|
||||
for (int dim = 0; dim < rank; ++dim) {
|
||||
output_div_pitches[dim].divmod(output_index, div, mod);
|
||||
output_index = mod;
|
||||
if (scales[dim] <= 1) { //downsample
|
||||
div = std::ceil(div / scales[dim]);
|
||||
} else { //upsample
|
||||
div = div / scales[dim];
|
||||
float orig_coord = transform_coordinate(static_cast<float>(div), scales[dim], static_cast<float>(output_shape[dim]),
|
||||
static_cast<float>(input_shape[dim]), roi[dim], roi[dim + rank]);
|
||||
if (extrapolation_enabled && !extrapolation_occured) {
|
||||
extrapolation_occured = (orig_coord < 0.f || orig_coord > static_cast<float>(input_shape[dim] - 1));
|
||||
}
|
||||
div = calc_nearest_pixel(orig_coord, scales[dim] < 1);
|
||||
if (div >= input_shape[dim]) div = input_shape[dim] - 1;
|
||||
if (div < 0) div = 0;
|
||||
input_index += input_pitches[dim] * div;
|
||||
}
|
||||
output_data[id] = input_data[input_index];
|
||||
output_data[id] = extrapolation_occured ? static_cast<T>(extrapolation_value) : input_data[input_index];
|
||||
}
|
||||
|
||||
// The following method supports a 4-D input in 'Linear mode'
|
||||
// that amounts to 'Bilinear' Upsampling/Resizing in the sense that it assumes
|
||||
// the scale values for the outermost 2 dimensions are 1.
|
||||
// This is the common use-case where the 4-D input (batched multi-channel images)
|
||||
// is usually of shape [N, C, H, W] and the scales are [1.0, 1.0, height_scale, width_scale]
|
||||
struct BilinearMappingInfo {
|
||||
int origin_;
|
||||
float weight_;
|
||||
int extrapolate_;
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
__global__ void _ResizeBilinear4DInputKernel(const int64_t input_dim2,
|
||||
const int64_t* input_pitches,
|
||||
const fast_divmod* output_div_pitches,
|
||||
const float* scales,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N) {
|
||||
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
|
||||
CUDA_LONG input_index = 0;
|
||||
|
||||
// For bilinear mode, scales[0]=scales[1]=1
|
||||
int mod;
|
||||
int index_of_dim0, index_of_dim1, index_of_dim2, index_of_dim3;
|
||||
output_div_pitches[0].divmod(id, index_of_dim0, mod);
|
||||
output_div_pitches[1].divmod(mod, index_of_dim1, mod);
|
||||
output_div_pitches[2].divmod(mod, index_of_dim2, mod);
|
||||
index_of_dim3 = mod;
|
||||
int index_of_input_dim2, index_of_input_dim3;
|
||||
float x_offset_0, y_offset_0, x_offset_1, y_offset_1;
|
||||
index_of_input_dim2 = static_cast<int64_t>(index_of_dim2 / scales[2]);
|
||||
index_of_input_dim3 = static_cast<int64_t>(index_of_dim3 / scales[3]);
|
||||
input_index = index_of_dim0 * input_pitches[0] +
|
||||
index_of_dim1 * input_pitches[1] +
|
||||
index_of_input_dim2 * input_pitches[2] +
|
||||
index_of_input_dim3;
|
||||
|
||||
T x00 = input_data[input_index];
|
||||
T x10, x01, x11;
|
||||
|
||||
bool end_of_dim2 = false, end_of_dim3 = false;
|
||||
if (index_of_input_dim2 == (input_dim2 - 1)) {
|
||||
// It's the end in dimension 2
|
||||
x01 = x00;
|
||||
end_of_dim2 = true;
|
||||
} else {
|
||||
x01 = input_data[input_index + input_pitches[2]];
|
||||
__global__ void _ResizeBilinearCoordinateMapping(
|
||||
int64_t input_height, int64_t input_width,
|
||||
int64_t output_height, int64_t output_width,
|
||||
float scale_height, float scale_width,
|
||||
float roi_height_start, float roi_height_end,
|
||||
float roi_width_start, float roi_width_end,
|
||||
const size_t SumHW, bool extrapolation_enabled,
|
||||
CudaFunctionOriginalCoordinate transform_coordinate,
|
||||
BilinearMappingInfo* dims_mapping) {
|
||||
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, SumHW);
|
||||
if (id < output_height) { // y = id
|
||||
float input_y = transform_coordinate(static_cast<float>(id), scale_height,
|
||||
static_cast<float>(output_height), static_cast<float>(input_height),
|
||||
roi_height_start, roi_height_end);
|
||||
dims_mapping[id].extrapolate_ = (int)(extrapolation_enabled && (input_y < 0 || input_y > static_cast<float>(input_height - 1)));
|
||||
input_y = max(0.0f, min(input_y, static_cast<float>(input_height - 1)));
|
||||
int y_int = static_cast<int>(input_y);
|
||||
dims_mapping[id].origin_ = y_int;
|
||||
dims_mapping[id].weight_ = (y_int >= input_height - 1) ? 0.5f : input_y - y_int;
|
||||
} else { //x = id - output_height
|
||||
float input_x = transform_coordinate(static_cast<float>(id - output_height), scale_width,
|
||||
static_cast<float>(output_width), static_cast<float>(input_width),
|
||||
roi_width_start, roi_width_end);
|
||||
dims_mapping[id].extrapolate_ = (int)(extrapolation_enabled && (input_x < 0 || input_x > static_cast<float>(input_width - 1)));
|
||||
input_x = max(0.0f, min(input_x, static_cast<float>(input_width - 1)));
|
||||
int x_int = static_cast<int>(input_x);
|
||||
dims_mapping[id].origin_ = x_int;
|
||||
dims_mapping[id].weight_ = (x_int >= input_width - 1) ? 0.5f : input_x - x_int;
|
||||
}
|
||||
|
||||
if (index_of_input_dim3 == (input_pitches[2] - 1)) {
|
||||
// It's the end in dimension 3
|
||||
x10 = x00;
|
||||
x11 = x01;
|
||||
end_of_dim3 = true;
|
||||
} else {
|
||||
x10 = input_data[input_index + 1];
|
||||
x11 = end_of_dim2 ? x10 : input_data[input_index + input_pitches[2] + 1];
|
||||
}
|
||||
|
||||
y_offset_0 = end_of_dim2 ? 0.5f : index_of_dim2 / scales[2] - index_of_input_dim2;
|
||||
y_offset_1 = 1.0f - y_offset_0;
|
||||
x_offset_0 = end_of_dim3 ? 0.5f : index_of_dim3 / scales[3] - index_of_input_dim3;
|
||||
x_offset_1 = 1.0f - x_offset_0;
|
||||
|
||||
output_data[id] =
|
||||
x00 * static_cast<T>(y_offset_1 * x_offset_1) +
|
||||
x01 * static_cast<T>(y_offset_0 * x_offset_1) +
|
||||
x10 * static_cast<T>(y_offset_1 * x_offset_0) +
|
||||
x11 * static_cast<T>(y_offset_0 * x_offset_0);
|
||||
}
|
||||
|
||||
// The following method supports a 2-D input in 'Linear mode'
|
||||
// The following method supports a N-D input in 'Linear mode'. Last two dimension is [H, W].
|
||||
// the scale values for the outer dimensions except last two are 1.
|
||||
template <typename T>
|
||||
__global__ void _ResizeBilinear2DInputKernel(const int64_t input_dim0,
|
||||
const int64_t* input_pitches,
|
||||
const fast_divmod* output_div_pitches,
|
||||
const float* scales,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N) {
|
||||
__global__ void _ResizeBilinearKernel(
|
||||
int64_t input_height, int64_t input_width,
|
||||
int64_t output_height, int64_t output_width,
|
||||
fast_divmod div_output_width, fast_divmod div_output_image,
|
||||
const T* input_data, T* output_data, const size_t N,
|
||||
float extrapolation_value,
|
||||
BilinearMappingInfo* dims_mapping) {
|
||||
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
|
||||
CUDA_LONG input_index = 0;
|
||||
int bxc, output_image_index;
|
||||
div_output_image.divmod(id, bxc, output_image_index);
|
||||
CUDA_LONG input_index = bxc * input_height * input_width;
|
||||
int output_y, output_x;
|
||||
div_output_width.divmod(output_image_index, output_y, output_x);
|
||||
|
||||
int mod;
|
||||
int index_of_dim0, index_of_dim1;
|
||||
output_div_pitches[0].divmod(id, index_of_dim0, mod);
|
||||
index_of_dim1 = mod;
|
||||
int index_of_input_dim0, index_of_input_dim1;
|
||||
float x_offset_0, y_offset_0, x_offset_1, y_offset_1;
|
||||
index_of_input_dim0 = static_cast<int64_t>(index_of_dim0 / scales[0]);
|
||||
index_of_input_dim1 = static_cast<int64_t>(index_of_dim1 / scales[1]);
|
||||
input_index = index_of_input_dim0 * input_pitches[0] + index_of_input_dim1;
|
||||
if (dims_mapping[output_y].extrapolate_ || dims_mapping[output_x + output_height].extrapolate_) {
|
||||
output_data[id] = extrapolation_value;
|
||||
return;
|
||||
}
|
||||
float y_offset_0 = dims_mapping[output_y].weight_;
|
||||
int y_int = dims_mapping[output_y].origin_;
|
||||
float x_offset_0 = dims_mapping[output_x + output_height].weight_;
|
||||
int x_int = dims_mapping[output_x + output_height].origin_;
|
||||
input_index += y_int * input_width + x_int;
|
||||
|
||||
T x00 = input_data[input_index];
|
||||
T x10, x01, x11;
|
||||
|
||||
bool end_of_dim0 = false, end_of_dim1 = false;
|
||||
if (index_of_input_dim0 == (input_dim0 - 1)) {
|
||||
// It's the end in dimension 0
|
||||
x01 = x00;
|
||||
end_of_dim0 = true;
|
||||
} else {
|
||||
x01 = input_data[input_index + input_pitches[0]];
|
||||
}
|
||||
|
||||
if (index_of_input_dim1 == (input_pitches[0] - 1)) {
|
||||
// It's the end in dimension 1
|
||||
x10 = x00;
|
||||
x11 = x01;
|
||||
end_of_dim1 = true;
|
||||
} else {
|
||||
x10 = input_data[input_index + 1];
|
||||
x11 = end_of_dim0 ? x10 : input_data[input_index + input_pitches[0] + 1];
|
||||
}
|
||||
|
||||
y_offset_0 = end_of_dim0 ? 0.5f : index_of_dim0 / scales[0] - index_of_input_dim0;
|
||||
y_offset_1 = 1.0f - y_offset_0;
|
||||
x_offset_0 = end_of_dim1 ? 0.5f : index_of_dim1 / scales[1] - index_of_input_dim1;
|
||||
x_offset_1 = 1.0f - x_offset_0;
|
||||
bool end_of_h = (y_int >= input_height - 1);
|
||||
bool end_of_w = (x_int >= input_width - 1);
|
||||
T x10 = end_of_w ? x00 : input_data[input_index + 1];
|
||||
T x01 = end_of_h ? x00 : input_data[input_index + input_width];
|
||||
T x11 = end_of_w ? x01 : (end_of_h ? x10 : input_data[input_index + input_width + 1]);
|
||||
|
||||
float y_offset_1 = 1.0f - y_offset_0;
|
||||
float x_offset_1 = 1.0f - x_offset_0;
|
||||
output_data[id] =
|
||||
x00 * static_cast<T>(y_offset_1 * x_offset_1) +
|
||||
x01 * static_cast<T>(y_offset_0 * x_offset_1) +
|
||||
|
|
@ -152,41 +250,216 @@ __global__ void _ResizeBilinear2DInputKernel(const int64_t input_dim0,
|
|||
}
|
||||
|
||||
template <typename T>
|
||||
void ResizeImpl(const onnxruntime::UpsampleMode upsample_mode,
|
||||
const size_t rank,
|
||||
const int64_t input_dim2,
|
||||
const int64_t* input_pitches,
|
||||
const fast_divmod* output_div_pitches,
|
||||
const float* scales_vals,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N) {
|
||||
__device__ __forceinline__ float CubicInterpolationRowwise(
|
||||
const T* image, int x, int y, int input_height, int input_width,
|
||||
float coeff0, float coeff1, float coeff2, float coeff3) {
|
||||
int row_index = max(0, min(y, input_height - 1)) * input_width;
|
||||
return coeff0 * static_cast<float>(image[row_index + max(0, min(x - 1, input_width - 1))])
|
||||
+ coeff1 * static_cast<float>(image[row_index + max(0, min(x, input_width - 1))])
|
||||
+ coeff2 * static_cast<float>(image[row_index + max(0, min(x + 1, input_width - 1))])
|
||||
+ coeff3 * static_cast<float>(image[row_index + max(0, min(x + 2, input_width - 1))]);
|
||||
}
|
||||
|
||||
struct CubicMappingInfo {
|
||||
int origin_;
|
||||
int extrapolate_;
|
||||
float coeff0_;
|
||||
float coeff1_;
|
||||
float coeff2_;
|
||||
float coeff3_;
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
__global__ void _ResizeCubicCoordinateMapping(
|
||||
int64_t input_height, int64_t input_width,
|
||||
int64_t output_height, int64_t output_width,
|
||||
float scale_height, float scale_width,
|
||||
float roi_height_start, float roi_height_end,
|
||||
float roi_width_start, float roi_width_end,
|
||||
const size_t SumHW, bool extrapolation_enabled,
|
||||
float cubic_coeff_a, bool exclude_outside,
|
||||
CudaFunctionOriginalCoordinate transform_coordinate,
|
||||
CubicMappingInfo* dims_mapping) {
|
||||
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, SumHW);
|
||||
auto& dm = dims_mapping[id];
|
||||
bool is_y_axis = (id < output_height);
|
||||
int max_input_coord = static_cast<int>(is_y_axis ? input_height : input_width);
|
||||
|
||||
float input_coordinat = transform_coordinate(
|
||||
static_cast<float>(is_y_axis ? id : id - output_height),
|
||||
(is_y_axis ? scale_height : scale_width),
|
||||
static_cast<float>(is_y_axis ? output_height : output_width),
|
||||
static_cast<float>(max_input_coord),
|
||||
(is_y_axis ? roi_height_start : roi_width_start),
|
||||
(is_y_axis ? roi_height_end : roi_width_end));
|
||||
int coord_int = static_cast<int>(floor(input_coordinat));
|
||||
float s_coord = abs(input_coordinat - coord_int);
|
||||
float coeff_sum = 1.0f;
|
||||
float coeff_0 = static_cast<float>(((cubic_coeff_a * (s_coord + 1) - 5 * cubic_coeff_a) * (s_coord + 1) + 8 * cubic_coeff_a) * (s_coord + 1) - 4 * cubic_coeff_a);
|
||||
float coeff_1 = static_cast<float>(((cubic_coeff_a + 2) * s_coord - (cubic_coeff_a + 3)) * s_coord * s_coord + 1);
|
||||
float coeff_2 = static_cast<float>(((cubic_coeff_a + 2) * (1 - s_coord) - (cubic_coeff_a + 3)) * (1 - s_coord) * (1 - s_coord) + 1);
|
||||
float coeff_3 = static_cast<float>(((cubic_coeff_a * (2 - s_coord) - 5 * cubic_coeff_a) * (2 - s_coord) + 8 * cubic_coeff_a) * (2 - s_coord) - 4 * cubic_coeff_a);
|
||||
if (exclude_outside) {
|
||||
coeff_0 = (coord_int - 1 < 0 || coord_int - 1 >= max_input_coord) ? 0.0 : coeff_0;
|
||||
coeff_1 = (coord_int + 0 < 0 || coord_int + 0 >= max_input_coord) ? 0.0 : coeff_1;
|
||||
coeff_2 = (coord_int + 1 < 0 || coord_int + 1 >= max_input_coord) ? 0.0 : coeff_2;
|
||||
coeff_3 = (coord_int + 2 < 0 || coord_int + 2 >= max_input_coord) ? 0.0 : coeff_3;
|
||||
coeff_sum = coeff_0 + coeff_1 + coeff_2 + coeff_3;
|
||||
}
|
||||
dm.origin_ = coord_int;
|
||||
dm.coeff0_ = coeff_0 / coeff_sum;
|
||||
dm.coeff1_ = coeff_1 / coeff_sum;
|
||||
dm.coeff2_ = coeff_2 / coeff_sum;
|
||||
dm.coeff3_ = coeff_3 / coeff_sum;
|
||||
dm.extrapolate_ = (int)(extrapolation_enabled && (input_coordinat < 0 || input_coordinat > static_cast<float>(max_input_coord - 1)));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void _ResizeBiCubicKernel(
|
||||
int64_t input_height, int64_t input_width, int64_t output_height, int64_t output_width,
|
||||
fast_divmod div_output_width, fast_divmod div_output_image,
|
||||
const T* input_data, T* output_data, const size_t N, float extrapolation_value,
|
||||
CubicMappingInfo* dims_mapping) {
|
||||
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
|
||||
int bxc, output_image_index, output_x, output_y;
|
||||
div_output_image.divmod(id, bxc, output_image_index);
|
||||
CUDA_LONG input_index = bxc * input_height * input_width;
|
||||
div_output_width.divmod(output_image_index, output_y, output_x);
|
||||
|
||||
CubicMappingInfo& y_info = dims_mapping[output_y];
|
||||
CubicMappingInfo& x_info = dims_mapping[output_x + output_height];
|
||||
if (y_info.extrapolate_ || x_info.extrapolate_) {
|
||||
output_data[id] = extrapolation_value;
|
||||
return;
|
||||
}
|
||||
|
||||
float w0 = x_info.coeff0_;
|
||||
float w1 = x_info.coeff1_;
|
||||
float w2 = x_info.coeff2_;
|
||||
float w3 = x_info.coeff3_;
|
||||
int x_int = x_info.origin_;
|
||||
int y_int = y_info.origin_;
|
||||
const T* image = input_data + input_index;
|
||||
output_data[id] = y_info.coeff0_ * CubicInterpolationRowwise(image, x_int, y_int - 1, input_height, input_width, w0, w1, w2, w3)
|
||||
+ y_info.coeff1_ * CubicInterpolationRowwise(image, x_int, y_int, input_height, input_width, w0, w1, w2, w3)
|
||||
+ y_info.coeff2_ * CubicInterpolationRowwise(image, x_int, y_int + 1, input_height, input_width, w0, w1, w2, w3)
|
||||
+ y_info.coeff3_ * CubicInterpolationRowwise(image, x_int, y_int + 2, input_height, input_width, w0, w1, w2, w3);
|
||||
}
|
||||
|
||||
size_t CalcResizeBufferSize(const onnxruntime::UpsampleMode upsample_mode,
|
||||
const std::vector<int64_t>& output_dims) {
|
||||
switch (upsample_mode) {
|
||||
case UpsampleMode::NN:
|
||||
return 0;
|
||||
case UpsampleMode::LINEAR:
|
||||
return sizeof(BilinearMappingInfo) * std::accumulate(output_dims.rbegin(), output_dims.rbegin() + 2, 0);
|
||||
case UpsampleMode::CUBIC:
|
||||
return sizeof(CubicMappingInfo) * std::accumulate(output_dims.rbegin(), output_dims.rbegin() + 2, 0);
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void ResizeImpl(
|
||||
const UpsampleMode upsample_mode,
|
||||
const int rank,
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& input_shape,
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& output_shape,
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& input_strides,
|
||||
CudaKernel::CudaAsyncBuffer<fast_divmod>& output_div_pitches,
|
||||
CudaKernel::CudaAsyncBuffer<float>& scales_vals,
|
||||
CudaKernel::CudaAsyncBuffer<float>& roi_vals,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N,
|
||||
bool extrapolation_enabled,
|
||||
float extrapolation_value,
|
||||
float cubic_coeff_a,
|
||||
bool exclude_outside,
|
||||
ResizeCoordinateTransformationMode coordinate_transform_mode,
|
||||
ResizeNearestMode nearest_mode,
|
||||
void* dims_mapping) {
|
||||
int blocksPerGrid = (int)(ceil(static_cast<float>(N) / GridDim::maxThreadsPerBlock));
|
||||
if (onnxruntime::UpsampleMode::NN == upsample_mode) {
|
||||
_ResizeNearestKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
|
||||
rank, input_pitches, output_div_pitches, scales_vals,
|
||||
input_data, output_data, N);
|
||||
} else if (onnxruntime::UpsampleMode::LINEAR == upsample_mode && rank == 4) {
|
||||
_ResizeBilinear4DInputKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
|
||||
input_dim2, input_pitches, output_div_pitches, scales_vals,
|
||||
input_data, output_data, N);
|
||||
} else if (onnxruntime::UpsampleMode::LINEAR == upsample_mode && rank == 2) {
|
||||
_ResizeBilinear2DInputKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
|
||||
input_dim2, input_pitches, output_div_pitches, scales_vals,
|
||||
input_data, output_data, N);
|
||||
CudaFunctionOriginalCoordinate transform_coordinate = GetDeviceOriginalCoordinateFunc(coordinate_transform_mode);
|
||||
CudaFunctionNearestPixel calc_nearest_pixel = GetDeviceNearstPixelFunction(nearest_mode);
|
||||
fast_divmod div_output_image = (rank > 2) ? output_div_pitches.CpuPtr()[rank - 3] : fast_divmod(gsl::narrow_cast<int>(N));
|
||||
int64_t output_height = output_shape.CpuPtr()[rank - 2];
|
||||
int64_t output_width = output_shape.CpuPtr()[rank - 1];
|
||||
int blocksPerDimsMappingGrid = (int)(ceil(static_cast<float>(output_height + output_width) / 32));
|
||||
|
||||
switch (upsample_mode) {
|
||||
case UpsampleMode::NN:
|
||||
input_shape.CopyToGpu();
|
||||
output_shape.CopyToGpu();
|
||||
roi_vals.CopyToGpu();
|
||||
scales_vals.CopyToGpu();
|
||||
input_strides.CopyToGpu();
|
||||
output_div_pitches.CopyToGpu();
|
||||
_ResizeNearestKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
|
||||
rank, input_shape.GpuPtr(), output_shape.GpuPtr(),
|
||||
input_strides.GpuPtr(), output_div_pitches.GpuPtr(),
|
||||
scales_vals.GpuPtr(), roi_vals.GpuPtr(),
|
||||
input_data, output_data, N,
|
||||
extrapolation_enabled, extrapolation_value,
|
||||
transform_coordinate, calc_nearest_pixel);
|
||||
return;
|
||||
case UpsampleMode::LINEAR:
|
||||
_ResizeBilinearCoordinateMapping<T><<<blocksPerDimsMappingGrid, 32, 0>>>(
|
||||
input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1],
|
||||
output_height, output_width,
|
||||
scales_vals.CpuPtr()[rank - 2], scales_vals.CpuPtr()[rank - 1],
|
||||
roi_vals.CpuPtr()[rank - 2], roi_vals.CpuPtr()[rank - 2 + rank],
|
||||
roi_vals.CpuPtr()[rank - 1], roi_vals.CpuPtr()[rank - 1 + rank],
|
||||
output_height + output_width, extrapolation_enabled, transform_coordinate,
|
||||
reinterpret_cast<BilinearMappingInfo*>(dims_mapping));
|
||||
_ResizeBilinearKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
|
||||
input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1],
|
||||
output_height, output_width,
|
||||
output_div_pitches.CpuPtr()[rank - 2], div_output_image,
|
||||
input_data, output_data, N, extrapolation_value,
|
||||
reinterpret_cast<BilinearMappingInfo*>(dims_mapping));
|
||||
return;
|
||||
case UpsampleMode::CUBIC:
|
||||
_ResizeCubicCoordinateMapping<T><<<blocksPerDimsMappingGrid, 32, 0>>>(
|
||||
input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1],
|
||||
output_height, output_width,
|
||||
scales_vals.CpuPtr()[rank - 2], scales_vals.CpuPtr()[rank - 1],
|
||||
roi_vals.CpuPtr()[rank - 2], roi_vals.CpuPtr()[rank - 2 + rank],
|
||||
roi_vals.CpuPtr()[rank - 1], roi_vals.CpuPtr()[rank - 1 + rank],
|
||||
output_height + output_width, extrapolation_enabled,
|
||||
cubic_coeff_a, exclude_outside, transform_coordinate,
|
||||
reinterpret_cast<CubicMappingInfo*>(dims_mapping));
|
||||
_ResizeBiCubicKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
|
||||
input_shape.CpuPtr()[rank - 2], input_shape.CpuPtr()[rank - 1],
|
||||
output_height, output_width,
|
||||
output_div_pitches.CpuPtr()[rank - 2], div_output_image,
|
||||
input_data, output_data, N, extrapolation_value,
|
||||
reinterpret_cast<CubicMappingInfo*>(dims_mapping));
|
||||
// CUDA_CALL(cudaGetLastError());
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
#define SPECIALIZED_IMPL(T) \
|
||||
template void ResizeImpl<T>(const onnxruntime::UpsampleMode upsample_mode, \
|
||||
const size_t rank, \
|
||||
const int64_t input_dim2, \
|
||||
const int64_t* input_pitches, \
|
||||
const fast_divmod* output_div_pitches, \
|
||||
const float* scales_vals, \
|
||||
const T* input_data, \
|
||||
T* output_data, \
|
||||
const size_t N);
|
||||
#define SPECIALIZED_IMPL(T) \
|
||||
template void ResizeImpl<T>( \
|
||||
const UpsampleMode upsample_mode, \
|
||||
const int rank, \
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& input_shape, \
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& output_shape, \
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& input_strides, \
|
||||
CudaKernel::CudaAsyncBuffer<fast_divmod>& output_div_pitches, \
|
||||
CudaKernel::CudaAsyncBuffer<float>& scales_vals, \
|
||||
CudaKernel::CudaAsyncBuffer<float>& roi_vals, \
|
||||
const T* input_data, \
|
||||
T* output_data, \
|
||||
const size_t N, \
|
||||
bool extrapolation_enabled, \
|
||||
float extrapolation_value, \
|
||||
float cubic_coeff_a, \
|
||||
bool exclude_outside, \
|
||||
ResizeCoordinateTransformationMode coordinate_transform_mode, \
|
||||
ResizeNearestMode nearest_mode, \
|
||||
void* dims_mapping);
|
||||
|
||||
SPECIALIZED_IMPL(float)
|
||||
SPECIALIZED_IMPL(double)
|
||||
|
|
|
|||
|
|
@ -6,20 +6,34 @@
|
|||
#include "core/providers/cuda/shared_inc/cuda_utils.h"
|
||||
#include "core/common/common.h"
|
||||
#include "core/providers/cpu/tensor/resize.h"
|
||||
#include "core/providers/cuda/cuda_common.h"
|
||||
|
||||
namespace onnxruntime {
|
||||
namespace cuda {
|
||||
|
||||
size_t CalcResizeBufferSize(const onnxruntime::UpsampleMode upsample_mode,
|
||||
const std::vector<int64_t>& output_dims);
|
||||
|
||||
template <typename T>
|
||||
void ResizeImpl(const onnxruntime::UpsampleMode upsample_mode,
|
||||
const size_t rank,
|
||||
const int64_t input_dim2,
|
||||
const int64_t* input_pitches,
|
||||
const fast_divmod* output_div_pitches,
|
||||
const float* scales_vals,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N);
|
||||
void ResizeImpl(
|
||||
const onnxruntime::UpsampleMode upsample_mode,
|
||||
const int rank,
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& input_shape,
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& output_shape,
|
||||
CudaKernel::CudaAsyncBuffer<int64_t>& input_strides,
|
||||
CudaKernel::CudaAsyncBuffer<fast_divmod>& output_div_pitches,
|
||||
CudaKernel::CudaAsyncBuffer<float>& scales_vals,
|
||||
CudaKernel::CudaAsyncBuffer<float>& roi,
|
||||
const T* input_data,
|
||||
T* output_data,
|
||||
const size_t N,
|
||||
bool extrapolation_enabled,
|
||||
float extrapolation_value,
|
||||
float cubic_coeff_a,
|
||||
bool exclude_outside,
|
||||
onnxruntime::ResizeCoordinateTransformationMode coordinate_transform_mode,
|
||||
onnxruntime::ResizeNearestMode nearest_mode,
|
||||
void* dims_mapping);
|
||||
|
||||
} // namespace cuda
|
||||
} // namespace onnxruntime
|
||||
|
|
|
|||
|
|
@ -31,34 +31,26 @@ REGISTER_KERNEL_TYPED(int32_t)
|
|||
REGISTER_KERNEL_TYPED(uint8_t)
|
||||
|
||||
template <typename T>
|
||||
Status Upsample<T>::BaseCompute(OpKernelContext* context, const std::vector<float>& scales) const {
|
||||
Status Upsample<T>::BaseCompute(OpKernelContext* context,
|
||||
const std::vector<float>& roi,
|
||||
const std::vector<float>& scales,
|
||||
const std::vector<int64_t>& output_dims) const {
|
||||
const Tensor* X = context->Input<Tensor>(0);
|
||||
|
||||
ORT_ENFORCE(nullptr != X);
|
||||
const std::vector<int64_t>& X_dims = X->Shape().GetDims();
|
||||
auto rank = X_dims.size();
|
||||
|
||||
ORT_ENFORCE(output_dims.size() == rank, "Rank of input and output tensor should be same.");
|
||||
if (rank == 0)
|
||||
return Status(ONNXRUNTIME, INVALID_ARGUMENT,
|
||||
is_resize ? "Resize: input tensor cannot be scalar." : "Upsample: input tensor cannot be scalar.");
|
||||
|
||||
return Status(ONNXRUNTIME, INVALID_ARGUMENT,
|
||||
is_resize_ ? "Resize: input tensor cannot be scalar." : "Upsample: input tensor cannot be scalar.");
|
||||
if (rank != scales.size())
|
||||
return Status(ONNXRUNTIME, INVALID_ARGUMENT,
|
||||
is_resize ? "Resize: input tensor's dimension does not match the scales." :
|
||||
"Upsample: input tensor's dimension does not match the scales.");
|
||||
return Status(ONNXRUNTIME, INVALID_ARGUMENT,
|
||||
is_resize_ ? "Resize: input tensor's dimension does not match the scales." : "Upsample: input tensor's dimension does not match the scales.");
|
||||
if (roi.size() != 2 * X->Shape().GetDims().size())
|
||||
return Status(ONNXRUNTIME, INVALID_ARGUMENT,
|
||||
"Resize: size of roi array should be 2 * N where N is the rank of input tensor X.");
|
||||
|
||||
if (UpsampleMode::LINEAR == mode_ && rank != 4 && rank != 2) {
|
||||
std::ostringstream oss;
|
||||
oss << "'Linear' mode only support 2-D inputs ('Bilinear') or 4-D inputs "
|
||||
"with the corresponding outermost 2 scale values being 1 in the ";
|
||||
oss << (is_resize ? "Resize operator" : "Upsample operator");
|
||||
return Status(ONNXRUNTIME, FAIL, oss.str());
|
||||
}
|
||||
|
||||
std::vector<int64_t> Y_dims;
|
||||
for (std::size_t i = 0; i < rank; i++) {
|
||||
Y_dims.push_back(static_cast<int64_t>(scales[i] * X_dims[i]));
|
||||
}
|
||||
Tensor* Y = context->Output(0, Y_dims);
|
||||
Tensor* Y = context->Output(0, output_dims);
|
||||
typedef typename ToCudaType<T>::MappedType CudaT;
|
||||
|
||||
// kernel
|
||||
|
|
@ -66,7 +58,7 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context, const std::vector<floa
|
|||
CudaAsyncBuffer<int64_t> input_strides(this, rank);
|
||||
gsl::span<int64_t> input_stride_span = input_strides.CpuSpan();
|
||||
|
||||
TensorPitches output_pitches(Y_dims);
|
||||
TensorPitches output_pitches(output_dims);
|
||||
CudaAsyncBuffer<fast_divmod> output_div_pitches(this, rank);
|
||||
gsl::span<fast_divmod> div_strides_span = output_div_pitches.CpuSpan();
|
||||
|
||||
|
|
@ -74,24 +66,29 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context, const std::vector<floa
|
|||
input_stride_span[i] = input_pitches[i];
|
||||
div_strides_span[i] = fast_divmod(gsl::narrow_cast<int>(output_pitches[i]));
|
||||
}
|
||||
input_strides.CopyToGpu();
|
||||
output_div_pitches.CopyToGpu();
|
||||
|
||||
size_t output_count = Y->Shape().Size();
|
||||
|
||||
if (is_resize) {
|
||||
if (is_resize_) {
|
||||
CudaAsyncBuffer<int64_t> input_shape(this, X_dims);
|
||||
CudaAsyncBuffer<int64_t> output_shape(this, output_dims);
|
||||
CudaAsyncBuffer<float> roi_vals(this, roi);
|
||||
CudaAsyncBuffer<float> scales_vals(this, scales);
|
||||
scales_vals.CopyToGpu();
|
||||
ResizeImpl(mode_,
|
||||
rank,
|
||||
(UpsampleMode::LINEAR == mode_) ? (rank == 2 ? X_dims[0] : X_dims[2]) : 0,
|
||||
input_strides.GpuPtr(),
|
||||
output_div_pitches.GpuPtr(),
|
||||
scales_vals.GpuPtr(),
|
||||
|
||||
size_t temp_buffer_size = CalcResizeBufferSize(mode_, output_dims);
|
||||
auto dims_mapping_buffer = GetScratchBuffer<unsigned char>(temp_buffer_size);
|
||||
void* dims_mapping = reinterpret_cast<void*>(dims_mapping_buffer.get());
|
||||
ResizeImpl(mode_, (int)rank, input_shape, output_shape,
|
||||
input_strides, output_div_pitches, scales_vals, roi_vals,
|
||||
reinterpret_cast<const CudaT*>(X->template Data<T>()),
|
||||
reinterpret_cast<CudaT*>(Y->template MutableData<T>()),
|
||||
output_count);
|
||||
output_count, use_extrapolation_, extrapolation_value_,
|
||||
cubic_coeff_a_, exclude_outside_,
|
||||
coordinate_transform_mode_, nearest_mode_,
|
||||
dims_mapping);
|
||||
} else {
|
||||
input_strides.CopyToGpu();
|
||||
output_div_pitches.CopyToGpu();
|
||||
|
||||
CudaAsyncBuffer<fast_divmod> scales_div(this, rank);
|
||||
gsl::span<fast_divmod> scales_div_span = scales_div.CpuSpan();
|
||||
|
||||
|
|
@ -102,7 +99,7 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context, const std::vector<floa
|
|||
|
||||
UpampleImpl(mode_,
|
||||
rank,
|
||||
(UpsampleMode::LINEAR == mode_) ? (rank == 2 ? X_dims[0] : X_dims[2]) : 0,
|
||||
(UpsampleMode::LINEAR == mode_) ? (rank == 2 ? X_dims[0] : X_dims[2]) : 0,
|
||||
input_strides.GpuPtr(),
|
||||
output_div_pitches.GpuPtr(),
|
||||
scales_div.GpuPtr(),
|
||||
|
|
@ -116,18 +113,52 @@ Status Upsample<T>::BaseCompute(OpKernelContext* context, const std::vector<floa
|
|||
|
||||
template <typename T>
|
||||
Status Upsample<T>::ComputeInternal(OpKernelContext* context) const {
|
||||
// Opset 7
|
||||
if (OpKernel::Node().InputDefs().size() == 1 || scales_cached_) {
|
||||
return BaseCompute(context, scales_);
|
||||
const Tensor* X = context->Input<Tensor>(0);
|
||||
ORT_ENFORCE(X != nullptr);
|
||||
|
||||
std::vector<int64_t> output_dims(X->Shape().GetDims().size());
|
||||
std::vector<float> roi_array(X->Shape().GetDims().size() * 2, 0.0f);
|
||||
if (!roi_cached_) {
|
||||
if (need_roi_input_) {
|
||||
ORT_ENFORCE(roi_input_idx_ > 0, "Invalid roi input index.");
|
||||
ParseRoiData(context->Input<Tensor>(roi_input_idx_), roi_array);
|
||||
}
|
||||
}
|
||||
const std::vector<float>& roi = roi_cached_ ? roi_ : roi_array;
|
||||
|
||||
if (OpKernel::Node().InputDefs().size() == 1) {
|
||||
// Compute output shape from scales and input dims
|
||||
ComputeOutputShape(scales_, X->Shape().GetDims(), output_dims);
|
||||
return BaseCompute(context, roi, scales_, output_dims);
|
||||
}
|
||||
|
||||
// Opset 9
|
||||
const Tensor* scales = context->Input<Tensor>(1);
|
||||
const auto* scales = context->Input<Tensor>(scales_input_idx_);
|
||||
const auto* sizes = context->Input<Tensor>(sizes_input_idx_);
|
||||
ORT_ENFORCE(scales != nullptr);
|
||||
int64_t scales_size = scales->Shape().Size();
|
||||
std::vector<float> scales_arrary(scales_size);
|
||||
ParseScalesData(scales, scales_arrary);
|
||||
return BaseCompute(context, scales_arrary);
|
||||
|
||||
if (scales_cached_) {
|
||||
ORT_ENFORCE(sizes == nullptr, "Only one of scales or sizes must be provided as input.");
|
||||
ComputeOutputShape(scales_, X->Shape().GetDims(), output_dims);
|
||||
return BaseCompute(context, roi, scales_, output_dims);
|
||||
}
|
||||
|
||||
std::vector<float> scales_array(X->Shape().GetDims().size());
|
||||
if (scales != nullptr && scales->Shape().Size() != 0) {
|
||||
// use scales input data
|
||||
ORT_ENFORCE(sizes == nullptr, "Only one of scales or sizes must be provided as input.");
|
||||
ParseScalesData(scales, scales_array);
|
||||
ComputeOutputShape(scales_array, X->Shape().GetDims(), output_dims);
|
||||
} else {
|
||||
// When sizes input is available directly populate it into the output_dims array.
|
||||
ORT_ENFORCE(sizes != nullptr && sizes->Shape().Size() != 0,
|
||||
"Either scales or sizes MUST be provided as input.");
|
||||
ORT_ENFORCE(sizes->Shape().Size() == output_dims.size(),
|
||||
"Resize: input tensor's rank does not match the output tensor's rank.");
|
||||
memcpy(output_dims.data(), sizes->template Data<int64_t>(), sizes->Shape().Size() * sizeof(int64_t));
|
||||
ParseScalesDataFromOutputSize(output_dims, X->Shape().GetDims(), scales_array);
|
||||
}
|
||||
|
||||
return BaseCompute(context, roi, scales_array, output_dims);
|
||||
}
|
||||
|
||||
} // namespace cuda
|
||||
|
|
|
|||
|
|
@ -17,7 +17,8 @@ class Upsample : public UpsampleBase, public CudaKernel {
|
|||
}
|
||||
|
||||
Status ComputeInternal(OpKernelContext* context) const override;
|
||||
Status BaseCompute(OpKernelContext* context, const std::vector<float>& scales) const;
|
||||
Status BaseCompute(OpKernelContext* context, const std::vector<float>& roi, const std::vector<float>& scales,
|
||||
const std::vector<int64_t>& output_dims) const;
|
||||
};
|
||||
|
||||
} // namespace cuda
|
||||
|
|
|
|||
|
|
@ -21,17 +21,5 @@ void UpampleImpl(const onnxruntime::UpsampleMode upsample_mode,
|
|||
T* output_data,
|
||||
const size_t N);
|
||||
|
||||
template <typename T>
|
||||
void ResizeImpl(
|
||||
const onnxruntime::UpsampleMode upsample_mode,
|
||||
int64_t batch_size,
|
||||
int64_t num_channels,
|
||||
int64_t input_height,
|
||||
int64_t input_width,
|
||||
float height_scale,
|
||||
float width_scale,
|
||||
const T* Xdata,
|
||||
T* Ydata,
|
||||
const size_t N);
|
||||
} // namespace cuda
|
||||
} // namespace onnxruntime
|
||||
|
|
|
|||
|
|
@ -441,7 +441,7 @@ TEST(ResizeOpTest, ResizeOpCubicDownSampleTest) {
|
|||
test.Run();
|
||||
}
|
||||
|
||||
TEST(ResizeOpTest, ResizeOpLineartDownSampleTest_exclude_outside) {
|
||||
TEST(ResizeOpTest, ResizeOpCubicDownSampleTest_exclude_outside) {
|
||||
OpTester test("Resize", 11);
|
||||
std::vector<float> roi{};
|
||||
std::vector<float> scales{0.8f, 0.8f};
|
||||
|
|
|
|||
Loading…
Reference in a new issue