[acc][mlir] Add 'if-condition' to 'atomic' operations. #164003
Merged
+70
−10
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
OpenACC 3.4 includes the ability to add an 'if' to an atomic operation, see: OpenACC/openacc-spec#511 This patch adds support to this for the dialect, so that Clang can use it soon.
llvmbot
added
mlir
flang
|
I think I got this right! Please feel free to add any other reviews here you think should see this. |
|
@llvm/pr-subscribers-mlir-openacc @llvm/pr-subscribers-openacc Author: Erich Keane (erichkeane) ChangesOpenACC 3.4 includes the ability to add an 'if' to an atomic operation, see: https://github.com/OpenACC/openacc-spec/pull/511 This patch adds support to this for the dialect, so that Clang can use it soon. Full diff: https://github.com/llvm/llvm-project/pull/164003.diff 4 Files Affected:
diff --git a/flang/lib/Lower/OpenACC.cpp b/flang/lib/Lower/OpenACC.cpp
index cfb18914e8126..e38dc567096da 100644
--- a/flang/lib/Lower/OpenACC.cpp
+++ b/flang/lib/Lower/OpenACC.cpp
@@ -327,7 +327,8 @@ genAtomicCaptureStatement(Fortran::lower::AbstractConverter &converter,
fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
mlir::acc::AtomicReadOp::create(firOpBuilder, loc, fromAddress, toAddress,
- mlir::TypeAttr::get(elementType));
+ mlir::TypeAttr::get(elementType),
+ /*IfCond=*/mlir::Value{});
}
/// Used to generate atomic.write operation which is created in existing
@@ -347,7 +348,8 @@ genAtomicWriteStatement(Fortran::lower::AbstractConverter &converter,
rhsExpr = firOpBuilder.createConvert(loc, varType, rhsExpr);
firOpBuilder.restoreInsertionPoint(insertionPoint);
- mlir::acc::AtomicWriteOp::create(firOpBuilder, loc, lhsAddr, rhsExpr);
+ mlir::acc::AtomicWriteOp::create(firOpBuilder, loc, lhsAddr, rhsExpr,
+ /*IfCond=*/mlir::Value{});
}
/// Used to generate atomic.update operation which is created in existing
@@ -463,7 +465,8 @@ static inline void genAtomicUpdateStatement(
mlir::Operation *atomicUpdateOp = nullptr;
atomicUpdateOp =
- mlir::acc::AtomicUpdateOp::create(firOpBuilder, currentLocation, lhsAddr);
+ mlir::acc::AtomicUpdateOp::create(firOpBuilder, currentLocation, lhsAddr,
+ /*IfCond=*/mlir::Value{});
llvm::SmallVector<mlir::Type> varTys = {varType};
llvm::SmallVector<mlir::Location> locs = {currentLocation};
@@ -588,7 +591,9 @@ void genAtomicCapture(Fortran::lower::AbstractConverter &converter,
fir::getBase(converter.genExprValue(assign2.lhs, stmtCtx)).getType();
mlir::Operation *atomicCaptureOp = nullptr;
- atomicCaptureOp = mlir::acc::AtomicCaptureOp::create(firOpBuilder, loc);
+ atomicCaptureOp =
+ mlir::acc::AtomicCaptureOp::create(firOpBuilder, loc,
+ /*IfCond=*/mlir::Value{});
firOpBuilder.createBlock(&(atomicCaptureOp->getRegion(0)));
mlir::Block &block = atomicCaptureOp->getRegion(0).back();
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 1eaa21b46554c..78166f3ffcb0c 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2787,10 +2787,16 @@ def AtomicReadOp : OpenACC_Op<"atomic.read", [AtomicReadOpInterface]> {
let arguments = (ins OpenACC_PointerLikeType:$x,
OpenACC_PointerLikeType:$v,
- TypeAttr:$element_type);
+ TypeAttr:$element_type,
+ Optional<I1>:$ifCond
+ );
let assemblyFormat = [{
+ oilist(
+ `if` `(` $ifCond `)`
+ )
$v `=` $x
- `:` type($v) `,` type($x) `,` $element_type attr-dict
+ `:` type($v) `,` type($x) `,` $element_type
+ attr-dict
}];
let hasVerifier = 1;
}
@@ -2809,8 +2815,12 @@ def AtomicWriteOp : OpenACC_Op<"atomic.write",[AtomicWriteOpInterface]> {
}];
let arguments = (ins OpenACC_PointerLikeType:$x,
- AnyType:$expr);
+ AnyType:$expr,
+ Optional<I1>:$ifCond);
let assemblyFormat = [{
+ oilist(
+ `if` `(` $ifCond `)`
+ )
$x `=` $expr
`:` type($x) `,` type($expr)
attr-dict
@@ -2850,10 +2860,15 @@ def AtomicUpdateOp : OpenACC_Op<"atomic.update",
let arguments = (ins Arg<OpenACC_PointerLikeType,
"Address of variable to be updated",
- [MemRead, MemWrite]>:$x);
+ [MemRead, MemWrite]>:$x,
+ Optional<I1>:$ifCond);
let regions = (region SizedRegion<1>:$region);
let assemblyFormat = [{
- $x `:` type($x) $region attr-dict
+ oilist(
+ `if` `(` $ifCond `)`
+ )
+ $x `:` type($x)
+ $region attr-dict
}];
let hasVerifier = 1;
let hasRegionVerifier = 1;
@@ -2896,8 +2911,13 @@ def AtomicCaptureOp : OpenACC_Op<"atomic.capture",
}];
+ let arguments = (ins Optional<I1>:$ifCond);
+
let regions = (region SizedRegion<1>:$region);
let assemblyFormat = [{
+ oilist(
+ `if` `(` $ifCond `)`
+ )
$region attr-dict
}];
let hasRegionVerifier = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 90cbbd86dc002..cffdf128d1e8e 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -3858,7 +3858,8 @@ LogicalResult AtomicUpdateOp::canonicalize(AtomicUpdateOp op,
}
if (Value writeVal = op.getWriteOpVal()) {
- rewriter.replaceOpWithNewOp<AtomicWriteOp>(op, op.getX(), writeVal);
+ rewriter.replaceOpWithNewOp<AtomicWriteOp>(op, op.getX(), writeVal,
+ op.getIfCond());
return success();
}
diff --git a/mlir/test/Dialect/OpenACC/ops.mlir b/mlir/test/Dialect/OpenACC/ops.mlir
index 1484d7efd87c2..8713689ed5799 100644
--- a/mlir/test/Dialect/OpenACC/ops.mlir
+++ b/mlir/test/Dialect/OpenACC/ops.mlir
@@ -1766,6 +1766,12 @@ acc.set default_async(%i32Value : i32)
func.func @acc_atomic_read(%v: memref<i32>, %x: memref<i32>) {
// CHECK: acc.atomic.read %[[v]] = %[[x]] : memref<i32>, memref<i32>, i32
acc.atomic.read %v = %x : memref<i32>, memref<i32>, i32
+
+ // CHECK-NEXT: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.read if(%[[IFCOND1]]) %[[v]] = %[[x]] : memref<i32>, memref<i32>, i32
+ %ifCond = arith.constant true
+ acc.atomic.read if(%ifCond) %v = %x : memref<i32>, memref<i32>, i32
+
return
}
@@ -1776,6 +1782,12 @@ func.func @acc_atomic_read(%v: memref<i32>, %x: memref<i32>) {
func.func @acc_atomic_write(%addr : memref<i32>, %val : i32) {
// CHECK: acc.atomic.write %[[ADDR]] = %[[VAL]] : memref<i32>, i32
acc.atomic.write %addr = %val : memref<i32>, i32
+
+ // CHECK-NEXT: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.write if(%[[IFCOND1]]) %[[ADDR]] = %[[VAL]] : memref<i32>, i32
+ %ifCond = arith.constant true
+ acc.atomic.write if(%ifCond) %addr = %val : memref<i32>, i32
+
return
}
@@ -1793,6 +1805,19 @@ func.func @acc_atomic_update(%x : memref<i32>, %expr : i32, %xBool : memref<i1>,
%newval = llvm.add %xval, %expr : i32
acc.yield %newval : i32
}
+
+ // CHECK: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.update if(%[[IFCOND1]]) %[[X]] : memref<i32>
+ // CHECK-NEXT: (%[[XVAL:.*]]: i32):
+ // CHECK-NEXT: %[[NEWVAL:.*]] = llvm.add %[[XVAL]], %[[EXPR]] : i32
+ // CHECK-NEXT: acc.yield %[[NEWVAL]] : i32
+ %ifCond = arith.constant true
+ acc.atomic.update if (%ifCond) %x : memref<i32> {
+ ^bb0(%xval: i32):
+ %newval = llvm.add %xval, %expr : i32
+ acc.yield %newval : i32
+ }
+
// CHECK: acc.atomic.update %[[XBOOL]] : memref<i1>
// CHECK-NEXT: (%[[XVAL:.*]]: i1):
// CHECK-NEXT: %[[NEWVAL:.*]] = llvm.and %[[XVAL]], %[[EXPRBOOL]] : i1
@@ -1902,6 +1927,17 @@ func.func @acc_atomic_capture(%v: memref<i32>, %x: memref<i32>, %expr: i32) {
acc.atomic.write %x = %expr : memref<i32>, i32
}
+ // CHECK: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.capture if(%[[IFCOND1]]) {
+ // CHECK-NEXT: acc.atomic.read %[[v]] = %[[x]] : memref<i32>, memref<i32>, i32
+ // CHECK-NEXT: acc.atomic.write %[[x]] = %[[expr]] : memref<i32>, i32
+ // CHECK-NEXT: }
+ %ifCond = arith.constant true
+ acc.atomic.capture if (%ifCond) {
+ acc.atomic.read %v = %x : memref<i32>, memref<i32>, i32
+ acc.atomic.write %x = %expr : memref<i32>, i32
+ }
+
return
}
|
|
@llvm/pr-subscribers-mlir Author: Erich Keane (erichkeane) ChangesOpenACC 3.4 includes the ability to add an 'if' to an atomic operation, see: https://github.com/OpenACC/openacc-spec/pull/511 This patch adds support to this for the dialect, so that Clang can use it soon. Full diff: https://github.com/llvm/llvm-project/pull/164003.diff 4 Files Affected:
diff --git a/flang/lib/Lower/OpenACC.cpp b/flang/lib/Lower/OpenACC.cpp
index cfb18914e8126..e38dc567096da 100644
--- a/flang/lib/Lower/OpenACC.cpp
+++ b/flang/lib/Lower/OpenACC.cpp
@@ -327,7 +327,8 @@ genAtomicCaptureStatement(Fortran::lower::AbstractConverter &converter,
fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
mlir::acc::AtomicReadOp::create(firOpBuilder, loc, fromAddress, toAddress,
- mlir::TypeAttr::get(elementType));
+ mlir::TypeAttr::get(elementType),
+ /*IfCond=*/mlir::Value{});
}
/// Used to generate atomic.write operation which is created in existing
@@ -347,7 +348,8 @@ genAtomicWriteStatement(Fortran::lower::AbstractConverter &converter,
rhsExpr = firOpBuilder.createConvert(loc, varType, rhsExpr);
firOpBuilder.restoreInsertionPoint(insertionPoint);
- mlir::acc::AtomicWriteOp::create(firOpBuilder, loc, lhsAddr, rhsExpr);
+ mlir::acc::AtomicWriteOp::create(firOpBuilder, loc, lhsAddr, rhsExpr,
+ /*IfCond=*/mlir::Value{});
}
/// Used to generate atomic.update operation which is created in existing
@@ -463,7 +465,8 @@ static inline void genAtomicUpdateStatement(
mlir::Operation *atomicUpdateOp = nullptr;
atomicUpdateOp =
- mlir::acc::AtomicUpdateOp::create(firOpBuilder, currentLocation, lhsAddr);
+ mlir::acc::AtomicUpdateOp::create(firOpBuilder, currentLocation, lhsAddr,
+ /*IfCond=*/mlir::Value{});
llvm::SmallVector<mlir::Type> varTys = {varType};
llvm::SmallVector<mlir::Location> locs = {currentLocation};
@@ -588,7 +591,9 @@ void genAtomicCapture(Fortran::lower::AbstractConverter &converter,
fir::getBase(converter.genExprValue(assign2.lhs, stmtCtx)).getType();
mlir::Operation *atomicCaptureOp = nullptr;
- atomicCaptureOp = mlir::acc::AtomicCaptureOp::create(firOpBuilder, loc);
+ atomicCaptureOp =
+ mlir::acc::AtomicCaptureOp::create(firOpBuilder, loc,
+ /*IfCond=*/mlir::Value{});
firOpBuilder.createBlock(&(atomicCaptureOp->getRegion(0)));
mlir::Block &block = atomicCaptureOp->getRegion(0).back();
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 1eaa21b46554c..78166f3ffcb0c 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2787,10 +2787,16 @@ def AtomicReadOp : OpenACC_Op<"atomic.read", [AtomicReadOpInterface]> {
let arguments = (ins OpenACC_PointerLikeType:$x,
OpenACC_PointerLikeType:$v,
- TypeAttr:$element_type);
+ TypeAttr:$element_type,
+ Optional<I1>:$ifCond
+ );
let assemblyFormat = [{
+ oilist(
+ `if` `(` $ifCond `)`
+ )
$v `=` $x
- `:` type($v) `,` type($x) `,` $element_type attr-dict
+ `:` type($v) `,` type($x) `,` $element_type
+ attr-dict
}];
let hasVerifier = 1;
}
@@ -2809,8 +2815,12 @@ def AtomicWriteOp : OpenACC_Op<"atomic.write",[AtomicWriteOpInterface]> {
}];
let arguments = (ins OpenACC_PointerLikeType:$x,
- AnyType:$expr);
+ AnyType:$expr,
+ Optional<I1>:$ifCond);
let assemblyFormat = [{
+ oilist(
+ `if` `(` $ifCond `)`
+ )
$x `=` $expr
`:` type($x) `,` type($expr)
attr-dict
@@ -2850,10 +2860,15 @@ def AtomicUpdateOp : OpenACC_Op<"atomic.update",
let arguments = (ins Arg<OpenACC_PointerLikeType,
"Address of variable to be updated",
- [MemRead, MemWrite]>:$x);
+ [MemRead, MemWrite]>:$x,
+ Optional<I1>:$ifCond);
let regions = (region SizedRegion<1>:$region);
let assemblyFormat = [{
- $x `:` type($x) $region attr-dict
+ oilist(
+ `if` `(` $ifCond `)`
+ )
+ $x `:` type($x)
+ $region attr-dict
}];
let hasVerifier = 1;
let hasRegionVerifier = 1;
@@ -2896,8 +2911,13 @@ def AtomicCaptureOp : OpenACC_Op<"atomic.capture",
}];
+ let arguments = (ins Optional<I1>:$ifCond);
+
let regions = (region SizedRegion<1>:$region);
let assemblyFormat = [{
+ oilist(
+ `if` `(` $ifCond `)`
+ )
$region attr-dict
}];
let hasRegionVerifier = 1;
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 90cbbd86dc002..cffdf128d1e8e 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -3858,7 +3858,8 @@ LogicalResult AtomicUpdateOp::canonicalize(AtomicUpdateOp op,
}
if (Value writeVal = op.getWriteOpVal()) {
- rewriter.replaceOpWithNewOp<AtomicWriteOp>(op, op.getX(), writeVal);
+ rewriter.replaceOpWithNewOp<AtomicWriteOp>(op, op.getX(), writeVal,
+ op.getIfCond());
return success();
}
diff --git a/mlir/test/Dialect/OpenACC/ops.mlir b/mlir/test/Dialect/OpenACC/ops.mlir
index 1484d7efd87c2..8713689ed5799 100644
--- a/mlir/test/Dialect/OpenACC/ops.mlir
+++ b/mlir/test/Dialect/OpenACC/ops.mlir
@@ -1766,6 +1766,12 @@ acc.set default_async(%i32Value : i32)
func.func @acc_atomic_read(%v: memref<i32>, %x: memref<i32>) {
// CHECK: acc.atomic.read %[[v]] = %[[x]] : memref<i32>, memref<i32>, i32
acc.atomic.read %v = %x : memref<i32>, memref<i32>, i32
+
+ // CHECK-NEXT: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.read if(%[[IFCOND1]]) %[[v]] = %[[x]] : memref<i32>, memref<i32>, i32
+ %ifCond = arith.constant true
+ acc.atomic.read if(%ifCond) %v = %x : memref<i32>, memref<i32>, i32
+
return
}
@@ -1776,6 +1782,12 @@ func.func @acc_atomic_read(%v: memref<i32>, %x: memref<i32>) {
func.func @acc_atomic_write(%addr : memref<i32>, %val : i32) {
// CHECK: acc.atomic.write %[[ADDR]] = %[[VAL]] : memref<i32>, i32
acc.atomic.write %addr = %val : memref<i32>, i32
+
+ // CHECK-NEXT: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.write if(%[[IFCOND1]]) %[[ADDR]] = %[[VAL]] : memref<i32>, i32
+ %ifCond = arith.constant true
+ acc.atomic.write if(%ifCond) %addr = %val : memref<i32>, i32
+
return
}
@@ -1793,6 +1805,19 @@ func.func @acc_atomic_update(%x : memref<i32>, %expr : i32, %xBool : memref<i1>,
%newval = llvm.add %xval, %expr : i32
acc.yield %newval : i32
}
+
+ // CHECK: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.update if(%[[IFCOND1]]) %[[X]] : memref<i32>
+ // CHECK-NEXT: (%[[XVAL:.*]]: i32):
+ // CHECK-NEXT: %[[NEWVAL:.*]] = llvm.add %[[XVAL]], %[[EXPR]] : i32
+ // CHECK-NEXT: acc.yield %[[NEWVAL]] : i32
+ %ifCond = arith.constant true
+ acc.atomic.update if (%ifCond) %x : memref<i32> {
+ ^bb0(%xval: i32):
+ %newval = llvm.add %xval, %expr : i32
+ acc.yield %newval : i32
+ }
+
// CHECK: acc.atomic.update %[[XBOOL]] : memref<i1>
// CHECK-NEXT: (%[[XVAL:.*]]: i1):
// CHECK-NEXT: %[[NEWVAL:.*]] = llvm.and %[[XVAL]], %[[EXPRBOOL]] : i1
@@ -1902,6 +1927,17 @@ func.func @acc_atomic_capture(%v: memref<i32>, %x: memref<i32>, %expr: i32) {
acc.atomic.write %x = %expr : memref<i32>, i32
}
+ // CHECK: %[[IFCOND1:.*]] = arith.constant true
+ // CHECK-NEXT: acc.atomic.capture if(%[[IFCOND1]]) {
+ // CHECK-NEXT: acc.atomic.read %[[v]] = %[[x]] : memref<i32>, memref<i32>, i32
+ // CHECK-NEXT: acc.atomic.write %[[x]] = %[[expr]] : memref<i32>, i32
+ // CHECK-NEXT: }
+ %ifCond = arith.constant true
+ acc.atomic.capture if (%ifCond) {
+ acc.atomic.read %v = %x : memref<i32>, memref<i32>, i32
+ acc.atomic.write %x = %expr : memref<i32>, i32
+ }
+
return
}
|
razvanlupusoru
approved these changes
Oct 17, 2025
flang/lib/Lower/OpenACC.cpp
Outdated
| mlir::acc::AtomicReadOp::create(firOpBuilder, loc, fromAddress, toAddress, | ||
| mlir::TypeAttr::get(elementType)); | ||
| mlir::TypeAttr::get(elementType), | ||
| /*IfCond=*/mlir::Value{}); |
There was a problem hiding this comment.
nit: Here and below - should it be ifCond instead of IfCond?
| "Address of variable to be updated", | ||
| [MemRead, MemWrite]>:$x); | ||
| [MemRead, MemWrite]>:$x, | ||
| Optional<I1>:$ifCond); |
There was a problem hiding this comment.
nit: mismatched indentation
There was a problem hiding this comment.
Ah! I missed you don't consider ins here for the indent, so I think I got this right now.
|
Your link to https://github.com/OpenACC/openacc-spec/pull/511 is not valid without access to that repository. Since OpenACC 3.4 is already released, it should be enough to refer to that. :) |
razvanlupusoru
changed the title
Ah, yes! I've improved the commit message, please let me know what you think. |
It's excellent! Thank you. |
| TypeAttr:$element_type); | ||
| TypeAttr:$element_type, | ||
| Optional<I1>:$ifCond | ||
| ); |
There was a problem hiding this comment.
nit: maybe move it up to line above?
|
Looks great. I already approved so it can be merged anytime! Thank you! |
|
It isn't clear to me what the windows failure is. As far as I can tell there are no build failures, and it hasn't finished testing, but is saying it is failed. I might have to keep an eye on this and poke at it again on Monday. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
OpenACC 3.4 includes the ability to add an 'if' to an atomic operation.
From the change log:
Added the if clause to the atomic construct to enable conditional atomic operations based867 on the parallelism strategy employedIn 2.12, the C/C++ grammar is changed to say:
#pragma acc atomic [ atomic-clause ] [ if( condition ) ] new-lineWith corresponding changes to the Fortran standard
This patch adds support to this for the dialect, so that Clang can use it soon.