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) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
