[clang] a87510a - [CIR][SPIR-V] Update HIP atomic scope checks (#231970)

[via cfe-commits]

Author: Akimasa Watanuki
Date: 2026-10-14T18:51:39+02:00
New Revision: a87510aa418e6d876280d2bc37b009e8fa5d0999

URL: https://github.com/llvm/llvm-project/commit/a87510aa418e6d876280d2bc37b009e8fa5d0999
DIFF: https://github.com/llvm/llvm-project/commit/a87510aa418e6d876280d2bc37b009e8fa5d0999.diff

LOG: [CIR][SPIR-V] Update HIP atomic scope checks (#231970)

PR #230829 made CIR translate AMDGPU scope names to their SPIR-V
equivalents, matching classic CodeGen. PR #231754 subsequently added HIP
atomic scope coverage with checks that still expected the AMDGPU names
for wavefront and agent scopes.

Expect subgroup and device for those scopes, share the LLVM checks
between CIR and classic CodeGen, and remove the obsolete FIXME.

Assisted-by: Codex / GPT-6

Added: 
    

Modified: 
    clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip

Removed: 
    


################################################################################
diff  --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip
index d12c56d1ff686..b9339e3e0f67a 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip
@@ -5,18 +5,16 @@
 
 // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
 // RUN:            -fcuda-is-device -emit-llvm %s -o %t-cir.ll
-// RUN: FileCheck --check-prefixes=LLVM,CIR-LLVM %s --input-file=%t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t-cir.ll
 
 // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \
 // RUN:            -fcuda-is-device -emit-llvm %s -o %t.ll
-// RUN: FileCheck --check-prefixes=LLVM,OGCG %s --input-file=%t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
 
 // RUN: %clang_cc1 -triple spir64-unknown-unknown -x hip -fclangir \
 // RUN:            -fcuda-is-device -emit-cir %s -o %t-spir.cir
 // RUN: FileCheck --check-prefix=SPIR %s --input-file=%t-spir.cir
 
-// FIXME: CIR emits AMDGPU scope names on SPIR-V for agent and wavefront.
-
 #define __device__ __attribute__((device))
 
 // CIR-LABEL: @_Z12singlethreadPff
@@ -33,8 +31,7 @@ __device__ float singlethread(float *p, float v) {
 // CIR-LABEL: @_Z9wavefrontPii
 // CIR: cir.atomic.fetch add relaxed syncscope(hip_wavefront)
 // LLVM-LABEL: @_Z9wavefrontPii
-// CIR-LLVM: atomicrmw add {{.*}} syncscope("wavefront") monotonic
-// OGCG: atomicrmw add {{.*}} syncscope("subgroup") monotonic
+// LLVM: atomicrmw add {{.*}} syncscope("subgroup") monotonic
 // SPIR-LABEL: @_Z9wavefrontPii
 // SPIR: cir.atomic.fetch add relaxed syncscope(system)
 __device__ int wavefront(int *p, int v) {
@@ -56,8 +53,7 @@ __device__ float workgroup(float *p, float v) {
 // CIR-LABEL: @_Z5agentPii
 // CIR: cir.atomic.fetch add relaxed syncscope(hip_agent)
 // LLVM-LABEL: @_Z5agentPii
-// CIR-LLVM: atomicrmw add {{.*}} syncscope("agent") monotonic
-// OGCG: atomicrmw add {{.*}} syncscope("device") monotonic
+// LLVM: atomicrmw add {{.*}} syncscope("device") monotonic
 // SPIR-LABEL: @_Z5agentPii
 // SPIR: cir.atomic.fetch add relaxed syncscope(system)
 __device__ int agent(int *p, int v) {