[clang] 0c55c0b - [CIR][SPIR-V] Split SPIR and SPIR-V target lowering to keep atomic sync scopes (#231754)

[via cfe-commits]

Author: Arseniy Obolenskiy
Date: 2026-10-14T13:12:23+02:00
New Revision: 0c55c0b546998db9ad64487decb117454dc34130

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

LOG: [CIR][SPIR-V] Split SPIR and SPIR-V target lowering to keep atomic sync scopes (#231754)

Split SPIR and SPIR-V target lowering so only spir maps every sync scope
to System, like classic codegen

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

Modified: 
    clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
    clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
    clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
index 03044286e8cf71..e5df841f1e84cd 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
@@ -55,6 +55,7 @@ createTargetLoweringInfo(LowerModule &lm) {
     return createNVPTXTargetLoweringInfo();
   case llvm::Triple::spir:
   case llvm::Triple::spir64:
+    return createCommonSPIRTargetLoweringInfo();
   case llvm::Triple::spirv:
   case llvm::Triple::spirv32:
   case llvm::Triple::spirv64:

diff  --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
index ecba391c0016f6..7bdca0a7d999f5 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
@@ -38,6 +38,8 @@ std::unique_ptr<TargetLoweringInfo> createAMDGPUTargetLoweringInfo();
 
 std::unique_ptr<TargetLoweringInfo> createNVPTXTargetLoweringInfo();
 
+std::unique_ptr<TargetLoweringInfo> createCommonSPIRTargetLoweringInfo();
+
 std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo();
 
 } // namespace cir

diff  --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
index 084d98f81bce27..ff9055272239fc 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
@@ -26,7 +26,7 @@ constexpr unsigned SPIRVAddrSpaceMap[] = {
     6, // GlobalHost
 };
 
-class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
+class CommonSPIRTargetLoweringInfo : public TargetLoweringInfo {
 public:
   unsigned getTargetAddrSpaceFromCIRAddrSpace(
       cir::LangAddressSpace addrSpace) const override {
@@ -40,8 +40,20 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
   }
 };
 
+class SPIRVTargetLoweringInfo : public CommonSPIRTargetLoweringInfo {
+public:
+  cir::SyncScopeKind
+  convertSyncScope(cir::SyncScopeKind syncScope) const override {
+    return syncScope;
+  }
+};
+
 } // namespace
 
+std::unique_ptr<TargetLoweringInfo> createCommonSPIRTargetLoweringInfo() {
+  return std::make_unique<CommonSPIRTargetLoweringInfo>();
+}
+
 std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo() {
   return std::make_unique<SPIRVTargetLoweringInfo>();
 }

diff  --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip
new file mode 100644
index 00000000000000..d12c56d1ff6864
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip
@@ -0,0 +1,77 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// 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: %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: %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
+// CIR: cir.atomic.fetch add relaxed syncscope(hip_single_thread)
+// LLVM-LABEL: @_Z12singlethreadPff
+// LLVM: atomicrmw fadd {{.*}} syncscope("singlethread") monotonic
+// SPIR-LABEL: @_Z12singlethreadPff
+// SPIR: cir.atomic.fetch add relaxed syncscope(system)
+__device__ float singlethread(float *p, float v) {
+  return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED,
+                                __HIP_MEMORY_SCOPE_SINGLETHREAD);
+}
+
+// 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
+// SPIR-LABEL: @_Z9wavefrontPii
+// SPIR: cir.atomic.fetch add relaxed syncscope(system)
+__device__ int wavefront(int *p, int v) {
+  return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED,
+                                __HIP_MEMORY_SCOPE_WAVEFRONT);
+}
+
+// CIR-LABEL: @_Z9workgroupPff
+// CIR: cir.atomic.fetch add relaxed syncscope(hip_workgroup)
+// LLVM-LABEL: @_Z9workgroupPff
+// LLVM: atomicrmw fadd {{.*}} syncscope("workgroup") monotonic
+// SPIR-LABEL: @_Z9workgroupPff
+// SPIR: cir.atomic.fetch add relaxed syncscope(system)
+__device__ float workgroup(float *p, float v) {
+  return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED,
+                                __HIP_MEMORY_SCOPE_WORKGROUP);
+}
+
+// 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
+// SPIR-LABEL: @_Z5agentPii
+// SPIR: cir.atomic.fetch add relaxed syncscope(system)
+__device__ int agent(int *p, int v) {
+  return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED,
+                                __HIP_MEMORY_SCOPE_AGENT);
+}
+
+// CIR-LABEL: @_Z6systemPff
+// CIR: cir.atomic.fetch add relaxed syncscope(hip_system)
+// LLVM-LABEL: @_Z6systemPff
+// LLVM: atomicrmw fadd {{.*}} monotonic
+// SPIR-LABEL: @_Z6systemPff
+// SPIR: cir.atomic.fetch add relaxed syncscope(system)
+__device__ float system(float *p, float v) {
+  return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED,
+                                __HIP_MEMORY_SCOPE_SYSTEM);
+}