From 0b4288b35e698411ec07a13e829e876166971f88 Mon Sep 17 00:00:00 2001 From: Katikala Date: Mon, 27 Jul 2026 10:47:17 -0600 Subject: [PATCH 1/9] feat(hip): add ReduceL2 ONNX op support Add end-to-end ReduceL2 lowering from onnx.ReduceL2 through hip.reduce_l2, HipToLLVM runtime dispatch, and a HIP kernel. Includes mock runtime support, LIT conversion/e2e tests, and numeric coverage in test_reduce.py. --- include/hip/Dialect/IR/HipBufferize.h | 1 + include/hip/Dialect/IR/HipOps.td | 39 +++++ lib/Conversion/HipToLLVM/HipToLLVMUtils.h | 1 + lib/Conversion/HipToLLVM/ReduceLowering.cpp | 2 + lib/Conversion/OnnxToHip/CMakeLists.txt | 1 + lib/Conversion/OnnxToHip/OnnxToHip.cpp | 1 + lib/Conversion/OnnxToHip/OnnxToHipUtils.h | 2 + .../OnnxToHip/ReduceL2Conversion.cpp | 100 ++++++++++++ lib/Dialect/IR/HipDialect.cpp | 14 ++ lib/Runtime/CMakeLists.txt | 2 + lib/Runtime/Kernels/hip/reduce_sum_kernel.hip | 150 ++++++++++++++++++ .../Kernels/include/hip_custom_kernels.h | 22 +++ lib/Runtime/hipdnn_ep_runtime.h | 13 ++ lib/Runtime/mock/mock_gpu.cpp | 21 +++ lib/Runtime/real/reduce_l2.cpp | 96 +++++++++++ .../hip-to-llvm/test_reduce_l2.mlir | 44 +++++ .../onnx-to-hip/test_reduce_l2.mlir | 57 +++++++ test/lit/e2e/test_reduce_l2_model.mlir | 20 +++ test/numeric/tests/test_reduce.py | 32 ++++ 19 files changed, 618 insertions(+) create mode 100644 lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp create mode 100644 lib/Runtime/real/reduce_l2.cpp create mode 100644 test/lit/Conversion/hip-to-llvm/test_reduce_l2.mlir create mode 100644 test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir create mode 100644 test/lit/e2e/test_reduce_l2_model.mlir diff --git a/include/hip/Dialect/IR/HipBufferize.h b/include/hip/Dialect/IR/HipBufferize.h index 78870ad24..ce7484245 100644 --- a/include/hip/Dialect/IR/HipBufferize.h +++ b/include/hip/Dialect/IR/HipBufferize.h @@ -152,6 +152,7 @@ registerHipBufferizableOpInterfaceModels(DialectRegistry ®istry) { ReduceMaxOp::attachInterface>(*ctx); ReduceMinOp::attachInterface>(*ctx); ReduceMeanOp::attachInterface>(*ctx); + ReduceL2Op::attachInterface>(*ctx); MatMulNBitsOp::attachInterface>( *ctx); QMoEOp::attachInterface>(*ctx); diff --git a/include/hip/Dialect/IR/HipOps.td b/include/hip/Dialect/IR/HipOps.td index 43dcc40ae..e3cab09d3 100644 --- a/include/hip/Dialect/IR/HipOps.td +++ b/include/hip/Dialect/IR/HipOps.td @@ -2731,6 +2731,45 @@ def Hip_ReduceMeanOp : Hip_DpsOp_Reduction<"reduce_mean"> { }]; } +def Hip_ReduceL2Op : Hip_DpsOp_Reduction<"reduce_l2"> { + let summary = "Reduce tensor by L2 norm along axes"; + let description = [{ + Computes the L2 norm (Euclidean norm) of input tensor elements along the + specified axes: output = sqrt(sum(x^2)). Implements the ONNX ReduceL2 + operator (opset 18+). + + Uses destination-passing style: output buffer is provided as argument. + The sum-of-squares and sqrt are performed inside the runtime kernel + (`reduce_size = num_input / num_output`), so this op is tolerant of a + dynamic reduce axis. + + When axes is not provided in ONNX, an empty tensor<0xi64> is passed. + Combined with noop_with_empty_axes attribute, this enables ONNX-compliant + optional axes semantics while keeping all operands required. + + Example: + ```mlir + hip.reduce_l2(%ctx) ins(%data, %axes : memref<128x3x256x32xf16, 1>, memref<1xi64, 1>) + outs(%output : memref<128x3x256x1xf16, 1>) {keepdims = 1 : i64, noop_with_empty_axes = 0 : i64} + ``` + }]; + + let arguments = (ins + Hip_ContextType:$ctx, + Hip_TensorOrMemRef:$data, + Hip_TensorOrMemRef:$axes, + Hip_TensorOrMemRef:$output, + DefaultValuedAttr:$keepdims, + DefaultValuedAttr:$noop_with_empty_axes + ); + + let assemblyFormat = [{ + `(` $ctx `)` `ins` `(` $data `,` $axes `:` type($data) `,` type($axes) `)` + `outs` `(` $output `:` type($output) `)` + attr-dict (`:` type($result_tensors)^)? + }]; +} + def Hip_ReduceMaxOp : Hip_DpsOp_Reduction<"reduce_max"> { let summary = "Reduce tensor by taking maximum along axes"; let description = [{ diff --git a/lib/Conversion/HipToLLVM/HipToLLVMUtils.h b/lib/Conversion/HipToLLVM/HipToLLVMUtils.h index a32c638cd..c8723634b 100644 --- a/lib/Conversion/HipToLLVM/HipToLLVMUtils.h +++ b/lib/Conversion/HipToLLVM/HipToLLVMUtils.h @@ -85,6 +85,7 @@ inline constexpr const char *kWrapPower = "wrap_power"; inline constexpr const char *kWrapRange = "wrap_range"; inline constexpr const char *kWrapReduceSum = "wrap_reduce_sum"; inline constexpr const char *kWrapReduceMean = "wrap_reduce_mean"; +inline constexpr const char *kWrapReduceL2 = "wrap_reduce_l2"; inline constexpr const char *kWrapReduceMax = "wrap_reduce_max"; inline constexpr const char *kWrapReduceMin = "wrap_reduce_min"; inline constexpr const char *kWrapGQA = "wrap_group_query_attention"; diff --git a/lib/Conversion/HipToLLVM/ReduceLowering.cpp b/lib/Conversion/HipToLLVM/ReduceLowering.cpp index 86659643d..df13b3451 100644 --- a/lib/Conversion/HipToLLVM/ReduceLowering.cpp +++ b/lib/Conversion/HipToLLVM/ReduceLowering.cpp @@ -168,6 +168,8 @@ void populateReduceLoweringPatterns(const LLVMTypeConverter &converter, "reduce_sum"); patterns.insert>(converter, kWrapReduceMean, "reduce_mean"); + patterns.insert>(converter, kWrapReduceL2, + "reduce_l2"); patterns.insert>(converter, kWrapReduceMax, "reduce_max"); patterns.insert>(converter, kWrapReduceMin, diff --git a/lib/Conversion/OnnxToHip/CMakeLists.txt b/lib/Conversion/OnnxToHip/CMakeLists.txt index 318d42bcf..36b33160c 100644 --- a/lib/Conversion/OnnxToHip/CMakeLists.txt +++ b/lib/Conversion/OnnxToHip/CMakeLists.txt @@ -16,6 +16,7 @@ add_library(OnnxToHip STATIC CastConversion.cpp ReduceSumConversion.cpp ReduceMeanConversion.cpp + ReduceL2Conversion.cpp MatMulNBitsConversion.cpp QMoEConversion.cpp GatherBlockQuantizedConversion.cpp diff --git a/lib/Conversion/OnnxToHip/OnnxToHip.cpp b/lib/Conversion/OnnxToHip/OnnxToHip.cpp index d586cbc32..e5cf292cd 100644 --- a/lib/Conversion/OnnxToHip/OnnxToHip.cpp +++ b/lib/Conversion/OnnxToHip/OnnxToHip.cpp @@ -478,6 +478,7 @@ static mlir::LogicalResult convertComputeOps(mlir::func::FuncOp funcOp, populateCastConversionPatterns(patterns, ctx); populateReduceSumConversionPatterns(patterns, ctx); populateReduceMeanConversionPatterns(patterns, ctx); + populateReduceL2ConversionPatterns(patterns, ctx); populateGatherConversionPatterns(patterns, ctx); populateCompressConversionPatterns(patterns, ctx); populateOneHotConversionPatterns(patterns, ctx); diff --git a/lib/Conversion/OnnxToHip/OnnxToHipUtils.h b/lib/Conversion/OnnxToHip/OnnxToHipUtils.h index 6ba6b329a..2d4b4b748 100644 --- a/lib/Conversion/OnnxToHip/OnnxToHipUtils.h +++ b/lib/Conversion/OnnxToHip/OnnxToHipUtils.h @@ -306,6 +306,8 @@ void populateReduceSumConversionPatterns(RewritePatternSet &patterns, MLIRContext *ctx); void populateReduceMeanConversionPatterns(RewritePatternSet &patterns, MLIRContext *ctx); +void populateReduceL2ConversionPatterns(RewritePatternSet &patterns, + MLIRContext *ctx); void populateMatMulNBitsConversionPatterns(RewritePatternSet &patterns, MLIRContext *ctx); void populateQMoEConversionPatterns(RewritePatternSet &patterns, diff --git a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp new file mode 100644 index 000000000..ce3b6acb4 --- /dev/null +++ b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp @@ -0,0 +1,100 @@ +/* + * Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. + * Licensed under the MIT License. + */ + +#include "OnnxToHipUtils.h" + +namespace mlir { +namespace hip { +namespace { + +/// onnx.ReduceL2 -> hip.reduce_l2 +/// +/// Direct, dim-tolerant conversion: sqrt(sum(x^2)) along the reduced axes +/// happens inside the runtime kernel, so no static reduce axis is required. +struct ReduceL2ToHip : public mlir::RewritePattern { + ReduceL2ToHip(mlir::MLIRContext *ctx) + : RewritePattern("onnx.ReduceL2", /*benefit=*/1, ctx) {} + + mlir::LogicalResult + matchAndRewrite(mlir::Operation *op, + mlir::PatternRewriter &rewriter) const override; +}; + +mlir::LogicalResult +ReduceL2ToHip::matchAndRewrite(mlir::Operation *op, + mlir::PatternRewriter &rewriter) const { + auto ctxOrFailure = getContextArg(op, rewriter); + if (mlir::failed(ctxOrFailure)) + return mlir::failure(); + mlir::Value context = *ctxOrFailure; + + mlir::Location loc = op->getLoc(); + mlir::Value data = op->getOperand(0); + + int64_t noopWithEmptyAxes = 0; + if (auto noopAttr = + op->getAttrOfType("noop_with_empty_axes")) { + noopWithEmptyAxes = noopAttr.getSInt(); + } + + int64_t keepdims = 1; + if (auto keepdimsAttr = op->getAttrOfType("keepdims")) { + keepdims = keepdimsAttr.getSInt(); + } + + bool axesStaticallyKnown = op->getNumOperands() <= 1; + auto inputType = mlir::dyn_cast(data.getType()); + llvm::SmallVector axesVec; + if (axesStaticallyKnown) { + if (auto axesAttr = op->getAttrOfType("axes")) { + for (auto a : axesAttr) + axesVec.push_back( + mlir::cast(a).getValue().getSExtValue()); + } else if (noopWithEmptyAxes == 0 && inputType) { + for (int64_t i : llvm::seq(inputType.getRank())) + axesVec.push_back(i); + } + } + + auto resultTypeOr = + inferReduceResultType(op, data, axesVec, axesStaticallyKnown, keepdims); + if (mlir::failed(resultTypeOr)) + return rewriter.notifyMatchFailure( + op, "ReduceL2: cannot infer unranked result (need ranked input and " + "static axes)"); + mlir::RankedTensorType resultType = *resultTypeOr; + mlir::Value init = createEmptyTensor(rewriter, loc, resultType, data); + + mlir::Value axesOperand; + if (op->getNumOperands() > 1) { + axesOperand = op->getOperand(1); + } else { + auto axesType = mlir::RankedTensorType::get( + {static_cast(axesVec.size())}, rewriter.getI64Type()); + auto axesAttr = + mlir::DenseIntElementsAttr::get(axesType, llvm::ArrayRef(axesVec)); + axesOperand = + mlir::arith::ConstantOp::create(rewriter, loc, axesType, axesAttr); + } + + auto keepdimsAttr = rewriter.getI64IntegerAttr(keepdims); + auto noopWithEmptyAxesAttr = rewriter.getI64IntegerAttr(noopWithEmptyAxes); + auto hipOp = mlir::hip::ReduceL2Op::create(rewriter, loc, context, data, + axesOperand, init, keepdimsAttr, + noopWithEmptyAxesAttr); + + rewriter.replaceOp(op, hipOp->getResult(0)); + return mlir::success(); +} + +} // namespace + +void populateReduceL2ConversionPatterns(RewritePatternSet &patterns, + MLIRContext *ctx) { + patterns.add(ctx); +} + +} // namespace hip +} // namespace mlir diff --git a/lib/Dialect/IR/HipDialect.cpp b/lib/Dialect/IR/HipDialect.cpp index 218a818b5..a8edeb6e5 100644 --- a/lib/Dialect/IR/HipDialect.cpp +++ b/lib/Dialect/IR/HipDialect.cpp @@ -1188,6 +1188,20 @@ void ReduceMeanOp::getEffects( emitDpsMemoryEffects(getDpsInputOperands(), getDpsInitsMutable(), effects); } +//===----------------------------------------------------------------------===// +// ReduceL2Op: ins(data, axes), outs(output) +//===----------------------------------------------------------------------===// + +MutableOperandRange ReduceL2Op::getDpsInitsMutable() { + return getOutputMutable(); +} + +void ReduceL2Op::getEffects( + SmallVectorImpl> + &effects) { + emitDpsMemoryEffects(getDpsInputOperands(), getDpsInitsMutable(), effects); +} + //===----------------------------------------------------------------------===// // ReduceMaxOp: ins(data, axes), outs(output) //===----------------------------------------------------------------------===// diff --git a/lib/Runtime/CMakeLists.txt b/lib/Runtime/CMakeLists.txt index 0cb86e9cc..fc7355c7e 100644 --- a/lib/Runtime/CMakeLists.txt +++ b/lib/Runtime/CMakeLists.txt @@ -317,6 +317,7 @@ if(NOT BUILD_MOCK_RUNTIME) compile_to_bitcode(real/one_hot.cpp runtime_one_hot.bc) compile_to_bitcode(real/reduce_sum.cpp runtime_reduce_sum.bc) compile_to_bitcode(real/reduce_mean.cpp runtime_reduce_mean.bc) + compile_to_bitcode(real/reduce_l2.cpp runtime_reduce_l2.bc) compile_to_bitcode(real/reduce_max.cpp runtime_reduce_max.bc) compile_to_bitcode(real/reduce_min.cpp runtime_reduce_min.bc) compile_to_bitcode(real/matmul_nbits.cpp runtime_matmul_nbits.bc) @@ -392,6 +393,7 @@ if(NOT BUILD_MOCK_RUNTIME) ${CMAKE_CURRENT_BINARY_DIR}/runtime_one_hot.bc ${CMAKE_CURRENT_BINARY_DIR}/runtime_reduce_sum.bc ${CMAKE_CURRENT_BINARY_DIR}/runtime_reduce_mean.bc + ${CMAKE_CURRENT_BINARY_DIR}/runtime_reduce_l2.bc ${CMAKE_CURRENT_BINARY_DIR}/runtime_reduce_max.bc ${CMAKE_CURRENT_BINARY_DIR}/runtime_reduce_min.bc ${CMAKE_CURRENT_BINARY_DIR}/runtime_matmul_nbits.bc diff --git a/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip b/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip index 1b90d6845..c130b7da0 100644 --- a/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip +++ b/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip @@ -541,6 +541,156 @@ extern "C" int hip_reduce_mean( return static_cast(err); } +// ============================================================================= +// FP16 / FP32 L2 norm reduction (ReduceL2) +// ============================================================================= +// +// Identical structure to reduce_mean_* but accumulates sum(x^2) and writes +// sqrt(sum) instead of mean. +__global__ void reduce_l2_f16_kernel( + const __half* __restrict__ data, + __half* __restrict__ output, + int64_t reduce_size, + int64_t inner, + int64_t num_output_elements) { + extern __shared__ float sdata_l2_f16[]; + + int64_t out_idx = blockIdx.x; + if (out_idx >= num_output_elements) return; + + int64_t oo = out_idx / inner; + int64_t ii = out_idx - oo * inner; + const __half* slice = data + oo * reduce_size * inner + ii; + int tid = threadIdx.x; + + float sum_sq = 0.0f; + for (int64_t i = tid; i < reduce_size; i += blockDim.x) { + float val = __half2float(slice[i * inner]); + sum_sq += val * val; + } + sdata_l2_f16[tid] = sum_sq; + __syncthreads(); + + for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { + if (static_cast(tid) < s) { + sdata_l2_f16[tid] += sdata_l2_f16[tid + s]; + } + __syncthreads(); + } + + if (tid == 0) { + output[out_idx] = __float2half(sqrtf(sdata_l2_f16[0])); + } +} + +__global__ void reduce_l2_f32_kernel( + const float* __restrict__ data, + float* __restrict__ output, + int64_t reduce_size, + int64_t inner, + int64_t num_output_elements) { + extern __shared__ float sdata_l2_f32[]; + + int64_t out_idx = blockIdx.x; + if (out_idx >= num_output_elements) return; + + int64_t oo = out_idx / inner; + int64_t ii = out_idx - oo * inner; + const float* slice = data + oo * reduce_size * inner + ii; + int tid = threadIdx.x; + + float sum_sq = 0.0f; + for (int64_t i = tid; i < reduce_size; i += blockDim.x) { + float val = slice[i * inner]; + sum_sq += val * val; + } + sdata_l2_f32[tid] = sum_sq; + __syncthreads(); + + for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { + if (static_cast(tid) < s) { + sdata_l2_f32[tid] += sdata_l2_f32[tid + s]; + } + __syncthreads(); + } + + if (tid == 0) { + output[out_idx] = sqrtf(sdata_l2_f32[0]); + } +} + +extern "C" int hip_reduce_l2( + void* stream, + const void* data, + void* output, + int64_t num_input_elements, + int64_t num_output_elements, + int64_t inner_size, + int hip_dtype) { + if (num_output_elements <= 0) return 0; + if (inner_size <= 0) inner_size = 1; + + hipStream_t hip_stream = static_cast(stream); + + if (num_input_elements % num_output_elements != 0) { + fprintf(stderr, + "[custom_kernels] hip_reduce_l2: input(%lld) not divisible " + "by output(%lld)\n", + (long long)num_input_elements, (long long)num_output_elements); + return -1; + } + + int64_t reduce_size = num_input_elements / num_output_elements; + + if (hip_dtype != HIP_DTYPE_FLOAT16 && hip_dtype != HIP_DTYPE_FLOAT32) { + fprintf(stderr, + "[custom_kernels] hip_reduce_l2: unsupported dtype=%d " + "(only FLOAT16/FLOAT32)\n", + hip_dtype); + return -1; + } + + int block_size = 256; + if (reduce_size < 256) { + block_size = 1; + while (block_size < reduce_size) block_size <<= 1; + if (block_size < 1) block_size = 1; + } + size_t shared_mem = block_size * sizeof(float); + + CUSTOM_KERNELS_DEBUG_LOG("[custom_kernels] hip_reduce_l2: dtype=%s, " + "input=%lld, output=%lld, reduce_size=%lld, block=%d\n", + hip_dtype == HIP_DTYPE_FLOAT16 ? "FLOAT16" : "FLOAT32", + (long long)num_input_elements, (long long)num_output_elements, + (long long)reduce_size, block_size); + + if (hip_dtype == HIP_DTYPE_FLOAT16) { + hipLaunchKernelGGL(reduce_l2_f16_kernel, + dim3(static_cast(num_output_elements)), + dim3(block_size), + shared_mem, hip_stream, + static_cast(data), + static_cast<__half*>(output), + reduce_size, inner_size, num_output_elements); + } else { + hipLaunchKernelGGL(reduce_l2_f32_kernel, + dim3(static_cast(num_output_elements)), + dim3(block_size), + shared_mem, hip_stream, + static_cast(data), + static_cast(output), + reduce_size, inner_size, num_output_elements); + } + + hipError_t err = hipGetLastError(); + if (err != hipSuccess) { + fprintf(stderr, + "[custom_kernels] hip_reduce_l2 launch failed: %s\n", + hipGetErrorString(err)); + } + return static_cast(err); +} + // ============================================================================= // ReduceMax / ReduceProd // diff --git a/lib/Runtime/Kernels/include/hip_custom_kernels.h b/lib/Runtime/Kernels/include/hip_custom_kernels.h index c6732a952..38bb7ce6b 100644 --- a/lib/Runtime/Kernels/include/hip_custom_kernels.h +++ b/lib/Runtime/Kernels/include/hip_custom_kernels.h @@ -975,6 +975,28 @@ HIP_KERNEL_API int hip_reduce_mean( int64_t inner_size, int hip_dtype); +/* ========================================================================= + * ReduceL2 (Parallel L2 Norm Reduction) + * ========================================================================= + * + * Same layout convention and `inner_size` semantics as hip_reduce_sum, but + * accumulates sum(x^2) in float and writes sqrt(sum) to the output. The + * reduction is performed in-kernel so the op needs no compile-time-static reduce + * dim and tolerates a dynamic reduce axis. + * + * Supported types: HIP_DTYPE_FLOAT16 and HIP_DTYPE_FLOAT32 (ONNX ReduceL2 is + * float-domain). Both accumulate in float. Other dtypes return -1. + * Returns: 0 on success, non-zero on failure + */ +HIP_KERNEL_API int hip_reduce_l2( + void* stream, + const void* data, + void* output, + int64_t num_input_elements, + int64_t num_output_elements, + int64_t inner_size, + int hip_dtype); + /* ========================================================================= * Pool — MaxPool / AveragePool / LpPool (1D / 2D / 3D) * ========================================================================= diff --git a/lib/Runtime/hipdnn_ep_runtime.h b/lib/Runtime/hipdnn_ep_runtime.h index 0bf4159f0..1edef1b9c 100644 --- a/lib/Runtime/hipdnn_ep_runtime.h +++ b/lib/Runtime/hipdnn_ep_runtime.h @@ -1045,6 +1045,19 @@ int wrap_reduce_mean(RuntimeState *state, void *data, void *axes, void *output, int64_t keepdims, int64_t noop_with_empty_axes, int64_t inner_size); +// ReduceL2 operation wrapper +// data_type: HIPDNN_EP_DATATYPE_* enum value identifying the element type. +// Supported types: HIPDNN_EP_DATATYPE_HALF, HIPDNN_EP_DATATYPE_FLOAT. +// Computes sqrt(sum(x^2)) in-kernel, so a dynamic reduce axis is tolerated. +// `inner_size` = product of input dims AFTER the reduced axis (1 for a +// trailing/contiguous reduce); enables strided reduction over a non-trailing +// axis. +int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, + int64_t data_num_elements, int64_t output_num_elements, + int64_t axes_num_elements, int64_t data_type, + int64_t keepdims, int64_t noop_with_empty_axes, + int64_t inner_size); + // ReduceMax operation wrapper // data_type: HIPDNN_EP_DATATYPE_* enum value identifying the element type. int wrap_reduce_max(RuntimeState *state, void *data, void *axes, void *output, diff --git a/lib/Runtime/mock/mock_gpu.cpp b/lib/Runtime/mock/mock_gpu.cpp index 155c54bad..a9c294d0d 100644 --- a/lib/Runtime/mock/mock_gpu.cpp +++ b/lib/Runtime/mock/mock_gpu.cpp @@ -1076,6 +1076,27 @@ int wrap_reduce_mean(RuntimeState *state, void *data, void *axes, void *output, return 0; } +int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, + int64_t data_num_elements, int64_t output_num_elements, + int64_t axes_num_elements, int64_t data_type, + int64_t keepdims, int64_t noop_with_empty_axes) { + if (!state) { + fprintf(stderr, "Invalid state in wrap_reduce_l2\n"); + return -1; + } + + MOCK_PRINT( + "[MOCK] wrap_reduce_l2(data_num_elements=%lld, " + "output_num_elements=%lld, axes_num_elements=%lld, data_type=%s(%lld), " + "keepdims=%lld, noop_with_empty_axes=%lld)\n", + (long long)data_num_elements, (long long)output_num_elements, + (long long)axes_num_elements, hipdnn_ep_datatype_name(data_type), + (long long)data_type, (long long)keepdims, + (long long)noop_with_empty_axes); + + return 0; +} + int wrap_cast(RuntimeState *state, void *input, void *output, int64_t num_elements, int64_t src_data_type, int64_t dst_data_type) { diff --git a/lib/Runtime/real/reduce_l2.cpp b/lib/Runtime/real/reduce_l2.cpp new file mode 100644 index 000000000..e1d07dd57 --- /dev/null +++ b/lib/Runtime/real/reduce_l2.cpp @@ -0,0 +1,96 @@ +/* + * Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. + * Licensed under the MIT License. + */ +#include "../debug_log.h" +#include "../hipdnn_ep_runtime.h" +#include "../op_profile.h" +#include "hip_custom_kernels.h" +#include "runtime_types.h" + +#include +#include + +#define HIP_CHECK(cmd) \ + do { \ + hipError_t error = (cmd); \ + if (error != hipSuccess) { \ + fprintf(stderr, "HIP error at %s:%d: %s\n", __FILE__, __LINE__, \ + hipGetErrorString(error)); \ + return -1; \ + } \ + } while (0) + +static int hipdnn_to_hip_dtype_l2(int64_t hipdnn_type) { + switch (hipdnn_type) { + case HIPDNN_EP_DATATYPE_HALF: + return HIP_DTYPE_FLOAT16; + case HIPDNN_EP_DATATYPE_FLOAT: + return HIP_DTYPE_FLOAT32; + default: + return -1; + } +} + +int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, + int64_t data_num_elements, int64_t output_num_elements, + int64_t axes_num_elements, int64_t data_type, + int64_t keepdims, int64_t noop_with_empty_axes, + int64_t inner_size) { + OP_PROFILE( + "reduce_l2", + [&] { + char b[64]; + snprintf(b, sizeof(b), "%lld->%lld", (long long)data_num_elements, + (long long)output_num_elements); + return std::string(b); + }, + state); + if (!state || !data || !output) { + RUNTIME_DEBUG_LOG("[REAL] wrap_reduce_l2: null argument\n"); + return -1; + } + + if (axes_num_elements == 0 && noop_with_empty_axes == 1) { + void *stream = hipdnn_ep_state_get_stream(state); + int64_t element_size_bytes = hipdnn_ep_datatype_size(data_type); + if (element_size_bytes < 0) { + fprintf(stderr, + "[REAL] wrap_reduce_l2: unsupported data_type=%lld for noop " + "memcpy\n", + (long long)data_type); + return -1; + } + int64_t total_bytes = data_num_elements * element_size_bytes; + RUNTIME_DEBUG_LOG( + "[REAL] wrap_reduce_l2: noop_with_empty_axes=1 with empty axes, " + "copying %lld bytes (data_type=%s)\n", + (long long)total_bytes, hipdnn_ep_datatype_name(data_type)); + HIP_CHECK(hipMemcpyAsync(output, data, total_bytes, hipMemcpyDeviceToDevice, + static_cast(stream))); + return 0; + } + + void *stream = hipdnn_ep_state_get_stream(state); + + int hip_dtype = hipdnn_to_hip_dtype_l2(data_type); + if (hip_dtype < 0) { + fprintf(stderr, + "[REAL] wrap_reduce_l2: unsupported data_type=%s(%lld) " + "(supported: f16, f32)\n", + hipdnn_ep_datatype_name(data_type), (long long)data_type); + return -1; + } + + RUNTIME_DEBUG_LOG( + "[REAL] wrap_reduce_l2: data_num=%lld, output_num=%lld, " + "axes_num=%lld, data_type=%s(%lld), keepdims=%lld, " + "noop_with_empty_axes=%lld, hip_dtype=%d -> calling hip_reduce_l2\n", + (long long)data_num_elements, (long long)output_num_elements, + (long long)axes_num_elements, hipdnn_ep_datatype_name(data_type), + (long long)data_type, (long long)keepdims, + (long long)noop_with_empty_axes, hip_dtype); + + return hip_reduce_l2(stream, data, output, data_num_elements, + output_num_elements, inner_size, hip_dtype); +} diff --git a/test/lit/Conversion/hip-to-llvm/test_reduce_l2.mlir b/test/lit/Conversion/hip-to-llvm/test_reduce_l2.mlir new file mode 100644 index 000000000..f102ec22a --- /dev/null +++ b/test/lit/Conversion/hip-to-llvm/test_reduce_l2.mlir @@ -0,0 +1,44 @@ +// Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +// Licensed under the MIT License. + +// ============================================================================ +// TEST PURPOSE: +// Verify HIP reduce_l2 operation is correctly lowered to an LLVM call to the +// wrap_reduce_l2 runtime function. +// ============================================================================ + +// RUN: hip-mlir-opt %s --convert-hip-to-llvm | FileCheck %s + +module { + func.func @reduce_l2_static_last_axis( + %ctx: !hip.context, + %input: memref<128x3x256x32xf16, 1>, + %axes: memref<1xi64, 1>, + %output: memref<128x3x256x1xf16, 1>) { + // CHECK-LABEL: llvm.func @reduce_l2_static_last_axis + + hip.reduce_l2(%ctx) ins(%input, %axes : memref<128x3x256x32xf16, 1>, memref<1xi64, 1>) + outs(%output : memref<128x3x256x1xf16, 1>) + {keepdims = 1 : i64, noop_with_empty_axes = 0 : i64} + + // CHECK: llvm.call @wrap_reduce_l2({{.*}}) : (!llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr, i64, i64, i64, i64, i64, i64, i64) -> i32 + + return + } + + func.func @reduce_l2_dynamic( + %ctx: !hip.context, + %input: memref, + %axes: memref<1xi64, 1>, + %output: memref) { + // CHECK-LABEL: llvm.func @reduce_l2_dynamic + + hip.reduce_l2(%ctx) ins(%input, %axes : memref, memref<1xi64, 1>) + outs(%output : memref) + {keepdims = 0 : i64, noop_with_empty_axes = 0 : i64} + + // CHECK: llvm.call @wrap_reduce_l2({{.*}}) : (!llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr, i64, i64, i64, i64, i64, i64, i64) -> i32 + + return + } +} diff --git a/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir b/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir new file mode 100644 index 000000000..fb7edb5d4 --- /dev/null +++ b/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir @@ -0,0 +1,57 @@ +// Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +// Licensed under the MIT License. + +// ============================================================================ +// TEST PURPOSE: +// Verify ONNX ReduceL2 is lowered DIRECTLY to the first-class hip.reduce_l2 +// op. The sqrt(sum(x^2)) happens in-kernel, so the conversion is dim-tolerant. +// +// This test validates: +// - onnx.ReduceL2 -> hip.reduce_l2 (no Mul/ReduceSum/Sqrt decomposition) +// - keepdims = 1 and keepdims = 0 both supported +// - SwinV2 attention L2-norm shape: [128x3x256x32] -> [128x3x256x1] +// ============================================================================ + +// RUN: hip-mlir-opt %s --hip-add-context-arg --convert-onnx-to-hip | FileCheck %s + +module { + func.func @main_graph(%arg0: tensor<1x4xf16>) -> tensor<1x4xf16> { + return %arg0 : tensor<1x4xf16> + } + + // SwinV2 Q/K L2 norm: trailing axis -1, keepdims=1. + func.func @test_reduce_l2_swinv2_attn(%data: tensor<128x3x256x32xf16>, %axes: tensor<1xi64>) -> tensor<128x3x256x1xf16> { + // CHECK-LABEL: func.func @test_reduce_l2_swinv2_attn + // CHECK-SAME: (%[[CTX:.*]]: !hip.context, %[[DATA:.*]]: tensor<128x3x256x32xf16>, %[[AXES:.*]]: tensor<1xi64>) -> tensor<128x3x256x1xf16> + + %output = "onnx.ReduceL2"(%data, %axes) {keepdims = 1 : si64, noop_with_empty_axes = 0 : si64} : (tensor<128x3x256x32xf16>, tensor<1xi64>) -> tensor<128x3x256x1xf16> + + // CHECK: tensor.empty() : tensor<128x3x256x1xf16> + // CHECK: hip.reduce_l2(%[[CTX]]) ins(%[[DATA]], %[[AXES]] : tensor<128x3x256x32xf16>, tensor<1xi64>) outs({{.*}} : tensor<128x3x256x1xf16>) + // CHECK-NOT: onnx.ReduceSum + // CHECK-NOT: onnx.Sqrt + // CHECK-NOT: hip.alloc + + return %output : tensor<128x3x256x1xf16> + } + + func.func @test_reduce_l2_keepdims(%data: tensor<1x128xf16>, %axes: tensor) -> tensor<1x1xf16> { + // CHECK-LABEL: func.func @test_reduce_l2_keepdims + + %output = "onnx.ReduceL2"(%data, %axes) {keepdims = 1 : si64, noop_with_empty_axes = 0 : si64} : (tensor<1x128xf16>, tensor) -> tensor<1x1xf16> + + // CHECK: hip.reduce_l2(%{{.*}}) ins(%{{.*}}, %{{.*}} : tensor<1x128xf16>, tensor) outs({{.*}} : tensor<1x1xf16>) + + return %output : tensor<1x1xf16> + } + + func.func @test_reduce_l2_no_keepdims(%data: tensor<4x8xf16>, %axes: tensor) -> tensor<4xf16> { + // CHECK-LABEL: func.func @test_reduce_l2_no_keepdims + + %output = "onnx.ReduceL2"(%data, %axes) {keepdims = 0 : si64, noop_with_empty_axes = 0 : si64} : (tensor<4x8xf16>, tensor) -> tensor<4xf16> + + // CHECK: hip.reduce_l2(%{{.*}}) ins(%{{.*}}, %{{.*}} : tensor<4x8xf16>, tensor) outs({{.*}} : tensor<4xf16>) {keepdims = 0 : i64} + + return %output : tensor<4xf16> + } +} diff --git a/test/lit/e2e/test_reduce_l2_model.mlir b/test/lit/e2e/test_reduce_l2_model.mlir new file mode 100644 index 000000000..69cce51ff --- /dev/null +++ b/test/lit/e2e/test_reduce_l2_model.mlir @@ -0,0 +1,20 @@ +// Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +// Licensed under the MIT License. + +// ============================================================================ +// TEST PURPOSE: +// Verify ReduceL2 E2E full pipeline: onnx.ReduceL2 -> hip.reduce_l2 -> wrap_reduce_l2 +// Pattern matches SwinV2 attention Q/K L2 normalization (axis=-1, keepdims=1). +// ============================================================================ + +// RUN: hip-mlir-opt %s --hipdnn-pipeline | FileCheck %s + +module { + func.func @main_graph(%arg0: tensor<128x3x256x32xf16>, %arg1: tensor<1xi64>) -> (tensor<128x3x256x1xf16>) { + // CHECK: llvm.func @wrap_reduce_l2 + // CHECK-NOT: onnx.ReduceL2 + + %0 = "onnx.ReduceL2"(%arg0, %arg1) {keepdims = 1 : si64, noop_with_empty_axes = 0 : si64} : (tensor<128x3x256x32xf16>, tensor<1xi64>) -> tensor<128x3x256x1xf16> + return %0 : tensor<128x3x256x1xf16> + } +} diff --git a/test/numeric/tests/test_reduce.py b/test/numeric/tests/test_reduce.py index 1542e4851..4be6b5628 100644 --- a/test/numeric/tests/test_reduce.py +++ b/test/numeric/tests/test_reduce.py @@ -178,6 +178,38 @@ def test_reduce_mean_strided_channel_axis(self, model_runner): compare_outputs(actual, expected, atol=2e-2, rtol=1e-2) +# --------------------------------------------------------------------------- +# ReduceL2 -- first-class hip.reduce_l2 op. SwinV2 attention Q/K L2 norm uses +# axis=-1, keepdims=1 over head_dim (e.g. 32). Runtime dtype: f16/f32. +# --------------------------------------------------------------------------- + + +class TestReduceL2: + @pytest.mark.parametrize( + "shape,axes,keepdims", + [ + ([4, 8], [1], 1), + ([4, 8], [1], 0), + ([128, 3, 256, 32], [-1], 1), + ], + ) + def test_reduce_l2(self, model_runner, shape, axes, keepdims): + model = _make_reduce_model("ReduceL2", np.float16, shape, axes, keepdims) + rng = np.random.default_rng(501) + x = rng.uniform(-3.0, 3.0, shape).astype(np.float16) + actual, expected = model_runner.run_sample(model, [x]) + compare_outputs(actual, expected, atol=2e-2, rtol=1e-2) + + def test_reduce_l2_swinv2_head_dim(self, model_runner): + """SwinV2 Q/K L2 norm: [128, 3, 256, 32] -> [128, 3, 256, 1].""" + shape = [128, 3, 256, 32] + model = _make_reduce_model("ReduceL2", np.float16, shape, [-1], 1) + rng = np.random.default_rng(502) + x = rng.uniform(-2.0, 2.0, shape).astype(np.float16) + actual, expected = model_runner.run_sample(model, [x]) + compare_outputs(actual, expected, atol=2e-2, rtol=1e-2) + + # --------------------------------------------------------------------------- # ReduceMax (added by qwen-vision-kernels PR) # Runtime dtypes: f16, i32, i64 (NO f32 -- not in the dispatch table) From 718fd835d77a41cc15704b71d46db0d35aabc9df Mon Sep 17 00:00:00 2001 From: Katikala Date: Mon, 27 Jul 2026 10:56:01 -0600 Subject: [PATCH 2/9] style: apply clang-format for ReduceL2 pre-commit Fix lintrunner/clang-format alignment in ReduceL2Conversion.cpp and wrap_reduce_l2 declaration in hipdnn_ep_runtime.h. --- lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp | 6 +++--- lib/Runtime/hipdnn_ep_runtime.h | 8 ++++---- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp index ce3b6acb4..eb432b1c9 100644 --- a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp +++ b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp @@ -81,9 +81,9 @@ ReduceL2ToHip::matchAndRewrite(mlir::Operation *op, auto keepdimsAttr = rewriter.getI64IntegerAttr(keepdims); auto noopWithEmptyAxesAttr = rewriter.getI64IntegerAttr(noopWithEmptyAxes); - auto hipOp = mlir::hip::ReduceL2Op::create(rewriter, loc, context, data, - axesOperand, init, keepdimsAttr, - noopWithEmptyAxesAttr); + auto hipOp = + mlir::hip::ReduceL2Op::create(rewriter, loc, context, data, axesOperand, + init, keepdimsAttr, noopWithEmptyAxesAttr); rewriter.replaceOp(op, hipOp->getResult(0)); return mlir::success(); diff --git a/lib/Runtime/hipdnn_ep_runtime.h b/lib/Runtime/hipdnn_ep_runtime.h index 1edef1b9c..abcff153a 100644 --- a/lib/Runtime/hipdnn_ep_runtime.h +++ b/lib/Runtime/hipdnn_ep_runtime.h @@ -1053,10 +1053,10 @@ int wrap_reduce_mean(RuntimeState *state, void *data, void *axes, void *output, // trailing/contiguous reduce); enables strided reduction over a non-trailing // axis. int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, - int64_t data_num_elements, int64_t output_num_elements, - int64_t axes_num_elements, int64_t data_type, - int64_t keepdims, int64_t noop_with_empty_axes, - int64_t inner_size); + int64_t data_num_elements, int64_t output_num_elements, + int64_t axes_num_elements, int64_t data_type, + int64_t keepdims, int64_t noop_with_empty_axes, + int64_t inner_size); // ReduceMax operation wrapper // data_type: HIPDNN_EP_DATATYPE_* enum value identifying the element type. From 8f40d596fcf9f2d6e50ff4e7cd3833ed2350a8cf Mon Sep 17 00:00:00 2001 From: Katikala Date: Tue, 28 Jul 2026 15:22:00 -0600 Subject: [PATCH 3/9] feat(hip): add dynamic-shape support for ReduceL2 conversion Port ReduceProd-style axes extraction and buildReduceL2Init so dynamic input dims lower correctly; add LIT and numeric coverage. Co-authored-by: Cursor --- .../OnnxToHip/ReduceL2Conversion.cpp | 160 ++++++++++++++++-- .../onnx-to-hip/test_reduce_l2.mlir | 12 ++ test/numeric/tests/test_reduce.py | 18 ++ 3 files changed, 172 insertions(+), 18 deletions(-) diff --git a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp index eb432b1c9..8b51b07c3 100644 --- a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp +++ b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp @@ -5,14 +5,131 @@ #include "OnnxToHipUtils.h" +#include "llvm/ADT/APInt.h" +#include "llvm/ADT/SmallSet.h" + namespace mlir { namespace hip { namespace { +/// Try to recognise \p v as a compile-time 1-D integer constant tensor. +/// Mirrors the helper in SliceConversion.cpp / ReduceProdConversion.cpp. +static mlir::DenseElementsAttr getCompileTimeConstantTensor(mlir::Value value) { + mlir::Operation *defOp = value.getDefiningOp(); + if (!defOp) + return nullptr; + if (auto cst = mlir::dyn_cast(defOp)) + return mlir::dyn_cast(cst.getValue()); + if (auto attr = defOp->getAttr("value")) + if (auto dense = mlir::dyn_cast(attr)) + return dense; + if (auto toTensor = mlir::dyn_cast(defOp)) { + auto bufDef = + toTensor.getBuffer().getDefiningOp(); + if (!bufDef) + return nullptr; + auto module = bufDef->getParentOfType(); + if (!module) + return nullptr; + auto global = + module.lookupSymbol(bufDef.getNameAttr()); + if (!global) + return nullptr; + return mlir::dyn_cast_or_null( + global.getInitialValueAttr()); + } + return nullptr; +} + +static mlir::LogicalResult +extractIntVector(mlir::Value v, llvm::SmallVectorImpl &out) { + if (!v) + return mlir::failure(); + auto dense = getCompileTimeConstantTensor(v); + if (!dense) + return mlir::failure(); + auto tensorType = mlir::dyn_cast(dense.getType()); + if (!tensorType || tensorType.getRank() != 1) + return mlir::failure(); + auto elemTy = tensorType.getElementType(); + if (!elemTy.isInteger(64) && !elemTy.isInteger(32)) + return mlir::failure(); + for (mlir::APInt entry : dense.getValues()) + out.push_back(entry.getSExtValue()); + return mlir::success(); +} + +/// Build the destination `tensor.empty` for ReduceL2. +/// +/// Output shape semantics match other ONNX reduce ops: +/// * keepdims=1: out_rank == in_rank; reduced axes become size 1. +/// * keepdims=0: out_rank == in_rank - #axes; reduced axes are dropped. +/// +/// For dynamic output dims we map back to the source `data` dim using the +/// known axes set (compile-time constant). When axes are not known, fall +/// back to positional alignment (correct for keepdims=1 and all-reduce). +static mlir::Value buildReduceL2Init(mlir::PatternRewriter &rewriter, + mlir::Location loc, + mlir::RankedTensorType resultType, + mlir::Value data, + llvm::ArrayRef axesVec, + bool axesKnown, int64_t keepdims) { + auto dataType = mlir::cast(data.getType()); + int64_t inRank = dataType.getRank(); + + llvm::SmallSet reducedAxes; + for (int64_t a : axesVec) { + if (a < 0) + a += inRank; + reducedAxes.insert(a); + } + + llvm::SmallVector outToIn(resultType.getRank(), -1); + if (axesKnown) { + if (keepdims) { + for (int64_t i = 0; i < resultType.getRank(); ++i) + outToIn[i] = reducedAxes.contains(i) ? -1 : i; + } else { + int64_t outIdx = 0; + for (int64_t i = 0; i < inRank; ++i) { + if (reducedAxes.contains(i)) + continue; + if (outIdx < resultType.getRank()) + outToIn[outIdx] = i; + ++outIdx; + } + } + } else { + for (int64_t i = 0; i < resultType.getRank(); ++i) + outToIn[i] = i < inRank ? i : -1; + } + + llvm::SmallVector dynSizes; + for (int64_t i = 0; i < resultType.getRank(); ++i) { + if (!resultType.isDynamicDim(i)) + continue; + int64_t inIdx = outToIn[i]; + if (inIdx < 0) { + dynSizes.push_back( + mlir::arith::ConstantIndexOp::create(rewriter, loc, 1)); + } else if (dataType.isDynamicDim(inIdx)) { + dynSizes.push_back( + mlir::tensor::DimOp::create(rewriter, loc, data, inIdx)); + } else { + dynSizes.push_back(mlir::arith::ConstantIndexOp::create( + rewriter, loc, dataType.getDimSize(inIdx))); + } + } + return mlir::tensor::EmptyOp::create(rewriter, loc, resultType.getShape(), + resultType.getElementType(), dynSizes); +} + /// onnx.ReduceL2 -> hip.reduce_l2 /// /// Direct, dim-tolerant conversion: sqrt(sum(x^2)) along the reduced axes -/// happens inside the runtime kernel, so no static reduce axis is required. +/// happens inside the runtime kernel, so no static reduce axis is required +/// at runtime. Compile-time shape inference still needs ranked input/output +/// or statically-known axes to build the DPS init tensor. struct ReduceL2ToHip : public mlir::RewritePattern { ReduceL2ToHip(mlir::MLIRContext *ctx) : RewritePattern("onnx.ReduceL2", /*benefit=*/1, ctx) {} @@ -44,40 +161,47 @@ ReduceL2ToHip::matchAndRewrite(mlir::Operation *op, keepdims = keepdimsAttr.getSInt(); } - bool axesStaticallyKnown = op->getNumOperands() <= 1; - auto inputType = mlir::dyn_cast(data.getType()); llvm::SmallVector axesVec; - if (axesStaticallyKnown) { + bool axesKnown = false; + mlir::Value axesOperand; + if (op->getNumOperands() > 1 && + !mlir::isa(op->getOperand(1).getType())) { + axesOperand = op->getOperand(1); + if (mlir::succeeded(extractIntVector(axesOperand, axesVec))) + axesKnown = true; + } else { if (auto axesAttr = op->getAttrOfType("axes")) { for (auto a : axesAttr) axesVec.push_back( mlir::cast(a).getValue().getSExtValue()); - } else if (noopWithEmptyAxes == 0 && inputType) { + axesKnown = true; + } else if (noopWithEmptyAxes == 0) { + auto inputType = mlir::cast(data.getType()); for (int64_t i : llvm::seq(inputType.getRank())) axesVec.push_back(i); + axesKnown = true; + } else { + axesKnown = true; // empty axes, noop } + auto axesType = mlir::RankedTensorType::get( + {static_cast(axesVec.size())}, rewriter.getI64Type()); + auto axesAttr = + mlir::DenseIntElementsAttr::get(axesType, llvm::ArrayRef(axesVec)); + axesOperand = + mlir::arith::ConstantOp::create(rewriter, loc, axesType, axesAttr); } auto resultTypeOr = - inferReduceResultType(op, data, axesVec, axesStaticallyKnown, keepdims); + inferReduceResultType(op, data, axesVec, axesKnown, keepdims); if (mlir::failed(resultTypeOr)) return rewriter.notifyMatchFailure( op, "ReduceL2: cannot infer unranked result (need ranked input and " "static axes)"); mlir::RankedTensorType resultType = *resultTypeOr; - mlir::Value init = createEmptyTensor(rewriter, loc, resultType, data); - mlir::Value axesOperand; - if (op->getNumOperands() > 1) { - axesOperand = op->getOperand(1); - } else { - auto axesType = mlir::RankedTensorType::get( - {static_cast(axesVec.size())}, rewriter.getI64Type()); - auto axesAttr = - mlir::DenseIntElementsAttr::get(axesType, llvm::ArrayRef(axesVec)); - axesOperand = - mlir::arith::ConstantOp::create(rewriter, loc, axesType, axesAttr); - } + mlir::Value init = + buildReduceL2Init(rewriter, loc, resultType, data, axesVec, axesKnown, + keepdims); auto keepdimsAttr = rewriter.getI64IntegerAttr(keepdims); auto noopWithEmptyAxesAttr = rewriter.getI64IntegerAttr(noopWithEmptyAxes); diff --git a/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir b/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir index fb7edb5d4..8d8961f57 100644 --- a/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir +++ b/test/lit/Conversion/onnx-to-hip/test_reduce_l2.mlir @@ -10,6 +10,7 @@ // - onnx.ReduceL2 -> hip.reduce_l2 (no Mul/ReduceSum/Sqrt decomposition) // - keepdims = 1 and keepdims = 0 both supported // - SwinV2 attention L2-norm shape: [128x3x256x32] -> [128x3x256x1] +// - Dynamic input shapes with keepdims = 0 // ============================================================================ // RUN: hip-mlir-opt %s --hip-add-context-arg --convert-onnx-to-hip | FileCheck %s @@ -54,4 +55,15 @@ module { return %output : tensor<4xf16> } + + func.func @reduce_l2_dynamic(%data: tensor, %axes: tensor) -> tensor { + %output = "onnx.ReduceL2"(%data, %axes) {keepdims = 0 : si64, noop_with_empty_axes = 0 : si64} : (tensor, tensor) -> tensor + return %output : tensor + } + + // CHECK-LABEL: func.func @reduce_l2_dynamic + // CHECK-SAME: (%[[CTX:.*]]: !hip.context, %[[DATA:.*]]: tensor, %[[AXES:.*]]: tensor) -> tensor + // CHECK: %[[INIT:.*]] = tensor.empty(%{{.*}}, %{{.*}}) : tensor + // CHECK: hip.reduce_l2(%[[CTX]]) ins(%[[DATA]], %[[AXES]] : tensor, tensor) outs(%[[INIT]] : tensor) {keepdims = 0 : i64} + // CHECK-NOT: hip.alloc } diff --git a/test/numeric/tests/test_reduce.py b/test/numeric/tests/test_reduce.py index 4be6b5628..c85e59c8d 100644 --- a/test/numeric/tests/test_reduce.py +++ b/test/numeric/tests/test_reduce.py @@ -209,6 +209,24 @@ def test_reduce_l2_swinv2_head_dim(self, model_runner): actual, expected = model_runner.run_sample(model, [x]) compare_outputs(actual, expected, atol=2e-2, rtol=1e-2) + def test_reduce_l2_dynamic_last_axis(self, model_runner): + """Dynamic batch/seq dims with last-axis L2 norm: [?, ?, 512] -> [?, ?].""" + tp = np_to_onnx_type(np.float16) + X = helper.make_tensor_value_info("X", tp, [None, None, 512]) + axes_init = numpy_helper.from_array(np.array([-1], dtype=np.int64), name="axes") + Y = helper.make_tensor_value_info("Y", tp, [None, None]) + node = helper.make_node( + "ReduceL2", ["X", "axes"], ["Y"], keepdims=0, noop_with_empty_axes=0 + ) + model = make_model_from_nodes( + [node], [X], [Y], initializers=[axes_init], opset=18 + ) + shape = [4, 8, 512] + rng = np.random.default_rng(503) + x = rng.uniform(-3.0, 3.0, shape).astype(np.float16) + actual, expected = model_runner.run_sample(model, [x]) + compare_outputs(actual, expected, atol=2e-2, rtol=1e-2) + # --------------------------------------------------------------------------- # ReduceMax (added by qwen-vision-kernels PR) From 19190f6ce9c579840ec137f4f8555e0aa6e38c33 Mon Sep 17 00:00:00 2001 From: Katikala Date: Tue, 28 Jul 2026 15:24:11 -0600 Subject: [PATCH 4/9] style: apply clang-format for ReduceL2 dynamic-shape conversion Co-authored-by: Cursor --- lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp index 8b51b07c3..fb8d18ba0 100644 --- a/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp +++ b/lib/Conversion/OnnxToHip/ReduceL2Conversion.cpp @@ -199,9 +199,8 @@ ReduceL2ToHip::matchAndRewrite(mlir::Operation *op, "static axes)"); mlir::RankedTensorType resultType = *resultTypeOr; - mlir::Value init = - buildReduceL2Init(rewriter, loc, resultType, data, axesVec, axesKnown, - keepdims); + mlir::Value init = buildReduceL2Init(rewriter, loc, resultType, data, axesVec, + axesKnown, keepdims); auto keepdimsAttr = rewriter.getI64IntegerAttr(keepdims); auto noopWithEmptyAxesAttr = rewriter.getI64IntegerAttr(noopWithEmptyAxes); From 9369989ed5f4744f41a804b441be237f18703921 Mon Sep 17 00:00:00 2001 From: Katikala Date: Wed, 29 Jul 2026 13:31:40 -0600 Subject: [PATCH 5/9] fix(hip): address ReduceL2 PR review feedback Align mock wrap_reduce_l2 with the 11-arg ABI, harden the noop memcpy path, and clear stale HIP errors before hip_reduce_l2 kernel launch. Co-authored-by: Cursor --- lib/Runtime/Kernels/hip/reduce_sum_kernel.hip | 1 + lib/Runtime/mock/mock_gpu.cpp | 9 ++++++--- lib/Runtime/real/reduce_l2.cpp | 9 ++++++++- 3 files changed, 15 insertions(+), 4 deletions(-) diff --git a/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip b/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip index c130b7da0..6b939a16c 100644 --- a/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip +++ b/lib/Runtime/Kernels/hip/reduce_sum_kernel.hip @@ -664,6 +664,7 @@ extern "C" int hip_reduce_l2( (long long)num_input_elements, (long long)num_output_elements, (long long)reduce_size, block_size); + (void)hipGetLastError(); if (hip_dtype == HIP_DTYPE_FLOAT16) { hipLaunchKernelGGL(reduce_l2_f16_kernel, dim3(static_cast(num_output_elements)), diff --git a/lib/Runtime/mock/mock_gpu.cpp b/lib/Runtime/mock/mock_gpu.cpp index a9c294d0d..da43b14ac 100644 --- a/lib/Runtime/mock/mock_gpu.cpp +++ b/lib/Runtime/mock/mock_gpu.cpp @@ -1079,7 +1079,10 @@ int wrap_reduce_mean(RuntimeState *state, void *data, void *axes, void *output, int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, int64_t data_num_elements, int64_t output_num_elements, int64_t axes_num_elements, int64_t data_type, - int64_t keepdims, int64_t noop_with_empty_axes) { + int64_t keepdims, int64_t noop_with_empty_axes, + int64_t inner_size) { + (void)axes; + (void)inner_size; if (!state) { fprintf(stderr, "Invalid state in wrap_reduce_l2\n"); return -1; @@ -1088,11 +1091,11 @@ int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, MOCK_PRINT( "[MOCK] wrap_reduce_l2(data_num_elements=%lld, " "output_num_elements=%lld, axes_num_elements=%lld, data_type=%s(%lld), " - "keepdims=%lld, noop_with_empty_axes=%lld)\n", + "keepdims=%lld, noop_with_empty_axes=%lld, inner_size=%lld)\n", (long long)data_num_elements, (long long)output_num_elements, (long long)axes_num_elements, hipdnn_ep_datatype_name(data_type), (long long)data_type, (long long)keepdims, - (long long)noop_with_empty_axes); + (long long)noop_with_empty_axes, (long long)inner_size); return 0; } diff --git a/lib/Runtime/real/reduce_l2.cpp b/lib/Runtime/real/reduce_l2.cpp index e1d07dd57..2af563052 100644 --- a/lib/Runtime/real/reduce_l2.cpp +++ b/lib/Runtime/real/reduce_l2.cpp @@ -52,6 +52,13 @@ int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, } if (axes_num_elements == 0 && noop_with_empty_axes == 1) { + if (data_num_elements != output_num_elements) { + fprintf(stderr, + "[REAL] wrap_reduce_l2: noop path shape mismatch " + "data_num_elements=%lld output_num_elements=%lld\n", + (long long)data_num_elements, (long long)output_num_elements); + return -1; + } void *stream = hipdnn_ep_state_get_stream(state); int64_t element_size_bytes = hipdnn_ep_datatype_size(data_type); if (element_size_bytes < 0) { @@ -61,7 +68,7 @@ int wrap_reduce_l2(RuntimeState *state, void *data, void *axes, void *output, (long long)data_type); return -1; } - int64_t total_bytes = data_num_elements * element_size_bytes; + int64_t total_bytes = output_num_elements * element_size_bytes; RUNTIME_DEBUG_LOG( "[REAL] wrap_reduce_l2: noop_with_empty_axes=1 with empty axes, " "copying %lld bytes (data_type=%s)\n", From 8eee86a43b32a0264ad33386265e40a6a4b54cf4 Mon Sep 17 00:00:00 2001 From: Katikala Date: Thu, 30 Jul 2026 22:50:17 -0600 Subject: [PATCH 6/9] feat(onnx-to-hip): stamp Slice params before constant externalization Add SliceShapeFold pre-lowering pass to capture compile-time starts, ends, axes, and steps onto onnx.Slice as hipdnn.slice_* attributes before lowerOnnxConstants externalizes the operand constants. Update SliceDecompose to prefer these stamped attributes so slices decompose to tensor.extract_slice instead of falling through to the hip.slice runtime stub. Add LIT coverage for SwinV2 window and downsample stride slice patterns. Co-authored-by: Cursor --- lib/Conversion/OnnxToHip/CMakeLists.txt | 1 + lib/Conversion/OnnxToHip/OnnxToHip.cpp | 1 + lib/Conversion/OnnxToHip/OnnxToHipUtils.h | 8 + lib/Conversion/OnnxToHip/SliceConversion.cpp | 50 ++++-- lib/Conversion/OnnxToHip/SliceShapeFold.cpp | 163 ++++++++++++++++++ .../onnx-to-hip/test_slice_swinv2_window.mlir | 60 +++++++ 6 files changed, 273 insertions(+), 10 deletions(-) create mode 100644 lib/Conversion/OnnxToHip/SliceShapeFold.cpp create mode 100644 test/lit/Conversion/onnx-to-hip/test_slice_swinv2_window.mlir diff --git a/lib/Conversion/OnnxToHip/CMakeLists.txt b/lib/Conversion/OnnxToHip/CMakeLists.txt index 36b33160c..d033d10de 100644 --- a/lib/Conversion/OnnxToHip/CMakeLists.txt +++ b/lib/Conversion/OnnxToHip/CMakeLists.txt @@ -39,6 +39,7 @@ add_library(OnnxToHip STATIC GatherShapeFold.cpp ReshapeShapeFold.cpp PadShapeFold.cpp + SliceShapeFold.cpp ShapeConversion.cpp ReshapeConversion.cpp CausalConvWithStateConversion.cpp diff --git a/lib/Conversion/OnnxToHip/OnnxToHip.cpp b/lib/Conversion/OnnxToHip/OnnxToHip.cpp index e5cf292cd..a00bf93a5 100644 --- a/lib/Conversion/OnnxToHip/OnnxToHip.cpp +++ b/lib/Conversion/OnnxToHip/OnnxToHip.cpp @@ -790,6 +790,7 @@ void ConvertOnnxToHipPass::runOnOperation() { populateGatherShapeFoldPatterns(preLoweringPatterns, ctx); populateReshapeShapeFoldPatterns(preLoweringPatterns, ctx); populatePadShapeFoldPatterns(preLoweringPatterns, ctx); + populateSliceShapeFoldPatterns(preLoweringPatterns, ctx); populateFastGeluFusionPatterns(preLoweringPatterns, ctx); populateErfGeluFusionPatterns(preLoweringPatterns, ctx); populateProjectorOpsRewritePatterns(preLoweringPatterns, ctx); diff --git a/lib/Conversion/OnnxToHip/OnnxToHipUtils.h b/lib/Conversion/OnnxToHip/OnnxToHipUtils.h index 2d4b4b748..b847b7f26 100644 --- a/lib/Conversion/OnnxToHip/OnnxToHipUtils.h +++ b/lib/Conversion/OnnxToHip/OnnxToHipUtils.h @@ -469,6 +469,14 @@ void populateReshapeShapeFoldPatterns(RewritePatternSet &patterns, void populatePadShapeFoldPatterns(RewritePatternSet &patterns, MLIRContext *ctx); +/// Pre-lowering pattern set: stamp compile-time `onnx.Slice` starts/ends/axes/ +/// steps onto the op as `hipdnn.slice_*` attributes so SliceDecompose can +/// rewrite to `tensor.extract_slice` after `lowerOnnxConstants` externalizes +/// the operand constants. Sibling of PadShapeFold; must run BEFORE +/// lowerOnnxConstants. See SliceShapeFold.cpp. +void populateSliceShapeFoldPatterns(RewritePatternSet &patterns, + MLIRContext *ctx); + /// Pre-lowering pattern set: collapse ORT's inlined `FastGelu` primitive /// chain (Pow / Mul / Sum / Tanh) back into a single /// `onnx.Gelu(approximate="tanh")`. ORT inlines the Gelu function body diff --git a/lib/Conversion/OnnxToHip/SliceConversion.cpp b/lib/Conversion/OnnxToHip/SliceConversion.cpp index a449b4981..a67d3601b 100644 --- a/lib/Conversion/OnnxToHip/SliceConversion.cpp +++ b/lib/Conversion/OnnxToHip/SliceConversion.cpp @@ -68,13 +68,10 @@ static mlir::DenseElementsAttr getCompileTimeConstantTensor(mlir::Value value) { return nullptr; } -/// Extract a 1-D integer tensor constant into a SmallVector. -/// Returns failure if the tensor is missing, not 1-D, or not int32/int64. +/// Populate \p out from a dense 1-D integer tensor attribute. static mlir::LogicalResult -extractIntVector(mlir::Value v, llvm::SmallVectorImpl &out) { - if (!v) - return mlir::failure(); - auto dense = getCompileTimeConstantTensor(v); +denseIntVectorToSmallVector(mlir::DenseElementsAttr dense, + llvm::SmallVectorImpl &out) { if (!dense) return mlir::failure(); auto tensorType = mlir::dyn_cast(dense.getType()); @@ -83,11 +80,34 @@ extractIntVector(mlir::Value v, llvm::SmallVectorImpl &out) { auto elemTy = tensorType.getElementType(); if (!elemTy.isInteger(64) && !elemTy.isInteger(32)) return mlir::failure(); + out.clear(); for (mlir::APInt entry : dense.getValues()) out.push_back(entry.getSExtValue()); return mlir::success(); } +/// Extract a 1-D integer tensor constant into a SmallVector. +/// Returns failure if the tensor is missing, not 1-D, or not int32/int64. +static mlir::LogicalResult +extractIntVector(mlir::Value v, llvm::SmallVectorImpl &out) { + if (!v) + return mlir::failure(); + return denseIntVectorToSmallVector(getCompileTimeConstantTensor(v), out); +} + +/// Prefer compile-time slice params stamped by SliceShapeFold (captured +/// before constant externalization); fall back to reading inline operands. +static mlir::LogicalResult +extractSliceParamVector(mlir::Operation *op, llvm::StringRef attrName, + mlir::Value operand, + llvm::SmallVectorImpl &out) { + if (auto attr = op->getAttrOfType(attrName)) { + out.assign(attr.asArrayRef().begin(), attr.asArrayRef().end()); + return mlir::success(); + } + return extractIntVector(operand, out); +} + /// Normalise an ONNX Slice operand reference (`v`): if it is an `onnx.NoValue` /// placeholder (used for absent optional inputs), returns null Value. static mlir::Value normaliseOptional(mlir::Value v) { @@ -120,8 +140,10 @@ struct SliceDecompose : public mlir::RewritePattern { int64_t rank = dataType.getRank(); llvm::SmallVector startsVec, endsVec; - if (mlir::failed(extractIntVector(op->getOperand(1), startsVec)) || - mlir::failed(extractIntVector(op->getOperand(2), endsVec))) + if (mlir::failed(extractSliceParamVector(op, "hipdnn.slice_starts", + op->getOperand(1), startsVec)) || + mlir::failed(extractSliceParamVector(op, "hipdnn.slice_ends", + op->getOperand(2), endsVec))) return rewriter.notifyMatchFailure( op, "starts/ends are not compile-time constants"); @@ -129,9 +151,13 @@ struct SliceDecompose : public mlir::RewritePattern { if (op->getNumOperands() >= 4) { mlir::Value axes = normaliseOptional(op->getOperand(3)); if (axes) { - if (mlir::failed(extractIntVector(axes, axesVec))) + if (mlir::failed(extractSliceParamVector(op, "hipdnn.slice_axes", axes, + axesVec))) return rewriter.notifyMatchFailure( op, "axes is not a compile-time constant"); + } else if (auto attr = + op->getAttrOfType("hipdnn.slice_axes")) { + axesVec.assign(attr.asArrayRef().begin(), attr.asArrayRef().end()); } } if (axesVec.empty()) @@ -142,9 +168,13 @@ struct SliceDecompose : public mlir::RewritePattern { if (op->getNumOperands() == 5) { mlir::Value steps = normaliseOptional(op->getOperand(4)); if (steps) { - if (mlir::failed(extractIntVector(steps, stepsVec))) + if (mlir::failed(extractSliceParamVector(op, "hipdnn.slice_steps", steps, + stepsVec))) return rewriter.notifyMatchFailure( op, "steps is not a compile-time constant"); + } else if (auto attr = + op->getAttrOfType("hipdnn.slice_steps")) { + stepsVec.assign(attr.asArrayRef().begin(), attr.asArrayRef().end()); } } if (stepsVec.empty()) diff --git a/lib/Conversion/OnnxToHip/SliceShapeFold.cpp b/lib/Conversion/OnnxToHip/SliceShapeFold.cpp new file mode 100644 index 000000000..781c97ba1 --- /dev/null +++ b/lib/Conversion/OnnxToHip/SliceShapeFold.cpp @@ -0,0 +1,163 @@ +/* + * Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. + * Licensed under the MIT License. + */ +//===- SliceShapeFold.cpp - Pre-lowering Slice param stamping ------------===// +// +// Sibling pre-lowering fold to PadShapeFold. Captures the compile-time value +// of `onnx.Slice`'s starts/ends/axes/steps operands onto the op as attributes +// BEFORE `lowerOnnxConstants` externalizes the constants -- so SliceDecompose +// (which runs in `convertComputeOps`, after externalization) can rewrite to +// `tensor.extract_slice` without reading the operand. +// +// SwinV2 window/partition slices use small 1-element i64 constants for every +// param; production builds externalize even those into `memref.global` entries +// with null `initial_value` (bytes live in constants.bin / ORT mem). Without +// this stamp SliceDecompose fails `extractIntVector` and every slice falls +// through to the stub `hip.slice` runtime op. +// +// Before: +// %s = onnx.Constant {value = dense<0> : tensor<1xi64>} +// %e = onnx.Constant {value = dense<8> : tensor<1xi64>} +// %a = onnx.Constant {value = dense<1> : tensor<1xi64>} +// %out = onnx.Slice(%data, %s, %e, %a) +// +// After: +// %out = onnx.Slice(%data, %s, %e, %a) +// {hipdnn.slice_starts = array, +// hipdnn.slice_ends = array, +// hipdnn.slice_axes = array} +// +//===----------------------------------------------------------------------===// + +#include "OnnxToHipUtils.h" + +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/IR/BuiltinTypes.h" +#include "llvm/ADT/SmallVector.h" +#include "llvm/ADT/Statistic.h" +#include "llvm/Support/Debug.h" + +#include + +#define DEBUG_TYPE "slice-shape-fold" + +STATISTIC(NumSliceConstStamps, + "Number of onnx.Slice ops whose constant params were stamped as " + "attributes before externalization"); + +namespace mlir { +namespace hip { + +namespace { + +static mlir::Value normaliseOptional(mlir::Value v) { + if (!v) + return v; + auto defOp = v.getDefiningOp(); + if (defOp && defOp->getName().getStringRef() == "onnx.NoValue") + return mlir::Value(); + return v; +} + +/// Inline 1-D integer constant (`arith.constant` or `onnx.Constant`), or +/// std::nullopt. Runs before externalization so `bufferization.to_tensor` is +/// intentionally not handled here. +static std::optional> +getInlineIntVector(mlir::Value v) { + if (!v) + return std::nullopt; + mlir::Operation *defOp = v.getDefiningOp(); + if (!defOp) + return std::nullopt; + + mlir::DenseElementsAttr dense; + if (auto cst = mlir::dyn_cast(defOp)) + dense = mlir::dyn_cast(cst.getValue()); + else if (defOp->getName().getStringRef() == "onnx.Constant") + dense = defOp->getAttrOfType("value"); + + if (!dense) + return std::nullopt; + auto tensorType = mlir::dyn_cast(dense.getType()); + if (!tensorType || tensorType.getRank() != 1) + return std::nullopt; + auto elemTy = tensorType.getElementType(); + if (!elemTy.isInteger(64) && !elemTy.isInteger(32)) + return std::nullopt; + + llvm::SmallVector out; + for (mlir::APInt entry : dense.getValues()) + out.push_back(entry.getSExtValue()); + return out; +} + +struct SliceStampConstParams : public mlir::RewritePattern { + SliceStampConstParams(mlir::MLIRContext *ctx) + : RewritePattern("onnx.Slice", /*benefit=*/1, ctx) {} + + mlir::LogicalResult + matchAndRewrite(mlir::Operation *op, + mlir::PatternRewriter &rewriter) const override { + if (op->hasAttr("hipdnn.slice_starts")) + return rewriter.notifyMatchFailure(op, "slice.already_stamped"); + + if (op->getNumOperands() < 3 || op->getNumOperands() > 5) + return rewriter.notifyMatchFailure(op, "slice.arity"); + + auto startsVec = getInlineIntVector(op->getOperand(1)); + if (!startsVec) + return rewriter.notifyMatchFailure(op, "slice.starts_not_inline_const"); + auto endsVec = getInlineIntVector(op->getOperand(2)); + if (!endsVec) + return rewriter.notifyMatchFailure(op, "slice.ends_not_inline_const"); + + std::optional> axesVec; + if (op->getNumOperands() >= 4) { + mlir::Value axes = normaliseOptional(op->getOperand(3)); + if (axes) { + axesVec = getInlineIntVector(axes); + if (!axesVec) + return rewriter.notifyMatchFailure(op, "slice.axes_not_inline_const"); + } + } + + std::optional> stepsVec; + if (op->getNumOperands() == 5) { + mlir::Value steps = normaliseOptional(op->getOperand(4)); + if (steps) { + stepsVec = getInlineIntVector(steps); + if (!stepsVec) + return rewriter.notifyMatchFailure(op, "slice.steps_not_inline_const"); + } + } + + rewriter.modifyOpInPlace(op, [&] { + op->setAttr("hipdnn.slice_starts", + rewriter.getDenseI64ArrayAttr(*startsVec)); + op->setAttr("hipdnn.slice_ends", + rewriter.getDenseI64ArrayAttr(*endsVec)); + if (axesVec) + op->setAttr("hipdnn.slice_axes", + rewriter.getDenseI64ArrayAttr(*axesVec)); + if (stepsVec) + op->setAttr("hipdnn.slice_steps", + rewriter.getDenseI64ArrayAttr(*stepsVec)); + }); + + LLVM_DEBUG(llvm::dbgs() << "[" DEBUG_TYPE << "] stamped slice params (" + << startsVec->size() << " entries)\n"); + ++NumSliceConstStamps; + return mlir::success(); + } +}; + +} // namespace + +void populateSliceShapeFoldPatterns(mlir::RewritePatternSet &patterns, + MLIRContext *ctx) { + patterns.add(ctx); +} + +} // namespace hip +} // namespace mlir diff --git a/test/lit/Conversion/onnx-to-hip/test_slice_swinv2_window.mlir b/test/lit/Conversion/onnx-to-hip/test_slice_swinv2_window.mlir new file mode 100644 index 000000000..df1d3a72c --- /dev/null +++ b/test/lit/Conversion/onnx-to-hip/test_slice_swinv2_window.mlir @@ -0,0 +1,60 @@ +// Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. +// Licensed under the MIT License. + +// SwinV2 shifted-window slice from /layers.0/blocks.1/Slice_1: +// input 1x128x256x96 +// axis=1, start=0, end=8 -> 1x8x256x96 +// +// Uses onnx.Constant operands (externalized by convert-onnx-to-hip) so the +// SliceShapeFold -> SliceDecompose path is exercised end-to-end. + +// RUN: mkdir -p %t && hip-mlir-opt --hip-add-context-arg --convert-onnx-to-hip='externalize-min-num-elements=1 externalize-output-dir=%t' %s | FileCheck %s + +module { + func.func @main_graph(%arg0: tensor<4xf32>) -> tensor<4xf32> { + return %arg0 : tensor<4xf32> + } + + func.func @test_swinv2_window_slice(%input: tensor<1x128x256x96xf16>) + -> tensor<1x8x256x96xf16> { + // CHECK-LABEL: func.func @test_swinv2_window_slice + %starts = "onnx.Constant"() {value = dense<0> : tensor<1xi64>} + : () -> tensor<1xi64> + %ends = "onnx.Constant"() {value = dense<8> : tensor<1xi64>} + : () -> tensor<1xi64> + %axes = "onnx.Constant"() {value = dense<1> : tensor<1xi64>} + : () -> tensor<1xi64> + %r = "onnx.Slice"(%input, %starts, %ends, %axes) + : (tensor<1x128x256x96xf16>, tensor<1xi64>, tensor<1xi64>, + tensor<1xi64>) -> tensor<1x8x256x96xf16> + + // CHECK-NOT: onnx.Slice + // CHECK-NOT: hip.slice + // CHECK: tensor.extract_slice {{.*}}[0, 0, 0, 0] [1, 8, 256, 96] [1, 1, 1, 1] + + return %r : tensor<1x8x256x96xf16> + } + + // SwinV2 downsample strided slice from /layers.0/downsample/Slice_2: + // axis=1, start=1, end=INT64_MAX, step=2 on 1x128x256x96 -> 1x64x256x96 + func.func @test_swinv2_downsample_stride_slice(%input: tensor<1x128x256x96xf16>) + -> tensor<1x64x256x96xf16> { + // CHECK-LABEL: func.func @test_swinv2_downsample_stride_slice + %starts = "onnx.Constant"() {value = dense<1> : tensor<1xi64>} + : () -> tensor<1xi64> + %ends = "onnx.Constant"() {value = dense<9223372036854775807> : tensor<1xi64>} + : () -> tensor<1xi64> + %axes = "onnx.Constant"() {value = dense<1> : tensor<1xi64>} + : () -> tensor<1xi64> + %steps = "onnx.Constant"() {value = dense<2> : tensor<1xi64>} + : () -> tensor<1xi64> + %r = "onnx.Slice"(%input, %starts, %ends, %axes, %steps) + : (tensor<1x128x256x96xf16>, tensor<1xi64>, tensor<1xi64>, + tensor<1xi64>, tensor<1xi64>) -> tensor<1x64x256x96xf16> + + // CHECK-NOT: hip.slice + // CHECK: tensor.extract_slice {{.*}}[0, 1, 0, 0] [1, 64, 256, 96] [1, 2, 1, 1] + + return %r : tensor<1x64x256x96xf16> + } +} From 20d82028902037a866bf528f5a5ab438006d8fe4 Mon Sep 17 00:00:00 2001 From: Katikala Date: Thu, 30 Jul 2026 22:57:20 -0600 Subject: [PATCH 7/9] style: apply clang-format for Slice shape fold changes Co-authored-by: Cursor --- lib/Conversion/OnnxToHip/SliceConversion.cpp | 12 ++++++------ lib/Conversion/OnnxToHip/SliceShapeFold.cpp | 6 +++--- 2 files changed, 9 insertions(+), 9 deletions(-) diff --git a/lib/Conversion/OnnxToHip/SliceConversion.cpp b/lib/Conversion/OnnxToHip/SliceConversion.cpp index a67d3601b..57b4694bb 100644 --- a/lib/Conversion/OnnxToHip/SliceConversion.cpp +++ b/lib/Conversion/OnnxToHip/SliceConversion.cpp @@ -155,8 +155,8 @@ struct SliceDecompose : public mlir::RewritePattern { axesVec))) return rewriter.notifyMatchFailure( op, "axes is not a compile-time constant"); - } else if (auto attr = - op->getAttrOfType("hipdnn.slice_axes")) { + } else if (auto attr = op->getAttrOfType( + "hipdnn.slice_axes")) { axesVec.assign(attr.asArrayRef().begin(), attr.asArrayRef().end()); } } @@ -168,12 +168,12 @@ struct SliceDecompose : public mlir::RewritePattern { if (op->getNumOperands() == 5) { mlir::Value steps = normaliseOptional(op->getOperand(4)); if (steps) { - if (mlir::failed(extractSliceParamVector(op, "hipdnn.slice_steps", steps, - stepsVec))) + if (mlir::failed(extractSliceParamVector(op, "hipdnn.slice_steps", + steps, stepsVec))) return rewriter.notifyMatchFailure( op, "steps is not a compile-time constant"); - } else if (auto attr = - op->getAttrOfType("hipdnn.slice_steps")) { + } else if (auto attr = op->getAttrOfType( + "hipdnn.slice_steps")) { stepsVec.assign(attr.asArrayRef().begin(), attr.asArrayRef().end()); } } diff --git a/lib/Conversion/OnnxToHip/SliceShapeFold.cpp b/lib/Conversion/OnnxToHip/SliceShapeFold.cpp index 781c97ba1..3d3f8bb50 100644 --- a/lib/Conversion/OnnxToHip/SliceShapeFold.cpp +++ b/lib/Conversion/OnnxToHip/SliceShapeFold.cpp @@ -128,15 +128,15 @@ struct SliceStampConstParams : public mlir::RewritePattern { if (steps) { stepsVec = getInlineIntVector(steps); if (!stepsVec) - return rewriter.notifyMatchFailure(op, "slice.steps_not_inline_const"); + return rewriter.notifyMatchFailure(op, + "slice.steps_not_inline_const"); } } rewriter.modifyOpInPlace(op, [&] { op->setAttr("hipdnn.slice_starts", rewriter.getDenseI64ArrayAttr(*startsVec)); - op->setAttr("hipdnn.slice_ends", - rewriter.getDenseI64ArrayAttr(*endsVec)); + op->setAttr("hipdnn.slice_ends", rewriter.getDenseI64ArrayAttr(*endsVec)); if (axesVec) op->setAttr("hipdnn.slice_axes", rewriter.getDenseI64ArrayAttr(*axesVec)); From 7cb881094d2dd4122e7a4d5dce38d4bb441d0f1f Mon Sep 17 00:00:00 2001 From: Katikala Date: Thu, 30 Jul 2026 23:30:16 -0600 Subject: [PATCH 8/9] fix(ci): resolve GPU test assets after repo rename Probe legacy onnx-hipdnn-ep asset paths on the self-hosted GPU runner when runner.workspace/model is missing after the hip-ep rename. Also make perf-result publishing non-fatal for fork PRs that cannot post comments with GITHUB_TOKEN. Co-authored-by: Cursor --- .github/workflows/windows-build.yml | 59 ++++++++++++++++++++++------- 1 file changed, 46 insertions(+), 13 deletions(-) diff --git a/.github/workflows/windows-build.yml b/.github/workflows/windows-build.yml index 7e24e87e3..ddc0d66e8 100644 --- a/.github/workflows/windows-build.yml +++ b/.github/workflows/windows-build.yml @@ -700,6 +700,37 @@ jobs: if exist wheelpkg rd /s /q wheelpkg if exist therock-dist rd /s /q therock-dist + # Persistent GPU assets (model/, amdgpu-oga-models/, l2-test-models/) live + # outside the checkout and survive across runs. After the onnx-hipdnn-ep -> + # hip-ep rename, runner.workspace moved but the staged dirs did not; probe + # legacy paths before failing. Override anytime via GPU_TEST_ASSET_ROOT. + - name: Resolve GPU test asset root + shell: pwsh + run: | + $explicit = "${{ vars.GPU_TEST_ASSET_ROOT }}".Trim() + $candidates = [System.Collections.Generic.List[string]]::new() + if ($explicit) { $candidates.Add($explicit) } + $candidates.Add("${{ runner.workspace }}") + $parent = Split-Path "${{ runner.workspace }}" -Parent + $candidates.Add((Join-Path $parent "onnx-hipdnn-ep")) + $candidates.Add((Join-Path $parent "hip-ep")) + + $root = $null + foreach ($c in ($candidates | Select-Object -Unique)) { + if ((Test-Path (Join-Path $c "model")) -or + (Test-Path (Join-Path $c "amdgpu-oga-models")) -or + (Test-Path (Join-Path $c "l2-test-models"))) { + $root = $c + Write-Host "Using GPU test asset root: $root" + break + } + } + if (-not $root) { + $root = if ($explicit) { $explicit } else { "${{ runner.workspace }}" } + Write-Host "WARNING: no GPU asset dirs found; defaulting to: $root" + } + "GPU_TEST_ASSET_ROOT=$root" >> $env:GITHUB_ENV + - name: Download GPU test package uses: actions/download-artifact@v4 with: @@ -733,17 +764,14 @@ jobs: # Output is tee'd to results/ for QPS extraction. Non-zero exit code # sets FAILED but continues so all models are tested. # - # GPU-test asset dirs (model/, amdgpu-oga-models/, l2-test-models/) resolve - # under ${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}: set the repo/org - # Actions variable GPU_TEST_ASSET_ROOT to a stable path (and stage the dirs - # there) so persistent assets need not live under the repo-name-derived - # _work//, which strands them on a repo rename. Unset => runner.workspace - # (unchanged from today). + # GPU-test asset dirs resolve under $env:GPU_TEST_ASSET_ROOT (see + # "Resolve GPU test asset root" above). Set the repo/org Actions variable + # GPU_TEST_ASSET_ROOT to override the auto-detected path. - name: Run onnxruntime_perf_test (GPU) shell: cmd run: | set PATH=%CD%\therock-dist\bin;%CD%\gpu-test-package\bin;%PATH% - set MODEL_DIR=${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\model + set MODEL_DIR=%GPU_TEST_ASSET_ROOT%\model if not exist "%MODEL_DIR%" ( echo [ERROR] MODEL_DIR not found: %MODEL_DIR% exit /b 1 @@ -787,7 +815,7 @@ jobs: set THEROCK_DIST=%CD%\therock-dist set PATH=%CD%\therock-dist\bin;%CD%\gpu-test-package\bin;%PATH% set LIB=%CD%\gpu-test-package\lib;%CD%\therock-dist\lib - set MODEL_DIR=${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\model + set MODEL_DIR=%GPU_TEST_ASSET_ROOT%\model if not exist "%MODEL_DIR%" ( echo [SKIP] MODEL_DIR not found: %MODEL_DIR% exit /b 0 @@ -833,7 +861,7 @@ jobs: set THEROCK_DIST=%CD%\therock-dist set PATH=%CD%\therock-dist\bin;%CD%\gpu-test-package\bin;%PATH% set LIB=%CD%\gpu-test-package\lib;%CD%\therock-dist\lib - set MODEL_DIR=${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\model + set MODEL_DIR=%GPU_TEST_ASSET_ROOT%\model if not exist "%MODEL_DIR%" ( echo [SKIP] MODEL_DIR not found: %MODEL_DIR% exit /b 0 @@ -875,7 +903,7 @@ jobs: $pkgBin = Join-Path $PWD "gpu-test-package\bin" $env:PATH = "$PWD\therock-dist\bin;$pkgBin;$env:PATH" - $modelDir = "${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\model" + $modelDir = "$env:GPU_TEST_ASSET_ROOT\model" if (-not (Test-Path $modelDir)) { Write-Host "[SKIP] MODEL_DIR not found: $modelDir" exit 0 @@ -955,7 +983,7 @@ jobs: # generated once on the GPU runner and persists across runs, so just # consume it here (no re-copy / no config rewrite). It must be present # -- treat its absence as an error, not a skip. - $ogaModelRoot = "${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\amdgpu-oga-models" + $ogaModelRoot = "$env:GPU_TEST_ASSET_ROOT\amdgpu-oga-models" if (-not (Test-Path $ogaModelRoot)) { Write-Host "[ERROR] AMDGPU OGA model directory not found: $ogaModelRoot" exit 1 @@ -1049,7 +1077,7 @@ jobs: $ErrorActionPreference = 'Continue' $PSNativeCommandUseErrorActionPreference = $false - $model = "${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\amdgpu-oga-models\Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml" + $model = "$env:GPU_TEST_ASSET_ROOT\amdgpu-oga-models\Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml" if (-not (Test-Path $model)) { Write-Host "[SKIP] model not found: $model"; exit 0 } # The morphizen wheel now bundles the whole AMD GPU umbrella chain @@ -1098,7 +1126,7 @@ jobs: rem The EP DLL is now hipgpu.dll; hip-onnx-runner's legacy default name rem is onnxruntime_morphizen_ep.dll, so point it at the new DLL. set MORPHIZEN_EP_LIB=%CD%\gpu-test-package\bin\hipgpu.dll - set L2_MODEL_DIR=${{ vars.GPU_TEST_ASSET_ROOT || runner.workspace }}\l2-test-models + set L2_MODEL_DIR=%GPU_TEST_ASSET_ROOT%\l2-test-models if not exist "%L2_MODEL_DIR%" ( echo [SKIP] L2 model directory not found: %L2_MODEL_DIR% exit /b 0 @@ -1227,6 +1255,8 @@ jobs: env: GH_TOKEN: ${{ github.token }} run: | + $ErrorActionPreference = 'Continue' + $PSNativeCommandUseErrorActionPreference = $false $resultsDir = "gpu-test-package/results" $runUrl = "https://github.com/${{ github.repository }}/actions/runs/${{ github.run_id }}" $rawSha = if ("${{ github.event_name }}" -eq "pull_request") { "${{ github.event.pull_request.head.sha }}" } else { "${{ github.sha }}" } @@ -1373,3 +1403,6 @@ jobs: Write-Host "WARNING: gh api unreachable, skipping PR comment" } } + + # Step summary is the source of truth; fork PRs may lack comment perms. + exit 0 From 62309997b63bc426c0b79e68c5ed4440d26d9e01 Mon Sep 17 00:00:00 2001 From: Katikala Date: Fri, 31 Jul 2026 00:26:25 -0600 Subject: [PATCH 9/9] fix(ci): make L2 result publishing non-fatal for fork PRs Mirror the perf publish step: fork PRs cannot post comments with GITHUB_TOKEN, so write results to the step summary and exit cleanly. Co-authored-by: Cursor --- .github/workflows/windows-build.yml | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/.github/workflows/windows-build.yml b/.github/workflows/windows-build.yml index ddc0d66e8..7cc68f564 100644 --- a/.github/workflows/windows-build.yml +++ b/.github/workflows/windows-build.yml @@ -1172,6 +1172,8 @@ jobs: env: GH_TOKEN: ${{ github.token }} run: | + $ErrorActionPreference = 'Continue' + $PSNativeCommandUseErrorActionPreference = $false $l2Dir = "l2-results" if (-not (Test-Path $l2Dir)) { Write-Host "No L2 results to publish (directory not found)" @@ -1246,6 +1248,9 @@ jobs: } } + # Step summary is the source of truth; fork PRs may lack comment perms. + exit 0 + # Parse QPS from result files and publish: # - PR trigger → post/update a PR comment via gh CLI (github-actions[bot]) # - other trigger → write to GITHUB_STEP_SUMMARY only