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