[flang-commits] [flang] [mlir] [mlir][gpu] Add stream-based async mode to gpu.launch (PR #213031)
via flang-commits
flang-commits at lists.llvm.org
Thu Jul 30 07:25:26 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-flang-fir-hlfir
Author: Ivan R. Ivanov (ivanradanov)
<details>
<summary>Changes</summary>
Currently, only gpu.launch_func supports stream based asynchronicity. This patch adds the same mode to gpu.launch, and makes kernel outlining preserve that when converting to gpu.launch_func.
This patch also documents the stream-based model in the operations description, which was previously missing.
Assisted-by: Claude Code
---
Patch is 30.37 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/213031.diff
11 Files Affected:
- (modified) flang/lib/Optimizer/CodeGen/TargetRewrite.cpp (+7-7)
- (modified) flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp (+8-8)
- (modified) flang/test/Fir/CUDA/cuda-gpu-launch-func.mlir (+3-3)
- (modified) flang/test/Fir/CUDA/cuda-launch.fir (+1-1)
- (modified) flang/test/Fir/CUDA/cuda-stream.mlir (+2-2)
- (modified) mlir/include/mlir/Dialect/GPU/IR/GPUOps.td (+22-18)
- (modified) mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp (+1-1)
- (modified) mlir/lib/Dialect/GPU/IR/GPUDialect.cpp (+56-51)
- (modified) mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp (+1-1)
- (modified) mlir/test/Dialect/GPU/invalid.mlir (+50-5)
- (modified) mlir/test/Dialect/GPU/ops.mlir (+11)
``````````diff
diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
index 37f502d1aa1e0..22435aafe352d 100644
--- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
+++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
@@ -536,13 +536,13 @@ class TargetRewrite : public fir::impl::TargetRewritePassBase<TargetRewrite> {
// TODO propagate/update call argument and result attributes.
if constexpr (std::is_same_v<std::decay_t<A>, mlir::gpu::LaunchFuncOp>) {
mlir::Value asyncToken = callOp.getAsyncToken();
- auto newCall = A::create(*rewriter, loc, callOp.getKernel(),
- callOp.getGridSizeOperandValues(),
- callOp.getBlockSizeOperandValues(),
- callOp.getDynamicSharedMemorySize(), newOpers,
- asyncToken ? asyncToken.getType() : nullptr,
- callOp.getAsyncDependencies(),
- /*clusterSize=*/std::nullopt);
+ auto newCall = A::create(
+ *rewriter, loc, callOp.getKernel(), callOp.getGridSizeOperandValues(),
+ callOp.getBlockSizeOperandValues(),
+ callOp.getDynamicSharedMemorySize(), newOpers,
+ asyncToken ? asyncToken.getType() : nullptr,
+ callOp.getAsyncDependencies(), callOp.getAsyncObject(),
+ /*clusterSize=*/std::nullopt);
if (callOp.getClusterSizeX())
newCall.getClusterSizeXMutable().assign(callOp.getClusterSizeX());
if (callOp.getClusterSizeY())
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
index 44f1b5154a2b0..dfd28ef96cf64 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFOpConversion.cpp
@@ -479,23 +479,23 @@ struct CUFLaunchOpConversion
args.push_back(arg);
}
mlir::Value dynamicShmemSize = op.getBytes() ? op.getBytes() : zero;
+ mlir::Type tokenType = nullptr;
+ SmallVector<Value, 1> tokens;
+ if (op.getStream()) {
+ tokens.push_back(
+ cuf::StreamCastOp::create(rewriter, loc, op.getStream()));
+ tokenType = tokens.front().getType();
+ }
auto gpuLaunchOp = mlir::gpu::LaunchFuncOp::create(
rewriter, loc, kernelName,
mlir::gpu::KernelDim3{gridSizeX, gridSizeY, gridSizeZ},
mlir::gpu::KernelDim3{blockSizeX, blockSizeY, blockSizeZ},
- dynamicShmemSize, args);
+ dynamicShmemSize, args, tokenType, tokens);
if (clusterDimX && clusterDimY && clusterDimZ) {
gpuLaunchOp.getClusterSizeXMutable().assign(clusterDimX);
gpuLaunchOp.getClusterSizeYMutable().assign(clusterDimY);
gpuLaunchOp.getClusterSizeZMutable().assign(clusterDimZ);
}
- if (op.getStream()) {
- mlir::OpBuilder::InsertionGuard guard(rewriter);
- rewriter.setInsertionPoint(gpuLaunchOp);
- mlir::Value stream =
- cuf::StreamCastOp::create(rewriter, loc, op.getStream());
- gpuLaunchOp.getAsyncDependenciesMutable().append(stream);
- }
if (procAttr)
gpuLaunchOp->setAttr(cuf::getProcAttrName(), procAttr);
else
diff --git a/flang/test/Fir/CUDA/cuda-gpu-launch-func.mlir b/flang/test/Fir/CUDA/cuda-gpu-launch-func.mlir
index 47e6dfd220081..03e441ada81c9 100644
--- a/flang/test/Fir/CUDA/cuda-gpu-launch-func.mlir
+++ b/flang/test/Fir/CUDA/cuda-gpu-launch-func.mlir
@@ -180,7 +180,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<!llvm.ptr<272>, d
%2 = llvm.mlir.constant(0 : i32) : i32
%3 = llvm.mlir.constant(10 : index) : i64
%token = cuf.stream_cast %stream : !llvm.ptr
- gpu.launch_func [%token] @cuda_device_mod::@_QMmod1Psub1 blocks in (%3, %3, %0) threads in (%3, %3, %0) : i64 dynamic_shared_memory_size %2 {cuf.proc_attr = #cuf.cuda_proc<global>}
+ %res_token = gpu.launch_func async [%token] @cuda_device_mod::@_QMmod1Psub1 blocks in (%3, %3, %0) threads in (%3, %3, %0) : i64 dynamic_shared_memory_size %2 {cuf.proc_attr = #cuf.cuda_proc<global>}
llvm.return
}
gpu.binary @cuda_device_mod [#gpu.object<#nvvm.target, "">]
@@ -213,7 +213,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<i1, dense<8> : ve
%13 = llvm.mlir.addressof @_QQclX91d13f6e74caa2f03965d7a7c6a8fdd5 : !llvm.ptr
%14 = llvm.call @_FortranACUFMemAlloc(%2, %11, %13, %6) : (i64, i32, !llvm.ptr, i32) -> !llvm.ptr
%token = cuf.stream_cast %stream : !llvm.ptr
- gpu.launch_func [%token] @cuda_device_mod::@_QMmod1Psub1 blocks in (%7, %7, %7) threads in (%12, %7, %7) : i64 dynamic_shared_memory_size %11 args(%14 : !llvm.ptr) {cuf.proc_attr = #cuf.cuda_proc<grid_global>}
+ %res_token = gpu.launch_func async [%token] @cuda_device_mod::@_QMmod1Psub1 blocks in (%7, %7, %7) threads in (%12, %7, %7) : i64 dynamic_shared_memory_size %11 args(%14 : !llvm.ptr) {cuf.proc_attr = #cuf.cuda_proc<grid_global>}
llvm.return
}
llvm.func @_QMmod1Psub1(!llvm.ptr) -> ()
@@ -244,7 +244,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<!llvm.ptr<272>, d
%3 = llvm.mlir.constant(10 : index) : i64
%stream = llvm.alloca %0 x i64 : (i64) -> !llvm.ptr
%token = cuf.stream_cast %stream : !llvm.ptr
- %4 = gpu.launch_func async [%token] @cuda_device_mod::@_QMmod1Psub1 blocks in (%3, %3, %0) threads in (%3, %3, %0) : i64 dynamic_shared_memory_size %2 {cuf.proc_attr = #cuf.cuda_proc<global>}
+ %res_token = gpu.launch_func async [%token] @cuda_device_mod::@_QMmod1Psub1 blocks in (%3, %3, %0) threads in (%3, %3, %0) : i64 dynamic_shared_memory_size %2 {cuf.proc_attr = #cuf.cuda_proc<global>}
llvm.return
}
gpu.binary @cuda_device_mod [#gpu.object<#nvvm.target, "">]
diff --git a/flang/test/Fir/CUDA/cuda-launch.fir b/flang/test/Fir/CUDA/cuda-launch.fir
index 92db6ecaadc45..e5ea7bb2479ac 100644
--- a/flang/test/Fir/CUDA/cuda-launch.fir
+++ b/flang/test/Fir/CUDA/cuda-launch.fir
@@ -155,4 +155,4 @@ module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_e
// CHECK: %[[STREAM:.*]] = fir.alloca i64 {bindc_name = "stream", uniq_name = "_QMtest_callFhostEstream"}
// CHECK: %[[DECL_STREAM:.*]]:2 = hlfir.declare %[[STREAM]] {uniq_name = "_QMtest_callFhostEstream"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)
// CHECK: %[[TOKEN:.*]] = cuf.stream_cast %[[DECL_STREAM]]#0 : !fir.ref<i64>
-// CHECK: gpu.launch_func [%[[TOKEN]]] @cuda_device_mod::@_QMdevptrPtest
+// CHECK: %[[RES_TOKEN:.*]] = gpu.launch_func async [%[[TOKEN]]] @cuda_device_mod::@_QMdevptrPtest
diff --git a/flang/test/Fir/CUDA/cuda-stream.mlir b/flang/test/Fir/CUDA/cuda-stream.mlir
index a501603fd35d1..1c463ea39458e 100644
--- a/flang/test/Fir/CUDA/cuda-stream.mlir
+++ b/flang/test/Fir/CUDA/cuda-stream.mlir
@@ -10,7 +10,7 @@ module attributes {gpu.container_module} {
%0 = fir.alloca i64
%1 = arith.constant 1 : index
%asyncTok = cuf.stream_cast %0 : !fir.ref<i64>
- gpu.launch_func [%asyncTok] @cuda_device_mod::@_QMmod1Psub1 blocks in (%1, %1, %1) threads in (%1, %1, %1) args() {cuf.proc_attr = #cuf.cuda_proc<grid_global>}
+ %res_token = gpu.launch_func async [%asyncTok] @cuda_device_mod::@_QMmod1Psub1 blocks in (%1, %1, %1) threads in (%1, %1, %1) args() {cuf.proc_attr = #cuf.cuda_proc<grid_global>}
return
}
}
@@ -18,4 +18,4 @@ module attributes {gpu.container_module} {
// CHECK-LABEL: func.func @_QMmod1Phost_sub()
// CHECK: %[[STREAM:.*]] = fir.alloca i64
// CHECK: %[[TOKEN:.*]] = cuf.stream_cast %[[STREAM]] : !fir.ref<i64>
-// CHECK: gpu.launch_func [%[[TOKEN]]] @cuda_device_mod::@_QMmod1Psub1
+// CHECK: %[[RES_TOKEN:.*]] = gpu.launch_func async [%[[TOKEN]]] @cuda_device_mod::@_QMmod1Psub1
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index a9032820f5646..91edd1dbd5dad 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -637,13 +637,22 @@ def GPU_LaunchFuncOp :GPU_Op<"launch_func", [
operation has a symbol attribute named `kernel` to identify the fully
specified kernel function to launch (both the gpu.module and func).
- The `gpu.launch_func` supports async dependencies: the kernel does not start
+ By the default, the host implicitly blocks until kernel execution has
+ completed.
+
+ Otherwise, the operation supports two async models.
+
+ The first one is dependency-based and is enabled when the `async` keyword is
+ present in text form, and corresponds to when the operation produces the
+ optional token result of type `!gpu.async.token`. Other async GPU ops can
+ take this token as dependency. In this case, the `gpu.launch_func` does not
+ block, and supports specifying async dependencies: the kernel does not start
executing until the ops producing those async dependencies have completed.
- By the default, the host implicitly blocks until kernel execution has
- completed. If the `async` keyword is present, the host does not block but
- instead a `!gpu.async.token` is returned. Other async GPU ops can take this
- token as dependency.
+ The second async model is stream-based. When `asyncObject` is present, the
+ launch operation is queued to execute on the queue represented by it.
+
+ The two async models are mutually exclusive.
The operation requires at least the grid and block sizes along the x,y,z
dimensions as arguments. When a lower-dimensional kernel is required,
@@ -734,16 +743,13 @@ def GPU_LaunchFuncOp :GPU_Op<"launch_func", [
"ValueRange":$kernelOperands,
CArg<"Type", "nullptr">:$asyncTokenType,
CArg<"ValueRange", "{}">:$asyncDependencies,
+ CArg<"Value", "nullptr">:$asyncObject,
CArg<"std::optional<KernelDim3>", "std::nullopt">:$clusterSize)>,
OpBuilder<(ins "SymbolRefAttr":$kernel, "KernelDim3":$gridSize,
"KernelDim3":$blockSize, "Value":$dynamicSharedMemorySize,
"ValueRange":$kernelOperands,
- "Type":$asyncTokenType,
+ CArg<"Type", "nullptr">:$asyncTokenType,
CArg<"ValueRange", "{}">:$asyncDependencies,
- CArg<"std::optional<KernelDim3>", "std::nullopt">:$clusterSize)>,
- OpBuilder<(ins "SymbolRefAttr":$kernel, "KernelDim3":$gridSize,
- "KernelDim3":$blockSize, "Value":$dynamicSharedMemorySize,
- "ValueRange":$kernelOperands,
CArg<"Value", "nullptr">:$asyncObject,
CArg<"std::optional<KernelDim3>", "std::nullopt">:$clusterSize)>
];
@@ -813,20 +819,17 @@ def GPU_LaunchOp : GPU_Op<"launch", [
UnitAttr:$cooperative,
OptionalAttr<FlatSymbolRefAttr>:$module,
OptionalAttr<FlatSymbolRefAttr>:$function,
- OptionalAttr<ConfinedAttr<I64Attr, [IntNonNegative]>>:$workgroup_attributions)>,
+ OptionalAttr<ConfinedAttr<I64Attr, [IntNonNegative]>>:$workgroup_attributions,
+ Optional<AnyType>:$asyncObject)>,
Results<(outs Optional<GPU_AsyncToken>:$asyncToken)> {
let summary = "GPU kernel launch operation";
let description = [{
Launch a kernel on the specified grid of thread blocks. The body of the
- kernel is defined by the single region that this operation contains. The
- operation takes an optional list of async dependencies followed by six
- operands and an optional operand.
+ kernel is defined by the single region that this operation contains.
- The `async` keyword indicates the kernel should be launched asynchronously;
- the operation returns a new !gpu.async.token when the keyword is specified.
- The kernel launched does not start executing until the ops producing its
- async dependencies (optional operands) have completed.
+ The async execution model is equivalent to the `gpu.launch_func` op, refer
+ to its description.
The first three operands (following any async dependencies) are grid sizes
along the x,y,z dimensions and the following three are block sizes along the
@@ -945,6 +948,7 @@ def GPU_LaunchOp : GPU_Op<"launch", [
CArg<"Value", "nullptr">:$dynamicSharedMemorySize,
CArg<"Type", "nullptr">:$asyncTokenType,
CArg<"ValueRange", "{}">:$asyncDependencies,
+ CArg<"Value", "nullptr">:$asyncObject,
CArg<"TypeRange", "{}">:$workgroupAttributions,
CArg<"TypeRange", "{}">:$privateAttributions,
CArg<"Value", "nullptr">:$clusterSizeX,
diff --git a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
index e0048876bf348..80b9a15ac9a41 100644
--- a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
+++ b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
@@ -1102,7 +1102,7 @@ LogicalResult LegalizeLaunchFuncOpPattern::matchAndRewrite(
adaptor.getBlockSizeZ()},
adaptor.getDynamicSharedMemorySize(),
llvmArgumentsWithSizes.empty() ? llvmArguments : llvmArgumentsWithSizes,
- stream, clusterSize);
+ nullptr, {}, stream, clusterSize);
if (launchOp.getCooperative())
newLaunchOp.setCooperative(true);
if (launchOp.getAsyncToken())
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index eaea0142e0438..22c5e3c9b86ad 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -715,7 +715,7 @@ void LaunchOp::build(OpBuilder &builder, OperationState &result,
Value getBlockSizeX, Value getBlockSizeY,
Value getBlockSizeZ, Value dynamicSharedMemorySize,
Type asyncTokenType, ValueRange asyncDependencies,
- TypeRange workgroupAttributions,
+ Value asyncObject, TypeRange workgroupAttributions,
TypeRange privateAttributions, Value clusterSizeX,
Value clusterSizeY, Value clusterSizeZ,
FlatSymbolRefAttr module, FlatSymbolRefAttr function) {
@@ -742,6 +742,8 @@ void LaunchOp::build(OpBuilder &builder, OperationState &result,
result.addOperands(clusterSizeZ);
if (dynamicSharedMemorySize)
result.addOperands(dynamicSharedMemorySize);
+ if (asyncObject)
+ result.addOperands(asyncObject);
// Add optional module and function attributes.
if (module)
@@ -763,12 +765,13 @@ void LaunchOp::build(OpBuilder &builder, OperationState &result,
for (Type argTy : privateAttributions)
body->addArgument(argTy, result.location);
// Fill OperandSegmentSize Attribute.
- SmallVector<int32_t, 11> segmentSizes(11, 1);
+ SmallVector<int32_t, 12> segmentSizes(12, 1);
segmentSizes.front() = asyncDependencies.size();
- segmentSizes.back() = dynamicSharedMemorySize ? 1 : 0;
segmentSizes[7] = clusterSizeX ? 1 : 0;
segmentSizes[8] = clusterSizeY ? 1 : 0;
segmentSizes[9] = clusterSizeZ ? 1 : 0;
+ segmentSizes[10] = dynamicSharedMemorySize ? 1 : 0;
+ segmentSizes[11] = asyncObject ? 1 : 0;
result.addAttribute(getOperandSegmentSizeAttr(),
builder.getDenseI32ArrayAttr(segmentSizes));
}
@@ -830,7 +833,23 @@ std::optional<KernelDim3> LaunchOp::getClusterSizeOperandValues() {
return KernelDim3{operands[6], operands[7], operands[8]};
}
+template <typename OpTy>
+static LogicalResult verifyLaunchAsyncModel(OpTy op) {
+ if (!op.getAsyncDependencies().empty() && !op.getAsyncToken())
+ return op.emitOpError("dependency operands require the dependency-based "
+ "async model i.e. returning a token");
+ if (op.getAsyncToken() && op.getAsyncObject())
+ return op.emitOpError("stream-based and dependency-based async models are "
+ "mutually exclusive");
+ if (op.getNumResults() == 0 && op.getAsyncToken())
+ return op.emitOpError("needs to be named when async keyword is specified");
+ return success();
+}
+
LogicalResult LaunchOp::verify() {
+ if (verifyLaunchAsyncModel(*this).failed())
+ return failure();
+
if (!(hasClusterSize()) &&
(getClusterSizeX() || getClusterSizeY() || getClusterSizeZ()))
return emitOpError() << "cluster size must be all present";
@@ -877,9 +896,6 @@ LogicalResult LaunchOp::verifyRegions() {
}
}
- if (getNumResults() == 0 && getAsyncToken())
- return emitOpError("needs to be named when async keyword is specified");
-
return success();
}
@@ -896,6 +912,9 @@ static void printSizeAssignment(OpAsmPrinter &p, KernelDim3 size,
}
void LaunchOp::print(OpAsmPrinter &p) {
+ if (auto asyncObject = getAsyncObject()) {
+ p << " <" << asyncObject << " : " << asyncObject.getType() << ">";
+ }
if (getAsyncToken()) {
p << " async";
if (!getAsyncDependencies().empty())
@@ -986,8 +1005,9 @@ parseSizeAssignment(OpAsmParser &parser,
}
/// Parses a Launch operation.
-/// operation ::= `gpu.launch` (`async` `[` ssa-id-list `]`)?
-/// `clusters` `(` ssa-id-list `)` `in` ssa-reassignment (Optional)
+/// operation ::= `gpu.launch` (`<` ssa-use `:` type `>`)?
+/// (`async` `[` ssa-id-list `]`)?
+/// (`clusters` `(` ssa-id-list `)` `in` ssa-reassignment)?
/// `blocks` `(` ssa-id-list `)` `in` ssa-reassignment
/// `threads` `(` ssa-id-list `)` `in` ssa-reassignment
/// (`dynamic_shared_memory_size` ssa-use)?
@@ -1005,6 +1025,17 @@ ParseResult LaunchOp::parse(OpAsmParser &parser, OperationState &result) {
SmallVector<OpAsmParser::UnresolvedOperand, 16> regionArgs(
LaunchOp::kNumConfigRegionAttributes);
+ // Parse optional asyncObject: < value : type >
+ OpAsmParser::UnresolvedOperand asyncObjectOperand;
+ Type asyncObjectType;
+ bool hasAsyncObject = false;
+ if (succeeded(parser.parseOptionalLess())) {
+ hasAsyncObject = true;
+ if (parser.parseOperand(asyncObjectOperand) || parser.parseColon() ||
+ parser.parseType(asyncObjectType) || parser.parseGreater())
+ return failure();
+ }
+
// Parse optional async dependencies.
SmallVector<OpAsmParser::UnresolvedOperand, 4> asyncDependencies;
Type asyncTokenType;
@@ -1067,6 +1098,11 @@ ParseResult LaunchOp::parse(OpAsmParser &parser, OperationState &result) {
return failure();
}
+ // Resolve the asyncObject operand
+ if (hasAsyncObject && parser.resolveOperand(asyncObjectOperand,
+ asyncObjectType, result.operands))
+ return failure();
+
// Parse optional module attribute.
StringRef moduleAttrName = getModuleAttrName(result.name);
if (succeeded(parser.parseOptionalKeyword(moduleAttrName))) {
@@ -1135,7 +1171,7 @@ ParseResult LaunchOp::parse(OpAsmParser &parser, OperationState &result) {
parser.parseOptionalAttrDict(result.attributes))
return failure();
- SmallVector<int32_t, 11> segmentSizes(11, 1);
+ SmallVector<int32_t, 12> segmentSizes(12, 1);
segmentSizes.front() = asyncDependencies.size();
if (!hasCluster) {
@@ -1143,7 +1179,8 @@ ParseResult LaunchOp::parse(OpAsmParser &parser, OperationState &result) {
segmentSizes[8] = 0;
segmentSizes[9] = 0;
}
- segmentSizes.back() = hasDynamicSharedMemorySize ? 1 : 0;
+ segmentSizes[10] = hasDynamicSharedMemorySize ? 1 : 0;
+ segmentSizes[11] = hasAsyncObject ? 1 : 0;
result.addAttribute(LaunchOp::getOperandSegmentSizeAttr(),
parser.getBuilder().getDenseI32ArrayAttr(segmentSizes));
return success();
@@ -1216,7 +1253,7 @@ void LaunchFuncOp::build(OpBuilder &builder, OperationState &result,
SymbolRefAttr kernelSymbol, KernelDim3 gridSize,
KernelDim3 getBlockSize, Value dynamicSharedMemorySize,
ValueRange kernelOperands, Type asyncTokenType,
- ValueRange asyncDependencies,
+ ValueRange asyncDependencies, Value asyncObject,
std::optional<KernelDim3> clusterSize) {
assert(kernelSymbol.getNestedReferences().size() == 1 &&
"expected a symbol reference with a single nested reference");
@@ -1232,6 +1269,8 @@ void LaunchFuncOp::build(OpBuilder &builder, OperationState &result,
if (dynamicSharedMemorySize)
result.addOperands(dynamicSharedMemorySize);
result.addOperands(kernelOperands);
+ if (asyncObject)
+ result.addOperands(asyncObject);
Properties &prop = result.getOrAddProperties...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/213031
More information about the flang-commits
mailing list