[clang] [CIR][SPIR-V] Keep atomic sync scopes (PR #231754)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 14 00:56:12 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/231754
>From 61b036d241af76e088cbbcc4da1fd668d3939040 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 14 Oct 2026 07:37:40 +0200
Subject: [PATCH 1/2] [CIR][SPIR-V] Keep atomic sync scopes
TargetLowering turned every scope into System on spirv targets
Keep System only for spir, like classic codegen
---
.../Transforms/TargetLowering/LowerModule.cpp | 2 +-
.../TargetLowering/TargetLoweringInfo.h | 2 +-
.../TargetLowering/Targets/SPIRV.cpp | 16 +++-
.../CodeGenHIP/amdgcnspirv-atomic-scope.hip | 77 +++++++++++++++++++
4 files changed, 93 insertions(+), 4 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/amdgcnspirv-atomic-scope.hip
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
index 03044286e8cf7..391d5723f73a1 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
@@ -58,7 +58,7 @@ createTargetLoweringInfo(LowerModule &lm) {
case llvm::Triple::spirv:
case llvm::Triple::spirv32:
case llvm::Triple::spirv64:
- return createSPIRVTargetLoweringInfo();
+ return createSPIRVTargetLoweringInfo(triple.isSPIRV());
default:
assert(!cir::MissingFeatures::targetLoweringInfo());
return std::make_unique<TargetLoweringInfo>();
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
index ecba391c0016f..54931d4b30a8e 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
@@ -38,7 +38,7 @@ std::unique_ptr<TargetLoweringInfo> createAMDGPUTargetLoweringInfo();
std::unique_ptr<TargetLoweringInfo> createNVPTXTargetLoweringInfo();
-std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo();
+std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo(bool isSPIRV);
} // 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 084d98f81bce2..0ec89939f57fe 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
@@ -27,7 +27,12 @@ constexpr unsigned SPIRVAddrSpaceMap[] = {
};
class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
+ // SPIR keeps the default system scope, like CommonSPIRTargetCodeGenInfo.
+ bool isSPIRV;
+
public:
+ explicit SPIRVTargetLoweringInfo(bool isSPIRV) : isSPIRV(isSPIRV) {}
+
unsigned getTargetAddrSpaceFromCIRAddrSpace(
cir::LangAddressSpace addrSpace) const override {
// TODO(cir): SYCL and CUDA/HIP device code map Default to Generic
@@ -38,12 +43,19 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
"Unknown CIR address space for SPIR-V target");
return SPIRVAddrSpaceMap[idx];
}
+
+ cir::SyncScopeKind
+ convertSyncScope(cir::SyncScopeKind syncScope) const override {
+ return isSPIRV ? syncScope
+ : TargetLoweringInfo::convertSyncScope(syncScope);
+ }
};
} // namespace
-std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo() {
- return std::make_unique<SPIRVTargetLoweringInfo>();
+std::unique_ptr<TargetLoweringInfo>
+createSPIRVTargetLoweringInfo(bool isSPIRV) {
+ return std::make_unique<SPIRVTargetLoweringInfo>(isSPIRV);
}
} // namespace cir
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 0000000000000..d12c56d1ff686
--- /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);
+}
>From e7303076d9fa75f328ba2dda51dcc36f2ed67bef Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 14 Oct 2026 09:55:59 +0200
Subject: [PATCH 2/2] split classes
---
.../Transforms/TargetLowering/LowerModule.cpp | 3 ++-
.../TargetLowering/TargetLoweringInfo.h | 4 +++-
.../TargetLowering/Targets/SPIRV.cpp | 22 +++++++++----------
3 files changed, 16 insertions(+), 13 deletions(-)
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
index 391d5723f73a1..e5df841f1e84c 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp
@@ -55,10 +55,11 @@ 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:
- return createSPIRVTargetLoweringInfo(triple.isSPIRV());
+ return createSPIRVTargetLoweringInfo();
default:
assert(!cir::MissingFeatures::targetLoweringInfo());
return std::make_unique<TargetLoweringInfo>();
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
index 54931d4b30a8e..7bdca0a7d999f 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h
@@ -38,7 +38,9 @@ std::unique_ptr<TargetLoweringInfo> createAMDGPUTargetLoweringInfo();
std::unique_ptr<TargetLoweringInfo> createNVPTXTargetLoweringInfo();
-std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo(bool isSPIRV);
+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 0ec89939f57fe..ff9055272239f 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
@@ -26,13 +26,8 @@ constexpr unsigned SPIRVAddrSpaceMap[] = {
6, // GlobalHost
};
-class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
- // SPIR keeps the default system scope, like CommonSPIRTargetCodeGenInfo.
- bool isSPIRV;
-
+class CommonSPIRTargetLoweringInfo : public TargetLoweringInfo {
public:
- explicit SPIRVTargetLoweringInfo(bool isSPIRV) : isSPIRV(isSPIRV) {}
-
unsigned getTargetAddrSpaceFromCIRAddrSpace(
cir::LangAddressSpace addrSpace) const override {
// TODO(cir): SYCL and CUDA/HIP device code map Default to Generic
@@ -43,19 +38,24 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
"Unknown CIR address space for SPIR-V target");
return SPIRVAddrSpaceMap[idx];
}
+};
+class SPIRVTargetLoweringInfo : public CommonSPIRTargetLoweringInfo {
+public:
cir::SyncScopeKind
convertSyncScope(cir::SyncScopeKind syncScope) const override {
- return isSPIRV ? syncScope
- : TargetLoweringInfo::convertSyncScope(syncScope);
+ return syncScope;
}
};
} // namespace
-std::unique_ptr<TargetLoweringInfo>
-createSPIRVTargetLoweringInfo(bool isSPIRV) {
- return std::make_unique<SPIRVTargetLoweringInfo>(isSPIRV);
+std::unique_ptr<TargetLoweringInfo> createCommonSPIRTargetLoweringInfo() {
+ return std::make_unique<CommonSPIRTargetLoweringInfo>();
+}
+
+std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo() {
+ return std::make_unique<SPIRVTargetLoweringInfo>();
}
} // namespace cir
More information about the cfe-commits
mailing list