[Mlir-commits] [mlir] [mlir][acc] Add attributes for parallelism dimensions (PR #182209)
Razvan Lupusoru
llvmlistbot at llvm.org
Wed Feb 18 18:18:51 PST 2026
https://github.com/razvanlupusoru created https://github.com/llvm/llvm-project/pull/182209
As OpenACC gets lowered to eventually mapping to GPU (via GPU dialect), we need to track parallelism assignment which we can use in how variables get privatized, how barriers and synchronizations are inserted to ensure appropriate OpenACC execution model, and for loop work-sharing. This adds GPUParallelDimAttr and GPUParallelDimsAttr for this.
>From 14646d609b29711c2afb34e57bf05ec0fe508591 Mon Sep 17 00:00:00 2001
From: Razvan Lupusoru <rlupusoru at nvidia.com>
Date: Wed, 18 Feb 2026 18:16:21 -0800
Subject: [PATCH] [mlir][acc] Add attributes for parallelism dimensions
As OpenACC gets lowered to eventually mapping to GPU (via
GPU dialect), we need to track parallelism assignment which
we can use in how variables get privatized, how barriers
and synchronizations are inserted to ensure appropriate
OpenACC execution model, and for loop work-sharing. This
adds GPUParallelDimAttr and GPUParallelDimsAttr for this.
---
mlir/include/mlir/Dialect/OpenACC/OpenACC.h | 1 +
.../Dialect/OpenACC/OpenACCCGAttributes.td | 94 ++++++
.../mlir/Dialect/OpenACC/OpenACCOps.td | 1 +
mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp | 283 ++++++++++++++++++
mlir/test/Dialect/OpenACC/invalid-cg.mlir | 21 ++
mlir/test/Dialect/OpenACC/ops-cg.mlir | 79 +++++
6 files changed, 479 insertions(+)
create mode 100644 mlir/include/mlir/Dialect/OpenACC/OpenACCCGAttributes.td
create mode 100644 mlir/test/Dialect/OpenACC/invalid-cg.mlir
create mode 100644 mlir/test/Dialect/OpenACC/ops-cg.mlir
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACC.h b/mlir/include/mlir/Dialect/OpenACC/OpenACC.h
index 2f6d1ee5bdd3e..cdc98a1ed2a4f 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACC.h
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACC.h
@@ -21,6 +21,7 @@
#include "mlir/Bytecode/BytecodeOpInterface.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/OpenACC/OpenACCOpsDialect.h.inc"
#include "mlir/Dialect/OpenACC/OpenACCOpsEnums.h.inc"
#include "mlir/Dialect/OpenACC/OpenACCOpsInterfaces.h.inc"
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCCGAttributes.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCCGAttributes.td
new file mode 100644
index 0000000000000..3bc38be9f348a
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCCGAttributes.td
@@ -0,0 +1,94 @@
+//===- OpenACCCGAttributes.td - OpenACC codegen attributes ------*- tablegen -*-===//
+//
+// Part of the MLIR Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Defines OpenACC codegen attributes (GPU parallel dimensions used for
+// privatization, barrier management, and loop work-sharing).
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OPENACCCG_ATTRIBUTES
+#define OPENACCCG_ATTRIBUTES
+
+def OpenACC_GPUParallelDimAttr : OpenACC_Attr<"GPUParallelDim", "par_dim"> {
+ let summary = "GPU parallel dimension for use in OpenACC parallelism assignment.";
+ let description = [{
+ Identifies a single GPU parallel dimension. Used for privatization scope,
+ barrier placement, and loop work-sharing in OpenACC codegen.
+ }];
+ let parameters = (ins "::mlir::IntegerAttr":$value);
+ let extraClassDeclaration = [{
+ // GPU dimension predicates
+ bool isSeq() const;
+ bool isThreadX() const;
+ bool isThreadY() const;
+ bool isThreadZ() const;
+ bool isBlockX() const;
+ bool isBlockY() const;
+ bool isBlockZ() const;
+ bool isAnyThread() const;
+ bool isAnyBlock() const;
+
+ // Attribute creation from gpu::Processor
+ static GPUParallelDimAttr get(::mlir::MLIRContext *, ::mlir::gpu::Processor);
+
+ // Get the underlying gpu::Processor
+ ::mlir::gpu::Processor getProcessor() const;
+
+ // Factory functions for each dimension
+ static GPUParallelDimAttr seqDim(::mlir::MLIRContext *);
+ static GPUParallelDimAttr threadXDim(::mlir::MLIRContext *);
+ static GPUParallelDimAttr threadYDim(::mlir::MLIRContext *);
+ static GPUParallelDimAttr threadZDim(::mlir::MLIRContext *);
+ static GPUParallelDimAttr blockXDim(::mlir::MLIRContext *);
+ static GPUParallelDimAttr blockYDim(::mlir::MLIRContext *);
+ static GPUParallelDimAttr blockZDim(::mlir::MLIRContext *);
+
+ // Factory functions by index (0=X, 1=Y, 2=Z)
+ static GPUParallelDimAttr threadDim(::mlir::MLIRContext *, unsigned index);
+ static GPUParallelDimAttr blockDim(::mlir::MLIRContext *, unsigned index);
+
+ // Ordering (for nesting validation: seq < thread < block)
+ int getOrder() const;
+
+ // Get the next higher/lower parallel dimension in the hierarchy
+ GPUParallelDimAttr getOneHigher() const;
+ GPUParallelDimAttr getOneLower() const;
+ }];
+ let hasCustomAssemblyFormat = 1;
+}
+
+def OpenACC_GPUParallelDimsAttr : OpenACC_Attr<"GPUParallelDims", "par_dims"> {
+ let summary = "List of GPU parallel dimensions for use in OpenACC parallelism assignment.";
+ let description = [{
+ Ordered list of GPU parallel dimensions. Used for privatization scope, barrier
+ placement, and loop work-sharing in OpenACC codegen.
+ }];
+ let parameters = (ins ArrayRefParameter<"::mlir::acc::GPUParallelDimAttr">:$array);
+ let extraClassDeclaration = [{
+ // Factory for sequential execution
+ static GPUParallelDimsAttr seq(::mlir::MLIRContext *);
+
+ // Check if this represents sequential execution (single seq dimension)
+ bool isSeq() const;
+
+ // Check if this represents parallel execution (not sequential)
+ bool isParallel() const;
+
+ // Check if this has multiple dimensions
+ bool isMultiDim() const;
+
+ // GPU dimension predicates
+ bool hasAnyBlockLevel() const;
+ bool hasOnlyBlockLevel() const;
+ bool hasOnlyThreadYLevel() const;
+ bool hasOnlyThreadXLevel() const;
+ }];
+ let hasCustomAssemblyFormat = 1;
+}
+
+#endif // OPENACCCG_ATTRIBUTES
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index b64cf50b4dba7..ebb73db77d1e4 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -3640,6 +3640,7 @@ def OpenACC_WaitOp : OpenACC_Op<"wait", [AttrSizedOperandSegments]> {
let hasVerifier = 1;
}
+include "mlir/Dialect/OpenACC/OpenACCCGAttributes.td"
include "mlir/Dialect/OpenACC/OpenACCCGOps.td"
#endif // OPENACC_OPS
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp
index 2753750128699..d2c1da8229433 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp
@@ -14,10 +14,14 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Region.h"
#include "mlir/Interfaces/ControlFlowInterfaces.h"
#include "mlir/Support/LogicalResult.h"
+#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
using namespace mlir;
@@ -119,6 +123,80 @@ static void addResultEffect(
effects.emplace_back(EffectTy::get(), mlir::cast<mlir::OpResult>(result));
}
+static int64_t gpuProcessorIndex(gpu::Processor p) {
+ switch (p) {
+ case gpu::Processor::Sequential:
+ return 0;
+ case gpu::Processor::ThreadX:
+ return 1;
+ case gpu::Processor::ThreadY:
+ return 2;
+ case gpu::Processor::ThreadZ:
+ return 3;
+ case gpu::Processor::BlockX:
+ return 4;
+ case gpu::Processor::BlockY:
+ return 5;
+ case gpu::Processor::BlockZ:
+ return 6;
+ }
+ llvm_unreachable("unhandled gpu::Processor");
+}
+
+static gpu::Processor indexToGpuProcessor(int64_t idx) {
+ switch (idx) {
+ case 0:
+ return gpu::Processor::Sequential;
+ case 1:
+ return gpu::Processor::ThreadX;
+ case 2:
+ return gpu::Processor::ThreadY;
+ case 3:
+ return gpu::Processor::ThreadZ;
+ case 4:
+ return gpu::Processor::BlockX;
+ case 5:
+ return gpu::Processor::BlockY;
+ case 6:
+ return gpu::Processor::BlockZ;
+ default:
+ return gpu::Processor::Sequential;
+ }
+}
+
+static GPUParallelDimAttr intToParDim(MLIRContext *context, int64_t dimInt) {
+ return GPUParallelDimAttr::get(
+ context, IntegerAttr::get(IndexType::get(context), dimInt));
+}
+
+static GPUParallelDimAttr processorParDim(MLIRContext *context,
+ gpu::Processor proc) {
+ return GPUParallelDimAttr::get(
+ context,
+ IntegerAttr::get(IndexType::get(context), gpuProcessorIndex(proc)));
+}
+
+static ParseResult parseProcessorValue(AsmParser &parser,
+ GPUParallelDimAttr &dim) {
+ std::string keyword;
+ llvm::SMLoc loc = parser.getCurrentLocation();
+ if (failed(parser.parseKeywordOrString(&keyword)))
+ return failure();
+ auto maybeProcessor = gpu::symbolizeProcessor(keyword);
+ if (!maybeProcessor) {
+ return parser.emitError(loc)
+ << "expected one of ::mlir::gpu::Processor enum names";
+ }
+ dim = intToParDim(parser.getContext(), gpuProcessorIndex(*maybeProcessor));
+ return success();
+}
+
+static void printProcessorValue(AsmPrinter &printer,
+ const GPUParallelDimAttr &attr) {
+ gpu::Processor processor = indexToGpuProcessor(attr.getValue().getInt());
+ printer << gpu::stringifyProcessor(processor);
+}
+
} // namespace
//===----------------------------------------------------------------------===//
@@ -184,3 +262,208 @@ void ReductionCombineOp::getEffects(
effects.emplace_back(MemoryEffects::Write::get(), &getDestMemrefMutable(),
SideEffects::DefaultResource::get());
}
+
+//===----------------------------------------------------------------------===//
+// GPUParallelDimAttr
+//===----------------------------------------------------------------------===//
+
+GPUParallelDimAttr GPUParallelDimAttr::get(MLIRContext *context,
+ gpu::Processor proc) {
+ return processorParDim(context, proc);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::seqDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::Sequential);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::threadXDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::ThreadX);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::threadYDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::ThreadY);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::threadZDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::ThreadZ);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::blockXDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::BlockX);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::blockYDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::BlockY);
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::blockZDim(MLIRContext *context) {
+ return processorParDim(context, gpu::Processor::BlockZ);
+}
+
+Attribute GPUParallelDimAttr::parse(AsmParser &parser, Type type) {
+ GPUParallelDimAttr dim;
+ if (parser.parseLess() || parseProcessorValue(parser, dim) ||
+ parser.parseGreater()) {
+ parser.emitError(parser.getCurrentLocation(),
+ "expected format `<` processor_name `>`");
+ return {};
+ }
+ return dim;
+}
+
+void GPUParallelDimAttr::print(AsmPrinter &printer) const {
+ printer << "<";
+ printProcessorValue(printer, *this);
+ printer << ">";
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::threadDim(MLIRContext *context,
+ unsigned index) {
+ assert(index <= 2 && "thread dimension index must be 0, 1, or 2");
+ switch (index) {
+ case 0:
+ return threadXDim(context);
+ case 1:
+ return threadYDim(context);
+ case 2:
+ return threadZDim(context);
+ }
+ llvm_unreachable("validated thread dimension index");
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::blockDim(MLIRContext *context,
+ unsigned index) {
+ assert(index <= 2 && "block dimension index must be 0, 1, or 2");
+ switch (index) {
+ case 0:
+ return blockXDim(context);
+ case 1:
+ return blockYDim(context);
+ case 2:
+ return blockZDim(context);
+ }
+ llvm_unreachable("validated block dimension index");
+}
+
+gpu::Processor GPUParallelDimAttr::getProcessor() const {
+ return indexToGpuProcessor(getValue().getInt());
+}
+
+int GPUParallelDimAttr::getOrder() const {
+ return gpuProcessorIndex(getProcessor());
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::getOneHigher() const {
+ int order = getOrder();
+ if (order >= 6) // BlockZ is the highest
+ return *this;
+ return get(getContext(), indexToGpuProcessor(order + 1));
+}
+
+GPUParallelDimAttr GPUParallelDimAttr::getOneLower() const {
+ int order = getOrder();
+ if (order <= 0) // Sequential is the lowest
+ return *this;
+ return get(getContext(), indexToGpuProcessor(order - 1));
+}
+
+bool GPUParallelDimAttr::isSeq() const {
+ return getProcessor() == gpu::Processor::Sequential;
+}
+bool GPUParallelDimAttr::isThreadX() const {
+ return getProcessor() == gpu::Processor::ThreadX;
+}
+bool GPUParallelDimAttr::isThreadY() const {
+ return getProcessor() == gpu::Processor::ThreadY;
+}
+bool GPUParallelDimAttr::isThreadZ() const {
+ return getProcessor() == gpu::Processor::ThreadZ;
+}
+bool GPUParallelDimAttr::isBlockX() const {
+ return getProcessor() == gpu::Processor::BlockX;
+}
+bool GPUParallelDimAttr::isBlockY() const {
+ return getProcessor() == gpu::Processor::BlockY;
+}
+bool GPUParallelDimAttr::isBlockZ() const {
+ return getProcessor() == gpu::Processor::BlockZ;
+}
+bool GPUParallelDimAttr::isAnyThread() const {
+ return isThreadX() || isThreadY() || isThreadZ();
+}
+bool GPUParallelDimAttr::isAnyBlock() const {
+ return isBlockX() || isBlockY() || isBlockZ();
+}
+
+//===----------------------------------------------------------------------===//
+// GPUParallelDimsAttr
+//===----------------------------------------------------------------------===//
+
+GPUParallelDimsAttr GPUParallelDimsAttr::seq(MLIRContext *ctx) {
+ return GPUParallelDimsAttr::get(ctx, {GPUParallelDimAttr::seqDim(ctx)});
+}
+
+bool GPUParallelDimsAttr::isSeq() const {
+ assert(!getArray().empty() && "no par_dims found");
+ if (getArray().size() == 1) {
+ auto parDim = dyn_cast<GPUParallelDimAttr>(getArray()[0]);
+ assert(parDim && "expected GPUParallelDimAttr");
+ return parDim.isSeq();
+ }
+ return false;
+}
+
+bool GPUParallelDimsAttr::isParallel() const { return !isSeq(); }
+
+bool GPUParallelDimsAttr::isMultiDim() const { return getArray().size() > 1; }
+
+bool GPUParallelDimsAttr::hasAnyBlockLevel() const {
+ return llvm::any_of(
+ getArray(), [](const GPUParallelDimAttr &p) { return p.isAnyBlock(); });
+}
+
+bool GPUParallelDimsAttr::hasOnlyBlockLevel() const {
+ return !getArray().empty() &&
+ llvm::all_of(getArray(), [](const GPUParallelDimAttr &p) {
+ return p.isAnyBlock();
+ });
+}
+
+bool GPUParallelDimsAttr::hasOnlyThreadYLevel() const {
+ return !getArray().empty() &&
+ llvm::all_of(getArray(), [](const GPUParallelDimAttr &p) {
+ return p.isThreadY();
+ });
+}
+
+bool GPUParallelDimsAttr::hasOnlyThreadXLevel() const {
+ return !getArray().empty() &&
+ llvm::all_of(getArray(), [](const GPUParallelDimAttr &p) {
+ return p.isThreadX();
+ });
+}
+
+Attribute GPUParallelDimsAttr::parse(AsmParser &parser, Type type) {
+ auto delimiter = AsmParser::Delimiter::Square;
+ SmallVector<GPUParallelDimAttr> parDims;
+ auto parseParDim = [&]() -> ParseResult {
+ GPUParallelDimAttr dim;
+ if (parseProcessorValue(parser, dim))
+ return failure();
+ parDims.push_back(dim);
+ return success();
+ };
+ if (parser.parseCommaSeparatedList(delimiter, parseParDim,
+ "list of OpenACC GPU parallel dimensions"))
+ return {};
+ return GPUParallelDimsAttr::get(parser.getContext(), parDims);
+}
+
+void GPUParallelDimsAttr::print(AsmPrinter &printer) const {
+ printer << "[";
+ llvm::interleaveComma(getArray(), printer,
+ [&printer](const GPUParallelDimAttr &p) {
+ printProcessorValue(printer, p);
+ });
+ printer << "]";
+}
diff --git a/mlir/test/Dialect/OpenACC/invalid-cg.mlir b/mlir/test/Dialect/OpenACC/invalid-cg.mlir
new file mode 100644
index 0000000000000..bc2408ceafe8a
--- /dev/null
+++ b/mlir/test/Dialect/OpenACC/invalid-cg.mlir
@@ -0,0 +1,21 @@
+// RUN: mlir-opt -split-input-file -verify-diagnostics %s
+
+// -----
+
+%c0 = arith.constant 0 : index
+%c1 = arith.constant 1 : index
+%c4 = arith.constant 4 : index
+scf.parallel (%iv) = (%c0) to (%c4) step (%c1) {
+ scf.reduce
+// expected-error at +1 {{expected one of ::mlir::gpu::Processor enum names}}
+} {acc.par_dims = #acc<par_dims[thread_x, bad_dim]>}
+
+// -----
+
+%c0_2 = arith.constant 0 : index
+%c1_2 = arith.constant 1 : index
+%c4_2 = arith.constant 4 : index
+scf.parallel (%iv) = (%c0_2) to (%c4_2) step (%c1_2) {
+ scf.reduce
+// expected-error at +1 {{expected one of ::mlir::gpu::Processor enum names}}
+} {acc.par_dims = #acc<par_dims[gang]>}
diff --git a/mlir/test/Dialect/OpenACC/ops-cg.mlir b/mlir/test/Dialect/OpenACC/ops-cg.mlir
new file mode 100644
index 0000000000000..e6453da21ed79
--- /dev/null
+++ b/mlir/test/Dialect/OpenACC/ops-cg.mlir
@@ -0,0 +1,79 @@
+// RUN: mlir-opt -split-input-file %s | FileCheck %s --check-prefixes=CHECK
+// Verify the printed output can be parsed.
+// RUN: mlir-opt -split-input-file %s | mlir-opt -split-input-file | FileCheck %s --check-prefixes=CHECK
+// Verify the generic form can be parsed.
+// RUN: mlir-opt -split-input-file -mlir-print-op-generic %s | mlir-opt -split-input-file | FileCheck %s --check-prefixes=CHECK
+
+// -----
+
+// CHECK-LABEL: func @par_dim_sequential
+func.func @par_dim_sequential() {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ scf.parallel (%iv) = (%c0) to (%c4) step (%c1) {
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[sequential]>}
+ return
+}
+// CHECK: scf.parallel
+// CHECK: } {acc.par_dims = #acc<par_dims[sequential]>}
+
+// -----
+
+// CHECK-LABEL: func @par_dim_single_thread_x
+func.func @par_dim_single_thread_x() {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c128 = arith.constant 128 : index
+ scf.parallel (%iv) = (%c0) to (%c128) step (%c1) {
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[thread_x]>}
+ return
+}
+// CHECK: scf.parallel
+// CHECK: } {acc.par_dims = #acc<par_dims[thread_x]>}
+
+// -----
+
+// CHECK-LABEL: func @par_dims_block_thread
+func.func @par_dims_block_thread() {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c8 = arith.constant 8 : index
+ %c128 = arith.constant 128 : index
+ scf.parallel (%i, %j) = (%c0, %c0) to (%c8, %c128) step (%c1, %c1) {
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_x, thread_x]>}
+ return
+}
+// CHECK: scf.parallel
+// CHECK: } {acc.par_dims = #acc<par_dims[block_x, thread_x]>}
+
+// -----
+
+// All GPU parallel dimensions (par_dim values) in par_dims list
+func.func @par_dims_all_dims() {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ scf.parallel (%iv) = (%c0) to (%c1) step (%c1) {
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_x, block_y, block_z, thread_x, thread_y, thread_z]>}
+ return
+}
+// CHECK: acc.par_dims = #acc<par_dims[block_x, block_y, block_z, thread_x, thread_y, thread_z]>
+
+// -----
+
+// 2D grid: block_y and thread_y
+func.func @par_dims_2d_grid() {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c32 = arith.constant 32 : index
+ scf.parallel (%i, %j) = (%c0, %c0) to (%c4, %c32) step (%c1, %c1) {
+ scf.reduce
+ } {acc.par_dims = #acc<par_dims[block_y, thread_y]>}
+ return
+}
+// CHECK: acc.par_dims = #acc<par_dims[block_y, thread_y]>
More information about the Mlir-commits
mailing list