[Mlir-commits] [mlir] 812c6f8 - [mlir][acc] Add parallelism mapping policy interface (#182890)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Mon Feb 23 09:52:51 PST 2026
Author: Razvan Lupusoru
Date: 2026-02-23T09:52:46-08:00
New Revision: 812c6f8305d6e94ab0077e06bf0ea156fb81d86a
URL: https://github.com/llvm/llvm-project/commit/812c6f8305d6e94ab0077e06bf0ea156fb81d86a
DIFF: https://github.com/llvm/llvm-project/commit/812c6f8305d6e94ab0077e06bf0ea156fb81d86a.diff
LOG: [mlir][acc] Add parallelism mapping policy interface (#182890)
Add a header that defines the interface for mapping OpenACC parallelism
levels (gang, worker, vector) to target-specific parallel dimension
attributes. Alongside this,
DefaultACCToGPUMappingPolicy is introduced for an initial implementation
of ACC parallelism to GPU mapping.
Added:
mlir/include/mlir/Dialect/OpenACC/OpenACCParMapping.h
Modified:
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCParMapping.h b/mlir/include/mlir/Dialect/OpenACC/OpenACCParMapping.h
new file mode 100644
index 0000000000000..09b700fe8ad40
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCParMapping.h
@@ -0,0 +1,164 @@
+//===- OpenACCParMapping.h - OpenACC Parallelism Mapping -------*- C++ -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines the interface for mapping OpenACC parallelism levels
+// (gang, worker, vector) to target-specific parallel dimension attributes.
+//
+// Users can provide custom implementations of ACCParMappingPolicy to
+// support
diff erent mapping strategies and target attributes.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_OPENACC_OPENACCPARMAPPING_H_
+#define MLIR_DIALECT_OPENACC_OPENACCPARMAPPING_H_
+
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/MLIRContext.h"
+#include "llvm/Support/ErrorHandling.h"
+
+namespace mlir {
+namespace acc {
+
+//===----------------------------------------------------------------------===//
+// Utility functions
+//===----------------------------------------------------------------------===//
+
+/// Convert a gang dimension value (1, 2, or 3) to the corresponding ParLevel.
+/// Asserts if the value is not a valid gang dimension.
+inline ParLevel getGangParLevel(int64_t gangDimValue) {
+ assert((gangDimValue >= 1 && gangDimValue <= 3) &&
+ "gang dimension must be 1, 2, or 3");
+ switch (gangDimValue) {
+ case 1:
+ return ParLevel::gang_dim1;
+ case 2:
+ return ParLevel::gang_dim2;
+ case 3:
+ return ParLevel::gang_dim3;
+ }
+ llvm_unreachable("validated gang dimension");
+}
+
+//===----------------------------------------------------------------------===//
+// ACCParMappingPolicy
+//===----------------------------------------------------------------------===//
+
+/// Policy class that defines how OpenACC parallelism levels map to
+/// target-specific parallel dimension attributes. Implementations provide the
+/// actual mapping.
+///
+/// Template parameter ParDimAttrT specifies the attribute type returned by
+/// the mapping functions (e.g., mlir::acc::GPUParallelDimAttr for GPU targets).
+///
+/// This policy allows
diff erent mapping strategies:
+/// - Standard GPU mapping (gang->block, worker->threadY, vector->threadX)
+/// - Custom mappings for specific targets or optimization strategies
+///
+/// Pass an implementation to functions that need to perform the mapping.
+template <typename ParDimAttrT>
+class ACCParMappingPolicy {
+public:
+ virtual ~ACCParMappingPolicy() = default;
+
+ /// Map an OpenACC parallelism level to target dimension.
+ /// @param ctx The MLIR context
+ /// @param level The OpenACC parallelism level (gang_dim1, gang_dim2,
+ /// gang_dim3, worker, vector, or seq)
+ /// @return The corresponding parallel dimension attribute
+ virtual ParDimAttrT map(MLIRContext *ctx, ParLevel level) const = 0;
+
+ /// Convenience methods for specific parallelism levels.
+ ParDimAttrT gangDim(MLIRContext *ctx, ParLevel level) const {
+ assert((level == ParLevel::gang_dim1 || level == ParLevel::gang_dim2 ||
+ level == ParLevel::gang_dim3) &&
+ "gangDim requires a gang parallelism level");
+ return map(ctx, level);
+ }
+ ParDimAttrT workerDim(MLIRContext *ctx) const {
+ return map(ctx, ParLevel::worker);
+ }
+ ParDimAttrT vectorDim(MLIRContext *ctx) const {
+ return map(ctx, ParLevel::vector);
+ }
+ ParDimAttrT seqDim(MLIRContext *ctx) const { return map(ctx, ParLevel::seq); }
+
+ //===--------------------------------------------------------------------===//
+ // Predicate methods - check if an attribute matches a parallelism level
+ //===--------------------------------------------------------------------===//
+
+ /// Check if the attribute represents vector parallelism.
+ virtual bool isVector(ParDimAttrT attr) const = 0;
+
+ /// Check if the attribute represents worker parallelism.
+ virtual bool isWorker(ParDimAttrT attr) const = 0;
+
+ /// Check if the attribute represents gang parallelism (any gang dimension).
+ virtual bool isGang(ParDimAttrT attr) const = 0;
+
+ /// Check if the attribute represents sequential execution.
+ virtual bool isSeq(ParDimAttrT attr) const = 0;
+};
+
+//===----------------------------------------------------------------------===//
+// DefaultACCToGPUMappingPolicy
+//===----------------------------------------------------------------------===//
+
+/// Default policy that provides the standard GPU mapping:
+/// gang(dim:1) -> BlockX (gridDim.x / blockIdx.x)
+/// gang(dim:2) -> BlockY (gridDim.y / blockIdx.y)
+/// gang(dim:3) -> BlockZ (gridDim.z / blockIdx.z)
+/// worker -> ThreadY (blockDim.y / threadIdx.y)
+/// vector -> ThreadX (blockDim.x / threadIdx.x)
+/// seq -> Sequential
+class DefaultACCToGPUMappingPolicy
+ : public ACCParMappingPolicy<mlir::acc::GPUParallelDimAttr> {
+public:
+ mlir::acc::GPUParallelDimAttr map(MLIRContext *ctx,
+ ParLevel level) const override {
+ switch (level) {
+ case ParLevel::gang_dim1:
+ return mlir::acc::GPUParallelDimAttr::blockXDim(ctx);
+ case ParLevel::gang_dim2:
+ return mlir::acc::GPUParallelDimAttr::blockYDim(ctx);
+ case ParLevel::gang_dim3:
+ return mlir::acc::GPUParallelDimAttr::blockZDim(ctx);
+ case ParLevel::worker:
+ return mlir::acc::GPUParallelDimAttr::threadYDim(ctx);
+ case ParLevel::vector:
+ return mlir::acc::GPUParallelDimAttr::threadXDim(ctx);
+ case ParLevel::seq:
+ return mlir::acc::GPUParallelDimAttr::seqDim(ctx);
+ }
+ llvm_unreachable("Unknown ParLevel");
+ }
+
+ bool isVector(mlir::acc::GPUParallelDimAttr attr) const override {
+ return attr.isThreadX();
+ }
+
+ bool isWorker(mlir::acc::GPUParallelDimAttr attr) const override {
+ return attr.isThreadY();
+ }
+
+ bool isGang(mlir::acc::GPUParallelDimAttr attr) const override {
+ return attr.isAnyBlock();
+ }
+
+ bool isSeq(mlir::acc::GPUParallelDimAttr attr) const override {
+ return attr.isSeq();
+ }
+};
+
+/// Type alias for the GPU-specific mapping policy
+using ACCToGPUMappingPolicy =
+ ACCParMappingPolicy<mlir::acc::GPUParallelDimAttr>;
+
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_OPENACCPARMAPPING_H_
More information about the Mlir-commits
mailing list