[Mlir-commits] [mlir] [MLIR][GPU] Add verifier to gpu.barrier requiring kernel context (PR #188729)

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Mehdi Amini (joker-eph)

<details>
<summary>Changes</summary>

gpu.barrier is only meaningful inside a GPU kernel body but had no verifier enforcing this. Add a verifier to gpu.barrier that requires it to be nested inside either a gpu.func or gpu.launch, both of which establish a GPU kernel execution context.

Fixes #<!-- -->186642

Assisted-by: Claude Code

---
Full diff: https://github.com/llvm/llvm-project/pull/188729.diff


3 Files Affected:

- (modified) mlir/include/mlir/Dialect/GPU/IR/GPUOps.td (+1) 
- (modified) mlir/lib/Dialect/GPU/IR/GPUDialect.cpp (+11) 
- (modified) mlir/test/Dialect/GPU/invalid.mlir (+10) 


``````````diff
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index 36e0875f53b0a..ce8f234c140bc 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -1467,6 +1467,7 @@ def GPU_BarrierOp : GPU_Op<"barrier">,
   }];
   let assemblyFormat = "(`memfence` $address_spaces^)? attr-dict";
   let hasCanonicalizer = 1;
+  let hasVerifier = 1;
   let builders = [OpBuilder<(
                       ins CArg<"std::optional<::mlir::gpu::AddressSpace>",
                                "std::nullopt">:$addressSpace)>,
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index 5d409f71847c6..ffc3a489932c4 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -1529,6 +1529,17 @@ void BarrierOp::getCanonicalizationPatterns(RewritePatternSet &results,
   results.add(eraseRedundantGpuBarrierOps);
 }
 
+LogicalResult BarrierOp::verify() {
+  // gpu.barrier is only meaningful inside a GPU kernel body. Require that it
+  // be nested (directly or indirectly) inside either a gpu.func or a
+  // gpu.launch, both of which establish a kernel execution context.
+  Operation *op = getOperation();
+  if (!op->getParentOfType<GPUFuncOp>() && !op->getParentOfType<LaunchOp>())
+    return emitOpError(
+        "expected to be nested inside a gpu.func or gpu.launch region");
+  return success();
+}
+
 void BarrierOp::build(mlir::OpBuilder &odsBuilder,
                       mlir::OperationState &odsState,
                       std::optional<AddressSpace> addressSpace) {
diff --git a/mlir/test/Dialect/GPU/invalid.mlir b/mlir/test/Dialect/GPU/invalid.mlir
index bf862b2c5ae3c..d6f61363aac10 100644
--- a/mlir/test/Dialect/GPU/invalid.mlir
+++ b/mlir/test/Dialect/GPU/invalid.mlir
@@ -1119,3 +1119,13 @@ func.func @warp_execute_wrong_terminator() {
   }
   return
 }
+
+// -----
+
+// Regression test for https://github.com/llvm/llvm-project/issues/186642:
+// gpu.barrier outside a gpu.func or gpu.launch used to crash during LLVM IR
+// translation. The verifier should now emit a proper error instead.
+gpu.module @kernels {
+  // expected-error at +1 {{'gpu.barrier' op expected to be nested inside a gpu.func or gpu.launch region}}
+  gpu.barrier
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/188729