[Mlir-commits] [mlir] 5829652 - [MLIR][NVVM] Add clusterlaunchcontrol Ops (#156585)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Sep 17 00:19:16 PDT 2025
Author: Srinivasa Ravi
Date: 2025-09-17T12:49:12+05:30
New Revision: 5829652a6def4031853104b43ee6dfd331162352
URL: https://github.com/llvm/llvm-project/commit/5829652a6def4031853104b43ee6dfd331162352
DIFF: https://github.com/llvm/llvm-project/commit/5829652a6def4031853104b43ee6dfd331162352.diff
LOG: [MLIR][NVVM] Add clusterlaunchcontrol Ops (#156585)
This change adds the `clusterlaunchcontrol.try.cancel` and
`clusterlaunchcontrol.query.cancel` Ops to the NVVM dialect.
Tests are added in `clusterlaunchcontrol.mlir`.
PTX Reference:
https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-try-cancel
Added:
mlir/test/Target/LLVMIR/nvvm/clusterlaunchcontrol.mlir
Modified:
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 70ab1df876d35..797f8ada9f238 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -4431,6 +4431,116 @@ def NVVM_DotAccumulate2WayOp : NVVM_Op<"dot.accumulate.2way"> {
}];
}
+//===----------------------------------------------------------------------===//
+// NVVM clusterlaunchcontrol Ops.
+//===----------------------------------------------------------------------===//
+
+def NVVM_ClusterLaunchControlTryCancelOp
+ : NVVM_Op<"clusterlaunchcontrol.try.cancel", [NVVMRequiresSM<100>]> {
+ let summary = "Request atomically canceling the launch of a cluster that has not started running yet";
+ let description = [{
+ `clusterlaunchcontrol.try.cancel` requests atomically canceling the launch
+ of a cluster that has not started running yet. It asynchronously writes an
+ opaque response to shared memory indicating whether the operation succeeded
+ or failed.
+
+ Operand `smemAddress` specifies the naturally aligned address of the
+ 16-byte wide shared memory location where the request's response is written.
+
+ Operand `mbarrier` specifies the mbarrier object used to track the
+ completion of the asynchronous operation.
+
+ If `multicast` is specified, the response is asynchronously written to the
+ corresponding local shared memory location (specifed by `addr`) of each CTA
+ in the requesting cluster.
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-try-cancel)
+ }];
+
+ let arguments = (ins UnitAttr:$multicast,
+ LLVM_PointerShared: $smemAddress,
+ LLVM_PointerShared: $mbarrier);
+
+ let assemblyFormat = "(`multicast` $multicast^ `,`)? $smemAddress `,` $mbarrier attr-dict";
+
+ let extraClassDeclaration = [{
+ static mlir::NVVM::IDArgPair
+ getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
+ llvm::IRBuilderBase &builder);
+ }];
+
+ string llvmBuilder = [{
+ auto [id, args] =
+ NVVM::ClusterLaunchControlTryCancelOp::getIntrinsicIDAndArgs(
+ *op, moduleTranslation, builder);
+ createIntrinsicCall(builder, id, args);
+ }];
+}
+
+def ClusterLaunchControlIsCanceled
+ : I32EnumCase<"IS_CANCELED", 0, "is_canceled">;
+def ClusterLaunchControlGetFirstCTAIDX
+ : I32EnumCase<"GET_FIRST_CTA_ID_X", 1, "get_first_cta_id_x">;
+def ClusterLaunchControlGetFirstCTAIDY
+ : I32EnumCase<"GET_FIRST_CTA_ID_Y", 2, "get_first_cta_id_y">;
+def ClusterLaunchControlGetFirstCTAIDZ
+ : I32EnumCase<"GET_FIRST_CTA_ID_Z", 3, "get_first_cta_id_z">;
+
+def ClusterLaunchControlQueryType
+ : I32Enum<"ClusterLaunchControlQueryType",
+ "NVVM ClusterLaunchControlQueryType",
+ [ClusterLaunchControlIsCanceled, ClusterLaunchControlGetFirstCTAIDX,
+ ClusterLaunchControlGetFirstCTAIDY, ClusterLaunchControlGetFirstCTAIDZ]> {
+ let cppNamespace = "::mlir::NVVM";
+}
+
+def ClusterLaunchControlQueryTypeAttr
+ : EnumAttr<NVVM_Dialect,
+ ClusterLaunchControlQueryType, "cluster_launch_control_query_type"> {
+ let assemblyFormat = "$value";
+}
+
+def NVVM_ClusterLaunchControlQueryCancelOp
+ : NVVM_Op<"clusterlaunchcontrol.query.cancel", [NVVMRequiresSM<100>]> {
+ let summary = "Query the response of a clusterlaunchcontrol.try.cancel operation";
+ let description = [{
+ `clusterlaunchcontrol.query.cancel` queries the response of a
+ `clusterlaunchcontrol.try.cancel` operation specified by operand
+ `try_cancel_response`.
+
+ Operand `query_type` specifies the type of query to perform and can be one
+ of the following:
+ - `is_canceled` : Returns true if the try cancel request succeeded,
+ and false otherwise.
+ - `get_first_cta_id_{x/y/z}` : Returns the x, y, or z coordinate of the
+ first CTA in the canceled cluster. Behaviour is defined only if the try
+ cancel request succeeded.
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-query-cancel)
+ }];
+
+ let arguments = (ins ClusterLaunchControlQueryTypeAttr:$query_type,
+ I128:$try_cancel_response);
+ let results = (outs AnyTypeOf<[I1, I32]>:$res);
+
+ let assemblyFormat = "`query` `=` $query_type `,` $try_cancel_response attr-dict `:` type($res)";
+
+ let hasVerifier = 1;
+
+ let extraClassDeclaration = [{
+ static mlir::NVVM::IDArgPair
+ getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
+ llvm::IRBuilderBase &builder);
+ }];
+
+ string llvmBuilder = [{
+ auto [id, args] =
+ NVVM::ClusterLaunchControlQueryCancelOp::getIntrinsicIDAndArgs(
+ *op, moduleTranslation, builder);
+ $res = createIntrinsicCall(builder, id, args);
+ }];
+}
+
//===----------------------------------------------------------------------===//
// NVVM target attribute.
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 9dba0848eba94..13f1dd9a664e5 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1401,6 +1401,24 @@ LogicalResult NVVM::PrefetchOp::verify() {
return success();
}
+LogicalResult NVVM::ClusterLaunchControlQueryCancelOp::verify() {
+ switch (getQueryType()) {
+ case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
+ if (!getType().isInteger(1))
+ return emitOpError("is_canceled query type returns an i1");
+ break;
+ case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
+ case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
+ case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
+ if (!getType().isInteger(32)) {
+ return emitOpError("get_first_cta_id_x, get_first_cta_id_y, "
+ "get_first_cta_id_z query types return an i32");
+ }
+ break;
+ }
+ return success();
+}
+
/// Packs the given `field` into the `result`.
/// The `result` is 64-bits and each `field` can be 32-bits or narrower.
static llvm::Value *
@@ -2087,6 +2105,51 @@ bool NVVM::InlinePtxOp::getAsmValues(
return false; // No manual mapping needed
}
+NVVM::IDArgPair ClusterLaunchControlTryCancelOp::getIntrinsicIDAndArgs(
+ Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+ auto curOp = cast<NVVM::ClusterLaunchControlTryCancelOp>(op);
+ llvm::SmallVector<llvm::Value *> args;
+ args.push_back(mt.lookupValue(curOp.getSmemAddress()));
+ args.push_back(mt.lookupValue(curOp.getMbarrier()));
+
+ llvm::Intrinsic::ID intrinsicID =
+ curOp.getMulticast()
+ ? llvm::Intrinsic::
+ nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
+ : llvm::Intrinsic::nvvm_clusterlaunchcontrol_try_cancel_async_shared;
+
+ return {intrinsicID, args};
+}
+
+NVVM::IDArgPair ClusterLaunchControlQueryCancelOp::getIntrinsicIDAndArgs(
+ Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+ auto curOp = cast<NVVM::ClusterLaunchControlQueryCancelOp>(op);
+ llvm::SmallVector<llvm::Value *> args;
+ args.push_back(mt.lookupValue(curOp.getTryCancelResponse()));
+
+ llvm::Intrinsic::ID intrinsicID;
+
+ switch (curOp.getQueryType()) {
+ case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
+ intrinsicID =
+ llvm::Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled;
+ break;
+ case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
+ intrinsicID = llvm::Intrinsic::
+ nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x;
+ break;
+ case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
+ intrinsicID = llvm::Intrinsic::
+ nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y;
+ break;
+ case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
+ intrinsicID = llvm::Intrinsic::
+ nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z;
+ break;
+ }
+ return {intrinsicID, args};
+}
+
//===----------------------------------------------------------------------===//
// NVVMDialect initialization, type parsing, and registration.
//===----------------------------------------------------------------------===//
diff --git a/mlir/test/Target/LLVMIR/nvvm/clusterlaunchcontrol.mlir b/mlir/test/Target/LLVMIR/nvvm/clusterlaunchcontrol.mlir
new file mode 100644
index 0000000000000..6ff85fa38a0a0
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/clusterlaunchcontrol.mlir
@@ -0,0 +1,55 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+llvm.func @clusterlaunchcontrol_try_cancel(%addr: !llvm.ptr<3>, %mbar: !llvm.ptr<3>) {
+ // CHECK-LABEL: define void @clusterlaunchcontrol_try_cancel(ptr addrspace(3) %0, ptr addrspace(3) %1) {
+ // CHECK-NEXT: call void @llvm.nvvm.clusterlaunchcontrol.try_cancel.async.shared(ptr addrspace(3) %0, ptr addrspace(3) %1)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ nvvm.clusterlaunchcontrol.try.cancel %addr, %mbar
+ llvm.return
+}
+
+llvm.func @clusterlaunchcontrol_try_cancel_multicast(%addr: !llvm.ptr<3>, %mbar: !llvm.ptr<3>) {
+ // CHECK-LABEL: define void @clusterlaunchcontrol_try_cancel_multicast(ptr addrspace(3) %0, ptr addrspace(3) %1) {
+ // CHECK-NEXT: call void @llvm.nvvm.clusterlaunchcontrol.try_cancel.async.multicast.shared(ptr addrspace(3) %0, ptr addrspace(3) %1)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ nvvm.clusterlaunchcontrol.try.cancel multicast, %addr, %mbar
+ llvm.return
+}
+
+llvm.func @clusterlaunchcontrol_query_cancel_is_canceled(%try_cancel_response: i128) {
+ // CHECK-LABEL: define void @clusterlaunchcontrol_query_cancel_is_canceled(i128 %0) {
+ // CHECK-NEXT: %2 = call i1 @llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled(i128 %0)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ %res = nvvm.clusterlaunchcontrol.query.cancel query = is_canceled, %try_cancel_response : i1
+ llvm.return
+}
+
+llvm.func @clusterlaunchcontrol_query_cancel_get_first_cta_id_x(%try_cancel_response: i128) {
+ // CHECK-LABEL: define void @clusterlaunchcontrol_query_cancel_get_first_cta_id_x(i128 %0) {
+ // CHECK-NEXT: %2 = call i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.x(i128 %0)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ %res = nvvm.clusterlaunchcontrol.query.cancel query = get_first_cta_id_x, %try_cancel_response : i32
+ llvm.return
+}
+
+llvm.func @clusterlaunchcontrol_query_cancel_get_first_cta_id_y(%try_cancel_response: i128) {
+ // CHECK-LABEL: define void @clusterlaunchcontrol_query_cancel_get_first_cta_id_y(i128 %0) {
+ // CHECK-NEXT: %2 = call i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.y(i128 %0)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ %res = nvvm.clusterlaunchcontrol.query.cancel query = get_first_cta_id_y, %try_cancel_response : i32
+ llvm.return
+}
+
+llvm.func @clusterlaunchcontrol_query_cancel_get_first_cta_id_z(%try_cancel_response: i128) {
+ // CHECK-LABEL: define void @clusterlaunchcontrol_query_cancel_get_first_cta_id_z(i128 %0) {
+ // CHECK-NEXT: %2 = call i32 @llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid.z(i128 %0)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ %res = nvvm.clusterlaunchcontrol.query.cancel query = get_first_cta_id_z, %try_cancel_response : i32
+ llvm.return
+}
diff --git a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
index b35a6dbcca286..383f4829f3287 100644
--- a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
@@ -535,3 +535,19 @@ llvm.func @nanosleep() {
nvvm.nanosleep 100000000000000
llvm.return
}
+
+// -----
+
+llvm.func @clusterlaunchcontrol_query_cancel_is_canceled_invalid_return_type(%try_cancel_response: i128) {
+ // expected-error at +1 {{'nvvm.clusterlaunchcontrol.query.cancel' op is_canceled query type returns an i1}}
+ %res = nvvm.clusterlaunchcontrol.query.cancel query = is_canceled, %try_cancel_response : i32
+ llvm.return
+}
+
+// -----
+
+llvm.func @clusterlaunchcontrol_query_cancel_get_first_cta_id_invalid_return_type(%try_cancel_response: i128) {
+ // expected-error at +1 {{'nvvm.clusterlaunchcontrol.query.cancel' op get_first_cta_id_x, get_first_cta_id_y, get_first_cta_id_z query types return an i32}}
+ %res = nvvm.clusterlaunchcontrol.query.cancel query = get_first_cta_id_x, %try_cancel_response : i1
+ llvm.return
+}
More information about the Mlir-commits
mailing list