[OpenACC][CIR] Implement 'gang' lowering for 'loop' #138968
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This clause requires an entire additional collection to keep track of the gang 'kind' or 'type'. That work is maintained in the OpenACC dialect functions. Otherwise, this is effectively the same as the worker/vectors.
llvmbot
added
clang
|
@llvm/pr-subscribers-openacc @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThis clause requires an entire additional collection to keep track of the gang 'kind' or 'type'. That work is maintained in the OpenACC dialect functions. Otherwise, this is effectively the same as the worker/vectors. Full diff: https://github.com/llvm/llvm-project/pull/138968.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
index ef4f64a167742..f7fc3bb5c656c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
@@ -107,6 +107,18 @@ class OpenACCClauseCIREmitter final
.CaseLower("radeon", mlir::acc::DeviceType::Radeon);
}
+ mlir::acc::GangArgType decodeGangType(OpenACCGangKind GK) {
+ switch (GK) {
+ case OpenACCGangKind::Num:
+ return mlir::acc::GangArgType::Num;
+ case OpenACCGangKind::Dim:
+ return mlir::acc::GangArgType::Dim;
+ case OpenACCGangKind::Static:
+ return mlir::acc::GangArgType::Static;
+ }
+ llvm_unreachable("unknown gang kind");
+ }
+
public:
OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
CIRGen::CIRGenBuilderTy &builder,
@@ -424,6 +436,42 @@ class OpenACCClauseCIREmitter final
return clauseNotImplemented(clause);
}
}
+
+ void VisitGangClause(const OpenACCGangClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ if (clause.getNumExprs() == 0) {
+ operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ llvm::SmallVector<mlir::Value> values;
+ llvm::SmallVector<mlir::acc::GangArgType> argTypes;
+ for (unsigned I = 0; I < clause.getNumExprs(); ++I) {
+ auto [kind, expr] = clause.getExpr(I);
+ mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
+ argTypes.push_back(decodeGangType(kind));
+ if (kind == OpenACCGangKind::Dim) {
+ llvm::APInt curValue =
+ expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
+ // The value is 1, 2, or 3, but the type isn't necessarily smaller
+ // than 64.
+ curValue = curValue.sextOrTrunc(64);
+ values.push_back(
+ createConstantInt(exprLoc, 64, curValue.getSExtValue()));
+ } else if (isa<OpenACCAsteriskSizeExpr>(expr)) {
+ values.push_back(createConstantInt(exprLoc, 64, -1));
+ } else {
+ values.push_back(createIntExpr(expr));
+ }
+ }
+
+ operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
+ argTypes, values);
+ }
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
};
template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/loop.cpp b/clang/test/CIR/CodeGenOpenACC/loop.cpp
index d636d1b37d969..4b7a7e7366323 100644
--- a/clang/test/CIR/CodeGenOpenACC/loop.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/loop.cpp
@@ -323,4 +323,73 @@ extern "C" void acc_loop(int *A, int *B, int *C, int N) {
// CHECK: acc.yield
// CHECK-NEXT: } loc
}
+
+#pragma acc parallel
+ // CHECK: acc.parallel {
+ {
+#pragma acc loop gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang device_type(nvidia) gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang([#acc.device_type<none>, #acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(dim:1) device_type(nvidia) gang(dim:2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({dim=%[[ONE_CONST]] : i64}, {dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, dim: 1) device_type(nvidia, radeon) gang(static:*, dim : 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, dim=%[[ONE_CONST]] : i64}, {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>], {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
+#pragma acc kernels
+ // CHECK: acc.kernels {
+ {
+#pragma acc loop gang(num:N) device_type(nvidia, radeon) gang(num:N)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({num=%[[N_CONV]] : si32}, {num=%[[N_CONV2]] : si32} [#acc.device_type<nvidia>], {num=%[[N_CONV2]] : si32} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N) device_type(nvidia) gang(static:*)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32}, {static=%[[STAR_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, num: N + 1) device_type(nvidia) gang(static:*, num : N + 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[CIR_ONE_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO:.*]] = cir.binop(add, %[[N_LOAD3]], %[[CIR_TWO_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_TWO]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, num=%[[N_PLUS_ONE_CONV]] : si32}, {static=%[[STAR_CONST]] : i64, num=%[[N_PLUS_TWO_CONV]] : si32} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index ca564037fad19..5d5add6318e06 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2231,6 +2231,16 @@ def OpenACC_LoopOp : OpenACC_Op<"loop",
// device_types. This is for the case where there is no expression specified
// in a 'worker'.
void addEmptyWorker(MLIRContext *, llvm::ArrayRef<DeviceType>);
+
+ // Adds a collection of operands for a 'gang' clause that has various types
+ // corresponding to each operand.
+ void addGangOperands(MLIRContext *, llvm::ArrayRef<DeviceType>,
+ llvm::ArrayRef<GangArgType>, mlir::ValueRange);
+
+ // Add an empty value to the 'gang' list with a current list of
+ // device_types. This is for the case where there is no expression specified
+ // in a 'gang'.
+ void addEmptyGang(MLIRContext *, llvm::ArrayRef<DeviceType>);
}];
let hasCustomAssemblyFormat = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 9f4645a4a7ca8..561f1aef7d5c9 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -2748,6 +2748,49 @@ void acc::LoopOp::addEmptyWorker(
effectiveDeviceTypes));
}
+void acc::LoopOp::addEmptyGang(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setGangAttr(addDeviceTypeAffectedOperandHelper(context, getGangAttr(),
+ effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addGangOperands(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes,
+ llvm::ArrayRef<GangArgType> argTypes, mlir::ValueRange values) {
+ llvm::SmallVector<int32_t> segments;
+ if (getGangOperandsSegments())
+ llvm::copy(*getGangOperandsSegments(), std::back_inserter(segments));
+
+ unsigned beforeCount = segments.size();
+
+ setGangOperandsDeviceTypeAttr(addDeviceTypeAffectedOperandHelper(
+ context, getGangOperandsDeviceTypeAttr(), effectiveDeviceTypes, values,
+ getGangOperandsMutable(), segments));
+
+ setGangOperandsSegments(segments);
+
+ // This is a bit of extra work to make sure we update the 'types' correctly by
+ // adding to the types collection the correct number of times. We could
+ // potentially add something similar to the
+ // addDeviceTypeAffectedOperandHelper, but it seems that would be pretty
+ // excessive for a one-off case.
+ unsigned numAdded = segments.size() - beforeCount;
+
+ if (numAdded > 0) {
+ llvm::SmallVector<mlir::Attribute> gangTypes;
+ if (getGangOperandsArgTypeAttr())
+ llvm::copy(getGangOperandsArgTypeAttr(), std::back_inserter(gangTypes));
+
+ for (unsigned I = 0; I < numAdded; ++I)
+ llvm::transform(argTypes, std::back_inserter(gangTypes),
+ [=](mlir::acc::GangArgType gangTy) {
+ return mlir::acc::GangArgTypeAttr::get(context, gangTy);
+ });
+
+ setGangOperandsArgTypeAttr(mlir::ArrayAttr::get(context, gangTypes));
+ }
+}
+
//===----------------------------------------------------------------------===//
// DataOp
//===----------------------------------------------------------------------===//
|
|
@llvm/pr-subscribers-mlir-openacc Author: Erich Keane (erichkeane) ChangesThis clause requires an entire additional collection to keep track of the gang 'kind' or 'type'. That work is maintained in the OpenACC dialect functions. Otherwise, this is effectively the same as the worker/vectors. Full diff: https://github.com/llvm/llvm-project/pull/138968.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
index ef4f64a167742..f7fc3bb5c656c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
@@ -107,6 +107,18 @@ class OpenACCClauseCIREmitter final
.CaseLower("radeon", mlir::acc::DeviceType::Radeon);
}
+ mlir::acc::GangArgType decodeGangType(OpenACCGangKind GK) {
+ switch (GK) {
+ case OpenACCGangKind::Num:
+ return mlir::acc::GangArgType::Num;
+ case OpenACCGangKind::Dim:
+ return mlir::acc::GangArgType::Dim;
+ case OpenACCGangKind::Static:
+ return mlir::acc::GangArgType::Static;
+ }
+ llvm_unreachable("unknown gang kind");
+ }
+
public:
OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
CIRGen::CIRGenBuilderTy &builder,
@@ -424,6 +436,42 @@ class OpenACCClauseCIREmitter final
return clauseNotImplemented(clause);
}
}
+
+ void VisitGangClause(const OpenACCGangClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ if (clause.getNumExprs() == 0) {
+ operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ llvm::SmallVector<mlir::Value> values;
+ llvm::SmallVector<mlir::acc::GangArgType> argTypes;
+ for (unsigned I = 0; I < clause.getNumExprs(); ++I) {
+ auto [kind, expr] = clause.getExpr(I);
+ mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
+ argTypes.push_back(decodeGangType(kind));
+ if (kind == OpenACCGangKind::Dim) {
+ llvm::APInt curValue =
+ expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
+ // The value is 1, 2, or 3, but the type isn't necessarily smaller
+ // than 64.
+ curValue = curValue.sextOrTrunc(64);
+ values.push_back(
+ createConstantInt(exprLoc, 64, curValue.getSExtValue()));
+ } else if (isa<OpenACCAsteriskSizeExpr>(expr)) {
+ values.push_back(createConstantInt(exprLoc, 64, -1));
+ } else {
+ values.push_back(createIntExpr(expr));
+ }
+ }
+
+ operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
+ argTypes, values);
+ }
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
};
template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/loop.cpp b/clang/test/CIR/CodeGenOpenACC/loop.cpp
index d636d1b37d969..4b7a7e7366323 100644
--- a/clang/test/CIR/CodeGenOpenACC/loop.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/loop.cpp
@@ -323,4 +323,73 @@ extern "C" void acc_loop(int *A, int *B, int *C, int N) {
// CHECK: acc.yield
// CHECK-NEXT: } loc
}
+
+#pragma acc parallel
+ // CHECK: acc.parallel {
+ {
+#pragma acc loop gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang device_type(nvidia) gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang([#acc.device_type<none>, #acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(dim:1) device_type(nvidia) gang(dim:2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({dim=%[[ONE_CONST]] : i64}, {dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, dim: 1) device_type(nvidia, radeon) gang(static:*, dim : 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, dim=%[[ONE_CONST]] : i64}, {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>], {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
+#pragma acc kernels
+ // CHECK: acc.kernels {
+ {
+#pragma acc loop gang(num:N) device_type(nvidia, radeon) gang(num:N)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({num=%[[N_CONV]] : si32}, {num=%[[N_CONV2]] : si32} [#acc.device_type<nvidia>], {num=%[[N_CONV2]] : si32} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N) device_type(nvidia) gang(static:*)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32}, {static=%[[STAR_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, num: N + 1) device_type(nvidia) gang(static:*, num : N + 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[CIR_ONE_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO:.*]] = cir.binop(add, %[[N_LOAD3]], %[[CIR_TWO_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_TWO]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, num=%[[N_PLUS_ONE_CONV]] : si32}, {static=%[[STAR_CONST]] : i64, num=%[[N_PLUS_TWO_CONV]] : si32} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index ca564037fad19..5d5add6318e06 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2231,6 +2231,16 @@ def OpenACC_LoopOp : OpenACC_Op<"loop",
// device_types. This is for the case where there is no expression specified
// in a 'worker'.
void addEmptyWorker(MLIRContext *, llvm::ArrayRef<DeviceType>);
+
+ // Adds a collection of operands for a 'gang' clause that has various types
+ // corresponding to each operand.
+ void addGangOperands(MLIRContext *, llvm::ArrayRef<DeviceType>,
+ llvm::ArrayRef<GangArgType>, mlir::ValueRange);
+
+ // Add an empty value to the 'gang' list with a current list of
+ // device_types. This is for the case where there is no expression specified
+ // in a 'gang'.
+ void addEmptyGang(MLIRContext *, llvm::ArrayRef<DeviceType>);
}];
let hasCustomAssemblyFormat = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 9f4645a4a7ca8..561f1aef7d5c9 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -2748,6 +2748,49 @@ void acc::LoopOp::addEmptyWorker(
effectiveDeviceTypes));
}
+void acc::LoopOp::addEmptyGang(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setGangAttr(addDeviceTypeAffectedOperandHelper(context, getGangAttr(),
+ effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addGangOperands(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes,
+ llvm::ArrayRef<GangArgType> argTypes, mlir::ValueRange values) {
+ llvm::SmallVector<int32_t> segments;
+ if (getGangOperandsSegments())
+ llvm::copy(*getGangOperandsSegments(), std::back_inserter(segments));
+
+ unsigned beforeCount = segments.size();
+
+ setGangOperandsDeviceTypeAttr(addDeviceTypeAffectedOperandHelper(
+ context, getGangOperandsDeviceTypeAttr(), effectiveDeviceTypes, values,
+ getGangOperandsMutable(), segments));
+
+ setGangOperandsSegments(segments);
+
+ // This is a bit of extra work to make sure we update the 'types' correctly by
+ // adding to the types collection the correct number of times. We could
+ // potentially add something similar to the
+ // addDeviceTypeAffectedOperandHelper, but it seems that would be pretty
+ // excessive for a one-off case.
+ unsigned numAdded = segments.size() - beforeCount;
+
+ if (numAdded > 0) {
+ llvm::SmallVector<mlir::Attribute> gangTypes;
+ if (getGangOperandsArgTypeAttr())
+ llvm::copy(getGangOperandsArgTypeAttr(), std::back_inserter(gangTypes));
+
+ for (unsigned I = 0; I < numAdded; ++I)
+ llvm::transform(argTypes, std::back_inserter(gangTypes),
+ [=](mlir::acc::GangArgType gangTy) {
+ return mlir::acc::GangArgTypeAttr::get(context, gangTy);
+ });
+
+ setGangOperandsArgTypeAttr(mlir::ArrayAttr::get(context, gangTypes));
+ }
+}
+
//===----------------------------------------------------------------------===//
// DataOp
//===----------------------------------------------------------------------===//
|
|
@llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesThis clause requires an entire additional collection to keep track of the gang 'kind' or 'type'. That work is maintained in the OpenACC dialect functions. Otherwise, this is effectively the same as the worker/vectors. Full diff: https://github.com/llvm/llvm-project/pull/138968.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
index ef4f64a167742..f7fc3bb5c656c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
@@ -107,6 +107,18 @@ class OpenACCClauseCIREmitter final
.CaseLower("radeon", mlir::acc::DeviceType::Radeon);
}
+ mlir::acc::GangArgType decodeGangType(OpenACCGangKind GK) {
+ switch (GK) {
+ case OpenACCGangKind::Num:
+ return mlir::acc::GangArgType::Num;
+ case OpenACCGangKind::Dim:
+ return mlir::acc::GangArgType::Dim;
+ case OpenACCGangKind::Static:
+ return mlir::acc::GangArgType::Static;
+ }
+ llvm_unreachable("unknown gang kind");
+ }
+
public:
OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
CIRGen::CIRGenBuilderTy &builder,
@@ -424,6 +436,42 @@ class OpenACCClauseCIREmitter final
return clauseNotImplemented(clause);
}
}
+
+ void VisitGangClause(const OpenACCGangClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ if (clause.getNumExprs() == 0) {
+ operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ llvm::SmallVector<mlir::Value> values;
+ llvm::SmallVector<mlir::acc::GangArgType> argTypes;
+ for (unsigned I = 0; I < clause.getNumExprs(); ++I) {
+ auto [kind, expr] = clause.getExpr(I);
+ mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
+ argTypes.push_back(decodeGangType(kind));
+ if (kind == OpenACCGangKind::Dim) {
+ llvm::APInt curValue =
+ expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
+ // The value is 1, 2, or 3, but the type isn't necessarily smaller
+ // than 64.
+ curValue = curValue.sextOrTrunc(64);
+ values.push_back(
+ createConstantInt(exprLoc, 64, curValue.getSExtValue()));
+ } else if (isa<OpenACCAsteriskSizeExpr>(expr)) {
+ values.push_back(createConstantInt(exprLoc, 64, -1));
+ } else {
+ values.push_back(createIntExpr(expr));
+ }
+ }
+
+ operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
+ argTypes, values);
+ }
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
};
template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/loop.cpp b/clang/test/CIR/CodeGenOpenACC/loop.cpp
index d636d1b37d969..4b7a7e7366323 100644
--- a/clang/test/CIR/CodeGenOpenACC/loop.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/loop.cpp
@@ -323,4 +323,73 @@ extern "C" void acc_loop(int *A, int *B, int *C, int N) {
// CHECK: acc.yield
// CHECK-NEXT: } loc
}
+
+#pragma acc parallel
+ // CHECK: acc.parallel {
+ {
+#pragma acc loop gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang device_type(nvidia) gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang([#acc.device_type<none>, #acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(dim:1) device_type(nvidia) gang(dim:2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({dim=%[[ONE_CONST]] : i64}, {dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, dim: 1) device_type(nvidia, radeon) gang(static:*, dim : 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, dim=%[[ONE_CONST]] : i64}, {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>], {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
+#pragma acc kernels
+ // CHECK: acc.kernels {
+ {
+#pragma acc loop gang(num:N) device_type(nvidia, radeon) gang(num:N)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({num=%[[N_CONV]] : si32}, {num=%[[N_CONV2]] : si32} [#acc.device_type<nvidia>], {num=%[[N_CONV2]] : si32} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N) device_type(nvidia) gang(static:*)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32}, {static=%[[STAR_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, num: N + 1) device_type(nvidia) gang(static:*, num : N + 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[CIR_ONE_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO:.*]] = cir.binop(add, %[[N_LOAD3]], %[[CIR_TWO_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_TWO]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, num=%[[N_PLUS_ONE_CONV]] : si32}, {static=%[[STAR_CONST]] : i64, num=%[[N_PLUS_TWO_CONV]] : si32} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index ca564037fad19..5d5add6318e06 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2231,6 +2231,16 @@ def OpenACC_LoopOp : OpenACC_Op<"loop",
// device_types. This is for the case where there is no expression specified
// in a 'worker'.
void addEmptyWorker(MLIRContext *, llvm::ArrayRef<DeviceType>);
+
+ // Adds a collection of operands for a 'gang' clause that has various types
+ // corresponding to each operand.
+ void addGangOperands(MLIRContext *, llvm::ArrayRef<DeviceType>,
+ llvm::ArrayRef<GangArgType>, mlir::ValueRange);
+
+ // Add an empty value to the 'gang' list with a current list of
+ // device_types. This is for the case where there is no expression specified
+ // in a 'gang'.
+ void addEmptyGang(MLIRContext *, llvm::ArrayRef<DeviceType>);
}];
let hasCustomAssemblyFormat = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 9f4645a4a7ca8..561f1aef7d5c9 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -2748,6 +2748,49 @@ void acc::LoopOp::addEmptyWorker(
effectiveDeviceTypes));
}
+void acc::LoopOp::addEmptyGang(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setGangAttr(addDeviceTypeAffectedOperandHelper(context, getGangAttr(),
+ effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addGangOperands(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes,
+ llvm::ArrayRef<GangArgType> argTypes, mlir::ValueRange values) {
+ llvm::SmallVector<int32_t> segments;
+ if (getGangOperandsSegments())
+ llvm::copy(*getGangOperandsSegments(), std::back_inserter(segments));
+
+ unsigned beforeCount = segments.size();
+
+ setGangOperandsDeviceTypeAttr(addDeviceTypeAffectedOperandHelper(
+ context, getGangOperandsDeviceTypeAttr(), effectiveDeviceTypes, values,
+ getGangOperandsMutable(), segments));
+
+ setGangOperandsSegments(segments);
+
+ // This is a bit of extra work to make sure we update the 'types' correctly by
+ // adding to the types collection the correct number of times. We could
+ // potentially add something similar to the
+ // addDeviceTypeAffectedOperandHelper, but it seems that would be pretty
+ // excessive for a one-off case.
+ unsigned numAdded = segments.size() - beforeCount;
+
+ if (numAdded > 0) {
+ llvm::SmallVector<mlir::Attribute> gangTypes;
+ if (getGangOperandsArgTypeAttr())
+ llvm::copy(getGangOperandsArgTypeAttr(), std::back_inserter(gangTypes));
+
+ for (unsigned I = 0; I < numAdded; ++I)
+ llvm::transform(argTypes, std::back_inserter(gangTypes),
+ [=](mlir::acc::GangArgType gangTy) {
+ return mlir::acc::GangArgTypeAttr::get(context, gangTy);
+ });
+
+ setGangOperandsArgTypeAttr(mlir::ArrayAttr::get(context, gangTypes));
+ }
+}
+
//===----------------------------------------------------------------------===//
// DataOp
//===----------------------------------------------------------------------===//
|
|
@llvm/pr-subscribers-mlir Author: Erich Keane (erichkeane) ChangesThis clause requires an entire additional collection to keep track of the gang 'kind' or 'type'. That work is maintained in the OpenACC dialect functions. Otherwise, this is effectively the same as the worker/vectors. Full diff: https://github.com/llvm/llvm-project/pull/138968.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
index ef4f64a167742..f7fc3bb5c656c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h
@@ -107,6 +107,18 @@ class OpenACCClauseCIREmitter final
.CaseLower("radeon", mlir::acc::DeviceType::Radeon);
}
+ mlir::acc::GangArgType decodeGangType(OpenACCGangKind GK) {
+ switch (GK) {
+ case OpenACCGangKind::Num:
+ return mlir::acc::GangArgType::Num;
+ case OpenACCGangKind::Dim:
+ return mlir::acc::GangArgType::Dim;
+ case OpenACCGangKind::Static:
+ return mlir::acc::GangArgType::Static;
+ }
+ llvm_unreachable("unknown gang kind");
+ }
+
public:
OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
CIRGen::CIRGenBuilderTy &builder,
@@ -424,6 +436,42 @@ class OpenACCClauseCIREmitter final
return clauseNotImplemented(clause);
}
}
+
+ void VisitGangClause(const OpenACCGangClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+ if (clause.getNumExprs() == 0) {
+ operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
+ } else {
+ llvm::SmallVector<mlir::Value> values;
+ llvm::SmallVector<mlir::acc::GangArgType> argTypes;
+ for (unsigned I = 0; I < clause.getNumExprs(); ++I) {
+ auto [kind, expr] = clause.getExpr(I);
+ mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
+ argTypes.push_back(decodeGangType(kind));
+ if (kind == OpenACCGangKind::Dim) {
+ llvm::APInt curValue =
+ expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
+ // The value is 1, 2, or 3, but the type isn't necessarily smaller
+ // than 64.
+ curValue = curValue.sextOrTrunc(64);
+ values.push_back(
+ createConstantInt(exprLoc, 64, curValue.getSExtValue()));
+ } else if (isa<OpenACCAsteriskSizeExpr>(expr)) {
+ values.push_back(createConstantInt(exprLoc, 64, -1));
+ } else {
+ values.push_back(createIntExpr(expr));
+ }
+ }
+
+ operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
+ argTypes, values);
+ }
+ } else {
+ // TODO: When we've implemented this for everything, switch this to an
+ // unreachable. Combined constructs remain.
+ return clauseNotImplemented(clause);
+ }
+ }
};
template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/loop.cpp b/clang/test/CIR/CodeGenOpenACC/loop.cpp
index d636d1b37d969..4b7a7e7366323 100644
--- a/clang/test/CIR/CodeGenOpenACC/loop.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/loop.cpp
@@ -323,4 +323,73 @@ extern "C" void acc_loop(int *A, int *B, int *C, int N) {
// CHECK: acc.yield
// CHECK-NEXT: } loc
}
+
+#pragma acc parallel
+ // CHECK: acc.parallel {
+ {
+#pragma acc loop gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang device_type(nvidia) gang
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: acc.loop gang([#acc.device_type<none>, #acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(dim:1) device_type(nvidia) gang(dim:2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({dim=%[[ONE_CONST]] : i64}, {dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, dim: 1) device_type(nvidia, radeon) gang(static:*, dim : 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, dim=%[[ONE_CONST]] : i64}, {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>], {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
+#pragma acc kernels
+ // CHECK: acc.kernels {
+ {
+#pragma acc loop gang(num:N) device_type(nvidia, radeon) gang(num:N)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({num=%[[N_CONV]] : si32}, {num=%[[N_CONV2]] : si32} [#acc.device_type<nvidia>], {num=%[[N_CONV2]] : si32} [#acc.device_type<radeon>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N) device_type(nvidia) gang(static:*)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32}, {static=%[[STAR_CONST]] : i64} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+#pragma acc loop gang(static:N, num: N + 1) device_type(nvidia) gang(static:*, num : N + 2)
+ for(unsigned I = 0; I < N; ++I);
+ // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[CIR_ONE_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64
+ // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CIR_TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO:.*]] = cir.binop(add, %[[N_LOAD3]], %[[CIR_TWO_CONST]]) nsw : !s32i
+ // CHECK-NEXT: %[[N_PLUS_TWO_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_TWO]] : !s32i to si32
+ // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32, num=%[[N_PLUS_ONE_CONV]] : si32}, {static=%[[STAR_CONST]] : i64, num=%[[N_PLUS_TWO_CONV]] : si32} [#acc.device_type<nvidia>]) {
+ // CHECK: acc.yield
+ // CHECK-NEXT: } loc
+ }
}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index ca564037fad19..5d5add6318e06 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2231,6 +2231,16 @@ def OpenACC_LoopOp : OpenACC_Op<"loop",
// device_types. This is for the case where there is no expression specified
// in a 'worker'.
void addEmptyWorker(MLIRContext *, llvm::ArrayRef<DeviceType>);
+
+ // Adds a collection of operands for a 'gang' clause that has various types
+ // corresponding to each operand.
+ void addGangOperands(MLIRContext *, llvm::ArrayRef<DeviceType>,
+ llvm::ArrayRef<GangArgType>, mlir::ValueRange);
+
+ // Add an empty value to the 'gang' list with a current list of
+ // device_types. This is for the case where there is no expression specified
+ // in a 'gang'.
+ void addEmptyGang(MLIRContext *, llvm::ArrayRef<DeviceType>);
}];
let hasCustomAssemblyFormat = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 9f4645a4a7ca8..561f1aef7d5c9 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -2748,6 +2748,49 @@ void acc::LoopOp::addEmptyWorker(
effectiveDeviceTypes));
}
+void acc::LoopOp::addEmptyGang(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ setGangAttr(addDeviceTypeAffectedOperandHelper(context, getGangAttr(),
+ effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addGangOperands(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes,
+ llvm::ArrayRef<GangArgType> argTypes, mlir::ValueRange values) {
+ llvm::SmallVector<int32_t> segments;
+ if (getGangOperandsSegments())
+ llvm::copy(*getGangOperandsSegments(), std::back_inserter(segments));
+
+ unsigned beforeCount = segments.size();
+
+ setGangOperandsDeviceTypeAttr(addDeviceTypeAffectedOperandHelper(
+ context, getGangOperandsDeviceTypeAttr(), effectiveDeviceTypes, values,
+ getGangOperandsMutable(), segments));
+
+ setGangOperandsSegments(segments);
+
+ // This is a bit of extra work to make sure we update the 'types' correctly by
+ // adding to the types collection the correct number of times. We could
+ // potentially add something similar to the
+ // addDeviceTypeAffectedOperandHelper, but it seems that would be pretty
+ // excessive for a one-off case.
+ unsigned numAdded = segments.size() - beforeCount;
+
+ if (numAdded > 0) {
+ llvm::SmallVector<mlir::Attribute> gangTypes;
+ if (getGangOperandsArgTypeAttr())
+ llvm::copy(getGangOperandsArgTypeAttr(), std::back_inserter(gangTypes));
+
+ for (unsigned I = 0; I < numAdded; ++I)
+ llvm::transform(argTypes, std::back_inserter(gangTypes),
+ [=](mlir::acc::GangArgType gangTy) {
+ return mlir::acc::GangArgTypeAttr::get(context, gangTy);
+ });
+
+ setGangOperandsArgTypeAttr(mlir::ArrayAttr::get(context, gangTypes));
+ }
+}
+
//===----------------------------------------------------------------------===//
// DataOp
//===----------------------------------------------------------------------===//
|
|
Linux failure is an unrelated Flang failure. |
bcardosolopes
approved these changes
May 8, 2025
| } else { | ||
| llvm::SmallVector<mlir::Value> values; | ||
| llvm::SmallVector<mlir::acc::GangArgType> argTypes; | ||
| for (unsigned I = 0; I < clause.getNumExprs(); ++I) { |
andykaylor
approved these changes
May 8, 2025
| MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes, | ||
| llvm::ArrayRef<GangArgType> argTypes, mlir::ValueRange values) { | ||
| llvm::SmallVector<int32_t> segments; | ||
| if (getGangOperandsSegments()) |
There was a problem hiding this comment.
Can you save the return value here so you don't have to make this call again on the next line?
| if (getGangOperandsArgTypeAttr()) | ||
| llvm::copy(getGangOperandsArgTypeAttr(), std::back_inserter(gangTypes)); | ||
|
|
||
| for (unsigned I = 0; I < numAdded; ++I) |
There was a problem hiding this comment.
Suggested change
| for (unsigned I = 0; I < numAdded; ++I) | |
| for (unsigned i = 0; i < numAdded; ++i) |
| @@ -107,6 +107,18 @@ class OpenACCClauseCIREmitter final | |||
| .CaseLower("radeon", mlir::acc::DeviceType::Radeon); | |||
| } | |||
|
|
|||
| mlir::acc::GangArgType decodeGangType(OpenACCGangKind GK) { | |||
There was a problem hiding this comment.
Suggested change
| mlir::acc::GangArgType decodeGangType(OpenACCGangKind GK) { | |
| mlir::acc::GangArgType decodeGangType(OpenACCGangKind gk) { |
| } else { | ||
| llvm::SmallVector<mlir::Value> values; | ||
| llvm::SmallVector<mlir::acc::GangArgType> argTypes; | ||
| for (unsigned I = 0; I < clause.getNumExprs(); ++I) { |
There was a problem hiding this comment.
Suggested change
| for (unsigned I = 0; I < clause.getNumExprs(); ++I) { | |
| for (unsigned i = 0; i < clause.getNumExprs(); ++i) { |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This clause requires an entire additional collection to keep track of the gang 'kind' or 'type'. That work is maintained in the OpenACC dialect functions. Otherwise, this is effectively the same as the worker/vectors.