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

[Mehdi Amini]

https://github.com/joker-eph created https://github.com/llvm/llvm-project/pull/188729

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

>From 4559556cb22e319958e6ed5f34b86b6405c72138 Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Wed, 25 Mar 2026 16:18:26 -0700
Subject: [PATCH] [MLIR][GPU] Add verifier to gpu.barrier requiring kernel
 context

gpu.barrier is only meaningful inside a GPU kernel body but had no verifier
enforcing this. When gpu.barrier appears at the top level of a gpu.module
(outside any gpu.func or gpu.launch), --convert-gpu-to-nvvm converts it to
nvvm.barrier0 without error. Subsequent translation to LLVM IR via
--gpu-module-to-binary then crashes because createIntrinsicCall is invoked
with an IRBuilder that has no valid insertion point (not inside any LLVM IR
function), causing a null pointer dereference in BasicBlock::getModule().

Fix by adding 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
---
 mlir/include/mlir/Dialect/GPU/IR/GPUOps.td |  1 +
 mlir/lib/Dialect/GPU/IR/GPUDialect.cpp     | 11 +++++++++++
 mlir/test/Dialect/GPU/invalid.mlir         | 10 ++++++++++
 3 files changed, 22 insertions(+)

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
+}