Administrator/opencv-MIRROR
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-11-09 18:04:29 -06:00
Code Issues Packages Projects Releases Wiki Activity

refactored suggested changes

Browse Source
This commit is contained in:
sanjay
2026-11-06 12:56:53 +05:30
parent f75b1307f9
commit ad3e9194db
19 changed files with 174 additions and 39 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
modules/dnn/src/cuda/argmax.cu
View File

@@ -35,7 +35,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
std::int64_t best_index = 0;
for (int k = 1; k < axis_size; ++k) {
const T value = p_input[base + static_cast<std::int64_t>(k) * inner_size];
if constexpr (IsArgMax) {
if (IsArgMax) {
if (SelectLast ? value >= best_value : value > best_value) {
best_value = value;
best_index = k;

11
modules/dnn/src/cuda/functors.hpp
View File

@@ -864,12 +864,13 @@ struct PReluFunctor {
CUDA4DNN_DEVICE PReluFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) {
// an unsigned value is never negative, so PRelu passes it through
if constexpr (std::is_unsigned<T>::value)
return x;
else
return x >= T(0) ? x : T(x * y);
return apply(x, y, std::is_unsigned<T>{});
}
// an unsigned value is never negative, so PRelu passes it through; an overload
// instead of if constexpr (C++17), which .cu files built as C++14 can't use
CUDA4DNN_DEVICE static T apply(T x, T, std::true_type) { return x; }
CUDA4DNN_DEVICE static T apply(T x, T y, std::false_type) { return x >= T(0) ? x : T(x * y); }
};
// bitwise functors are instantiated for integer types only

6
modules/dnn/src/cuda/quantize_dequantize.cu
View File

@@ -53,7 +53,11 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
for (auto i : grid_stride_range(n)) {
std::size_t a_idx = (i / slice_size) % sz_a;
int zp = zero_point ? static_cast<int>(zero_point[a_idx]) : 0;
int rounded = __float2int_rn(input[i] / scale[a_idx]);
// clamp before converting: __float2int_rn turns +Inf or a huge value into
// INT_MIN, which would saturate to qmin instead of qmax
float scaled = input[i] / scale[a_idx];
scaled = (scaled != scaled) ? 0.f : fminf(fmaxf(scaled, -1024.f), 1024.f);
int rounded = __float2int_rn(scaled);
output[i] = saturate_quantized<T>(rounded + zp);
}
}

16
modules/dnn/src/cuda/range.cu
View File

@@ -30,13 +30,23 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
__device__ inline double as_double(std::int32_t v) { return static_cast<double>(v); }
__device__ inline double as_double(std::int64_t v) { return static_cast<double>(v); }
template <class T>
__device__ inline T range_value(T start, T delta, std::size_t i) {
return static_cast<T>(as_double(start) + static_cast<double>(i) * as_double(delta));
}
// int64 stays in integer arithmetic: double holds integers exactly only up to 2^53
__device__ inline std::int64_t range_value(std::int64_t start, std::int64_t delta, std::size_t i) {
return start + static_cast<std::int64_t>(i) * delta;
}
template <class T>
__global__ void range(Span<T> output, const T* start_p, const T* delta_p)
{
const double start = as_double(*start_p);
const double delta = as_double(*delta_p);
const T start = *start_p;
const T delta = *delta_p;
for (auto i : grid_stride_range(output.size()))
output[i] = static_cast<T>(start + static_cast<double>(i) * delta);
output[i] = range_value(start, delta, static_cast<std::size_t>(i));
}
}

8
modules/dnn/src/cuda/roi_align.cu
View File

@@ -75,6 +75,12 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
const float step_x = bin_size_w / gw;
const std::int64_t b = read_batch_index(batch_indices, batch_index_element_size, r);
// a kernel can't raise the CPU path's error for an index outside the batch,
// so it writes 0 instead of reading outside the input
if (b < 0 || b >= p.batch_size) {
output[id] = static_cast<T>(0.f);
continue;
}
const T* img = input.data().get() +
((b * p.channels) + c) * static_cast<std::int64_t>(p.height) * p.width;
@@ -133,7 +139,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
CV_Assert(batch_index_element_size == sizeof(std::int32_t) ||
batch_index_element_size == sizeof(std::int64_t));
CV_Assert(params.output_height > 0 && params.output_width > 0);
CV_Assert(params.height > 0 && params.width > 0 && params.channels > 0);
CV_Assert(params.batch_size > 0 && params.height > 0 && params.width > 0 && params.channels > 0);
if (output.size() == 0)
return;

42
modules/dnn/src/cuda/topk.cu
View File

@@ -32,22 +32,41 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
/* float is exact for __half, float, int8 and uint8; int32/int64 compare in
* their own type, since float holds integers exactly only up to 2^24. */
template <class T> struct TopKCompare { using type = float; };
template <> struct TopKCompare<std::int32_t> { using type = std::int32_t; };
template <> struct TopKCompare<std::int64_t> { using type = std::int64_t; };
/* Starting value that any real element beats: infinity for float types,
* and the lowest or highest integer for int32/int64. */
template <class C>
__device__ inline C topk_worst(bool largest) {
using limits = ::cuda::std::numeric_limits<C>;
if (limits::has_infinity)
return largest ? -limits::infinity() : limits::infinity();
return largest ? limits::lowest() : limits::max();
}
template <class C>
struct Candidate {
float value;
C value;
int index;
};
/* Strict weak ordering matching topk2_layer.cpp's comparators: better value
* first, and on an exact tie the smaller index. */
__device__ __forceinline__ bool better_than(const Candidate& a, const Candidate& b, bool largest) {
template <class C>
__device__ inline bool better_than(const Candidate<C>& a, const Candidate<C>& b, bool largest) {
if (a.value == b.value)
return a.index < b.index;
return largest ? (a.value > b.value) : (a.value < b.value);
}
template <class C>
struct MergeCandidate {
bool largest;
__device__ __forceinline__ Candidate operator()(const Candidate& a, const Candidate& b) const {
__device__ inline Candidate<C> operator()(const Candidate<C>& a, const Candidate<C>& b) const {
return better_than(a, b, largest) ? a : b;
}
};
@@ -56,13 +75,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
__global__ void topk(Span<T> values, Span<std::int64_t> indices, View<T> input,
int outer, int dim_axis, int inner, int k, bool largest)
{
using BlockReduce = cub::BlockReduce<Candidate, BlockSize>;
using C = typename TopKCompare<T>::type;
using Cand = Candidate<C>;
using BlockReduce = cub::BlockReduce<Cand, BlockSize>;
__shared__ typename BlockReduce::TempStorage reduce_storage;
__shared__ Candidate selected;
__shared__ Cand selected;
const int nslices = outer * inner;
const float worst = largest ? -::cuda::std::numeric_limits<float>::infinity()
: ::cuda::std::numeric_limits<float>::infinity();
const C worst = topk_worst<C>(largest);
for (int slice = blockIdx.x; slice < nslices; slice += gridDim.x) {
const int b = slice / inner;
@@ -72,11 +92,11 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
const std::int64_t out_base = static_cast<std::int64_t>(b) * k * inner + j;
for (int rank = 0; rank < k; ++rank) {
Candidate best{worst, ::cuda::std::numeric_limits<int>::max()};
Cand best{worst, ::cuda::std::numeric_limits<int>::max()};
for (int u = threadIdx.x; u < dim_axis; u += BlockSize) {
const float v = static_cast<float>(input.data().get()[in_base + static_cast<std::int64_t>(u) * inner]);
const Candidate cand{v, u};
const C v = static_cast<C>(input.data().get()[in_base + static_cast<std::int64_t>(u) * inner]);
const Cand cand{v, u};
/* Everything at or before the previous pick is spent. Ordering
* strictly means no element can be chosen twice, and duplicates
@@ -88,7 +108,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
best = cand;
}
const Candidate winner = BlockReduce(reduce_storage).Reduce(best, MergeCandidate{largest});
const Cand winner = BlockReduce(reduce_storage).Reduce(best, MergeCandidate<C>{largest});
if (threadIdx.x == 0) {
selected = winner;

2
modules/dnn/src/cuda4dnn/kernels/roi_align.hpp
View File

@@ -15,7 +15,7 @@
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
struct RoiAlignParams {
int channels{}, height{}, width{};
int batch_size{}, channels{}, height{}, width{};
int output_height{}, output_width{};
int sampling_ratio{};
float spatial_scale{1.f};

6
modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp
View File

@@ -77,7 +77,8 @@ namespace cv { namespace dnn { namespace cuda4dnn {
auto loops = input.size_range(0, axis);
auto norm_size = input.size_range(axis, input.rank());
if (norm_size == 1) {
// with a size-one axis the result is 0, or the bias when there is one
if (norm_size == 1 && bias.empty()) {
kernels::fill<T>(stream, output, 0.f);
return;
} else {
@@ -114,8 +115,9 @@ namespace cv { namespace dnn { namespace cuda4dnn {
auto loops = input.size_range(0, axis);
auto norm_size = input.size_range(axis, input.rank());
// with a size-one axis the result is 0, or the bias when there is one;
// the optional Mean/InvStdDev outputs still need the statistics
if (norm_size == 1 && outputs.size() == 1) {
if (norm_size == 1 && outputs.size() == 1 && bias.empty()) {
kernels::fill<T>(stream, output, 0.f);
return;
}

6
modules/dnn/src/cuda4dnn/primitives/one_hot.hpp
View File

@@ -25,7 +25,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
public:
OneHotOp(csl::Stream stream_, int axis_, double off_, double on_, bool runtimeValues_)
: stream(std::move(stream_)), axis(axis_),
offValue(static_cast<T>(static_cast<float>(off_))), onValue(static_cast<T>(static_cast<float>(on_))),
offValue(static_cast<T>(off_)), onValue(static_cast<T>(on_)),
runtimeValues(runtimeValues_)
{
}
@@ -50,8 +50,8 @@ namespace cv { namespace dnn { namespace cuda4dnn {
inputs[2].copyTo(values);
CV_Assert(values.total() == 2);
values.convertTo(values64, CV_64F);
offValue = static_cast<T>(static_cast<float>(values64.ptr<double>()[0]));
onValue = static_cast<T>(static_cast<float>(values64.ptr<double>()[1]));
offValue = static_cast<T>(values64.ptr<double>()[0]);
onValue = static_cast<T>(values64.ptr<double>()[1]);
}
std::size_t inner = 1;

1
modules/dnn/src/cuda4dnn/primitives/roi_align.hpp
View File

@@ -45,6 +45,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
* per-run fields are filled here and the attribute-derived ones come from
* the layer. */
kernels::RoiAlignParams p = params;
p.batch_size = inShape[0];
p.channels = inShape[1];
p.height = inShape[2];
p.width = inShape[3];

6
modules/dnn/src/layers/conv2_layer.cpp
View File

@@ -559,7 +559,7 @@ public:
static bool cudaPostKernel(ActivationFunc fn)
{
return fn == getActivationFunc(ACTIV_HARDSWISH) || fn == getActivationFunc(ACTIV_GELU) ||
fn == getActivationFunc(ACTIV_GELU_APPROX);
fn == getActivationFunc(ACTIV_GELU_APPROX) || fn == getActivationFunc(ACTIV_EXP);
}
//! Takes an expression the layer has a kernel for. Only a one-layer chain carries
@@ -1281,6 +1281,10 @@ public:
post = make_cuda_node<cuda4dnn::GeluOp>(targetId, context.stream);
} else if (activationFunc == getActivationFunc(ACTIV_GELU_APPROX)) {
post = make_cuda_node<cuda4dnn::GeluApproximationOp>(targetId, context.stream);
} else if (activationFunc == getActivationFunc(ACTIV_EXP)) {
const float scale = activParams.size() > 0 ? activParams[0] : 1.f;
const float shift = activParams.size() > 1 ? activParams[1] : 0.f;
post = make_cuda_node<cuda4dnn::ExpOp>(targetId, context.stream, scale, shift);
}
config.activation_type = ConvolutionConfiguration::ActivationType::IDENTITY;

6
modules/dnn/src/layers/fusion_cuda.hpp
View File

@@ -45,7 +45,7 @@ inline bool cudaSupportsFusion(const PreparedFusion& fusion, int channels)
return fn == getActivationFunc(ACTIV_GELU) || fn == getActivationFunc(ACTIV_GELU_APPROX) ||
fn == getActivationFunc(ACTIV_RELU) || fn == getActivationFunc(ACTIV_SWISH) ||
fn == getActivationFunc(ACTIV_SIGMOID) || fn == getActivationFunc(ACTIV_TANH) ||
fn == getActivationFunc(ACTIV_MISH);
fn == getActivationFunc(ACTIV_MISH) || fn == getActivationFunc(ACTIV_HARDSIGMOID);
if (channels < 0)
return true;
return (fusion.scaleBuf.empty() || (int)fusion.scaleBuf.total() == channels) &&
@@ -92,6 +92,10 @@ inline Ptr<BackendNode> makeFusedCudaNode(const PreparedFusion& fusion, const Pt
postNode = make_cuda_node<cuda4dnn::TanHOp>(target, stream);
} else if (fn == getActivationFunc(ACTIV_MISH)) {
postNode = make_cuda_node<cuda4dnn::MishOp>(target, stream);
} else if (fn == getActivationFunc(ACTIV_HARDSIGMOID)) {
const float alpha = fusion.activationParams.size() > 0 ? fusion.activationParams[0] : 0.2f;
const float beta = fusion.activationParams.size() > 1 ? fusion.activationParams[1] : 0.5f;
postNode = make_cuda_node<cuda4dnn::HardSigmoidOp>(target, stream, alpha, beta);
} else {
const float slope = fusion.activationParams.empty() ? 0.f : fusion.activationParams[0];
postNode = make_cuda_node<cuda4dnn::ReLUOp>(target, stream, slope);

50
modules/dnn/src/layers/matmulnbits_layer.cpp
View File

@@ -5,6 +5,12 @@
// Third party copyrights are property of their respective owners.
#include "../precomp.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/inner_product.hpp"
using namespace cv::dnn::cuda4dnn;
#endif
#include <opencv2/dnn/shape_utils.hpp>
#include <opencv2/core/hal/intrin.hpp>
@@ -33,6 +39,15 @@ class MatMulNBitsLayerImpl CV_FINAL : public MatMulNBitsLayer {
}
virtual bool supportBackend(int backendId) CV_OVERRIDE {
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || inputs.size() != 1 || blobs.size() != 2)
return false;
const int t = netimpl_->argData(inputs[0]).type;
return t < 0 || t == CV_32F;
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
@@ -168,6 +183,41 @@ class MatMulNBitsLayerImpl CV_FINAL : public MatMulNBitsLayer {
default: runImpl<8>(A, B, scales, Y); break;
}
}
#ifdef HAVE_CUDA
template<int BITS>
Mat dequantWeights() const {
const int n_blk = (K + block_size - 1) / block_size;
const int blob_bytes = block_size * BITS / 8;
const uchar* bq = blobs[0].ptr<const uchar>();
const float* sc = blobs[1].ptr<const float>();
Mat W(N, K, CV_32F);
parallel_for_(Range(0, N), [&](const Range& r) {
AutoBuffer<float, 1024> wcol(n_blk * block_size);
for (int n = r.start; n < r.end; n++) {
for (int blk = 0; blk < n_blk; blk++) {
const size_t b = static_cast<size_t>(n) * n_blk + blk;
dequantBlock<BITS>(bq + b * blob_bytes, blob_bytes, sc[b], wcol.data() + blk * block_size);
}
std::copy(wcol.data(), wcol.data() + K, W.ptr<float>(n));
}
});
return W;
}
Ptr<BackendNode> initCUDA(void* context_, InputArrayOfArrays inputs_, InputArrayOfArrays) CV_OVERRIDE {
auto context = reinterpret_cast<csl::CSLContext*>(context_);
std::vector<UMat> inps;
inputs_.getUMatVector(inps);
if (dequantW.empty())
dequantW = bits == 4 ? dequantWeights<4>() : dequantWeights<8>();
const std::size_t axis = inps[0].dims - 1;
return make_cuda_node<InnerProductOp>(preferableTarget, std::move(context->stream),
std::move(context->cublas_handle), axis, dequantW, Mat());
}
Mat dequantW;
#endif
};
Ptr<MatMulNBitsLayer> MatMulNBitsLayer::create(const LayerParams& params)

3
modules/dnn/src/layers/range_layer.cpp
View File

@@ -82,9 +82,6 @@ public:
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || this->inputs.size() != 3)
return false;
// a dynamic op on CUDA gets inputs 1.. downloaded before sizing; start must already be on the host
if (!netimpl_->isConstArg(this->inputs[0]))
return false;
// an intermediate arg has no type until the first forward; refuse only a known unsupported one
const int t = netimpl_->argData(this->inputs[0]).type;
return t < 0 || isCudaNodeType(t, false);

13
modules/dnn/src/net_impl2.cpp
View File

@@ -600,6 +600,13 @@ void Net::Impl::prepareForInference()
}
}
#ifdef HAVE_CUDA
static bool sizedFromHostInput(const std::string& type, size_t j, size_t ninp)
{
return j > 0 || ninp == 1 || type == "Range";
}
#endif
void Net::Impl::finalizeGraph(const Ptr<Graph>& graph, bool useCUDA)
{
GraphImpl* g = static_cast<GraphImpl*>(graph.get());
@@ -678,7 +685,7 @@ void Net::Impl::finalizeGraph(const Ptr<Graph>& graph, bool useCUDA)
const bool sizedOnHost = op->dynamicOutputShapes() && hostSizedDynamicOps.count(op->type);
for (size_t j = 0; j < ninp; j++) {
Arg inp = op->inputs[j];
if (inp.empty() || !hostArg[inp.idx] || (sizedOnHost && (j > 0 || ninp == 1)))
if (inp.empty() || !hostArg[inp.idx] || (sizedOnHost && sizedFromHostInput(op->type, j, ninp)))
continue;
unmarkHostOp((size_t)producer[inp.idx]);
changed = true;
@@ -1858,7 +1865,7 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
outOrigData.clear();
// a dynamic op on CUDA is sized from its value inputs, read on the host: all but the
// first, or the only one of a single-input op such as ConstantOfShape
// first, the only one of a single-input op such as ConstantOfShape, or all of Range's
const bool dynamicOnCuda = opBackend == DNN_BACKEND_CUDA && op->dynamicOutputShapes();
for (i = 0; i < ninputs; i++) {
@@ -1870,7 +1877,7 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
// A host Shape op reads only the header, which the host shadow already carries.
const bool headerOnly = opidx < gimpl->hostShapeOp_.size() && gimpl->hostShapeOp_[opidx] &&
op->type == "Shape";
if (!headerOnly && (opBackend != DNN_BACKEND_CUDA || (dynamicOnCuda && (i > 0 || ninputs == 1))))
if (!headerOnly && (opBackend != DNN_BACKEND_CUDA || (dynamicOnCuda && sizedFromHostInput(op->type, i, ninputs))))
cudaDownloadArg(inp, argTensor(inp));
#endif
const Mat& m = argTensor(inp);

11
modules/dnn/src/onnx/onnx_importer2.cpp
View File

@@ -2446,9 +2446,11 @@ void ONNXImporter2::parseUpsample(LayerParams& layerParams, const opencv_onnx::N
layerParams.set("zoom_factor_y", scales.at<float>(2));
layerParams.set("zoom_factor_x", scales.at<float>(3));
n_inputs = 1;
} else if (layerParams.get<String>("mode", "nearest") != "nearest") {
// runtime scales, or a 4-D tensor whose H and W the output takes (the simplifier's
// resize-to-reference subgraph); Resize2 handles both, the legacy Resize neither on CUDA
} else if (layerParams.get<String>("mode", "nearest") != "nearest" ||
(!layerParams.get<bool>("align_corners", false) &&
!layerParams.get<bool>("half_pixel_centers", false))) {
// runtime scales go to Resize2, which can stay on CUDA; nearest only in its plain
// asymmetric form, since that is the one Resize2 can match exactly
layerParams.type = "Resize2";
// keep the legacy layer's sampling, which only the two flags selected
if (!layerParams.has("coordinate_transformation_mode")) {
@@ -2458,6 +2460,9 @@ void ONNXImporter2::parseUpsample(LayerParams& layerParams, const opencv_onnx::N
halfPixel ? "half_pixel" :
layerParams.get<bool>("align_corners", false) ? "align_corners" : "asymmetric");
}
// the legacy nearest path (cv::resize INTER_NEAREST) rounds down
if (layerParams.get<String>("mode", "nearest") == "nearest" && !layerParams.has("nearest_mode"))
layerParams.set("nearest_mode", "floor");
}
}
replaceLayerParam(layerParams, "mode", "interpolation");

2
modules/dnn/src/op_cuda.cpp
View File

@@ -163,6 +163,8 @@ void registerCudaCommonExecs()
CV_DNN_REGISTER_EXEC_CLASS(DepthToSpace, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(SpaceToDepth, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(NormalizeBBox, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(Normalize, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(MatMulNBits, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(Pad2, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(Clip, DNN_BACKEND_CUDA, CUDALegacyExec);
CV_DNN_REGISTER_EXEC_CLASS(ConvTranspose2, DNN_BACKEND_CUDA, CUDALegacyExec);

12
modules/dnn/test/test_onnx_conformance.cpp
View File

@@ -1897,6 +1897,18 @@ public:
bool checkFallbacks(Net& net) const
{
// the check below reads the classic engine's layer list; for the new engine,
// ask which ops it actually placed on the CPU
if (net.getMainGraph())
{
if (backend != DNN_BACKEND_CUDA)
return false;
const std::vector<std::string> cpuOps = cudaFallbackOps(net);
for (const std::string& op : cpuOps)
std::cout << "FALLBACK: " << op << " runs on the CPU instead of CUDA" << endl;
return !cpuOps.empty();
}
// Check if all the layers are supported with current backend and target.
// Some layers might be fused so their timings equal to zero.
std::vector<double> timings;

10
modules/dnn/test/test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp
View File

@@ -6,6 +6,16 @@
"test_max_float64",
"test_min_float64",
"test_mod_mixed_sign_float64",
// New-engine CUDA: supportBackend() turns down Scatter/ScatterND reductions and blocked Q/DQ,
// so the whole graph would run on the CPU.
"test_scatter_elements_with_reduction_max",
"test_scatter_elements_with_reduction_min",
"test_scatternd_add",
"test_scatternd_max",
"test_scatternd_min",
"test_scatternd_multiply",
"test_dequantizelinear_blocked",
"test_quantizelinear_blocked",
// New-engine CUDA: arithmetic ops lack integer (int16/uint16/uint32/uint64) support and
// standalone BatchNorm with non-const params; denylisted for now (pass on CPU, wrong on CUDA).
"test_add_int16",