[Mlir-commits] [mlir] [mlir][gpu][SPIR-V] Emit diagnostic for gpu.barrier with unsupported cluster scope (PR #214458)

[Arseniy Obolenskiy]

https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/214458

>From 12513f32313fecab5239f614cecaba30a51b1611 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 6 Aug 2026 13:34:43 +0200
Subject: [PATCH] [mlir][SPIR-V] Emit diagnostic for gpu.barrier with
 unsupported cluster scope

---
 mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp  |  8 +++++---
 .../Conversion/GPUToSPIRV/gpu-to-spirv.mlir    | 18 ++++++++++++++++++
 2 files changed, 23 insertions(+), 3 deletions(-)

diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 609dbe57a9a46..03a7c4b4a2daa 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -471,9 +471,11 @@ LogicalResult GPUBarrierConversion::matchAndRewrite(
 
   // Map GPU scope to SPIR-V scope.
   auto spirvScope = mapGPUBarrierScopeToSPIRV(barrierOp.getScope());
-  if (failed(spirvScope))
-    return rewriter.notifyMatchFailure(
-        barrierOp, "cluster scope is not supported in SPIR-V");
+  if (failed(spirvScope)) {
+    barrierOp.emitError()
+        << "gpu.barrier with cluster scope is not supported in SPIR-V";
+    return failure();
+  }
 
   auto scopeAttr = spirv::ScopeAttr::get(context, *spirvScope);
   auto memoryScopeAttr =
diff --git a/mlir/test/Conversion/GPUToSPIRV/gpu-to-spirv.mlir b/mlir/test/Conversion/GPUToSPIRV/gpu-to-spirv.mlir
index 5bcd180a37c16..0bf547f225cb0 100644
--- a/mlir/test/Conversion/GPUToSPIRV/gpu-to-spirv.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/gpu-to-spirv.mlir
@@ -149,6 +149,24 @@ module attributes {
 
 // -----
 
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<
+    #spirv.vce<v1.1, [Kernel, Addresses, NamedBarrier], []>, #spirv.resource_limits<>>
+} {
+  gpu.module @kernels {
+    gpu.func @barrier_cluster_scope() kernel
+      attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
+      // expected-error @+2 {{gpu.barrier with cluster scope is not supported in SPIR-V}}
+      // expected-error @+1 {{failed to legalize operation 'gpu.barrier'}}
+      gpu.barrier scope <cluster>
+      gpu.return
+    }
+  }
+}
+
+// -----
+
 module attributes {
   gpu.container_module,
   spirv.target_env = #spirv.target_env<