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