From 759382e2cd1c8c900cd6b61e54fdcf7f58318639 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Thu, 5 Oct 2023 12:17:58 +0100 Subject: [PATCH 01/13] Add AMDGPU reflect pass for atomic xor AMDGPU reflect pass is needed to choose between safe and unsafe atomics at the libclc level. In the long run we will delete this patch as work is being done to ensure correct lowering of atomic instructions. See patches: https://github.com/llvm/llvm-project/pull/85052/ https://github.com/llvm/llvm-project/pull/69229/ This work is necessary as malloc shared atomics rely on PCIe atomics which can have patchy and unreliable support. We want to therefore be able to choose at compile time whether we should use safe atomics using CAS (which PCIe should support), or if we want to rely of the availability of the newest PCIe atomics, if malloc shared atomics are desired. Also changes the implementation of Or, And so that they can choose between the safe or unsafe version based on the AMDGPU reflect value. --- libclc/CMakeLists.txt | 9 +- .../libspirv/atomic/atomic_and.cl | 13 +-- .../libspirv/atomic/atomic_helpers.h | 28 ++++-- .../libspirv/atomic/atomic_or.cl | 13 +-- .../libspirv/atomic/atomic_safe.def | 9 ++ .../libspirv/atomic/atomic_xor.cl | 16 +--- llvm/lib/Target/AMDGPU/AMDGPU.h | 5 + llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp | 91 +++++++++++++++++++ .../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 1 + llvm/lib/Target/AMDGPU/CMakeLists.txt | 1 + .../CodeGen/AMDGPU/amdgpu-oclc-reflect.ll | 17 ++++ 11 files changed, 166 insertions(+), 37 deletions(-) create mode 100644 libclc/amdgcn-amdhsa/libspirv/atomic/atomic_safe.def create mode 100644 llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp create mode 100644 llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index ef29824b8416a..468b03d7403cf 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -312,6 +312,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) # Disables NVVM reflection to defer to after linking list( APPEND flags -Xclang -target-feature -Xclang +ptx72 -march=sm_86 -mllvm --nvvm-reflect-enable=false) + elseif( ARCH STREQUAL amdgcn ) + # AMDGCN needs libclc to be compiled to high bc version since all atomic + # clang builtins need to be accessible + list( APPEND flags -mcpu=gfx940 -mllvm --amdgpu-oclc-reflect-enable=false) elseif( ARCH STREQUAL x86_64) # TODO: This is used by SYCL Native Cpu, we should define an option to set this flags list( APPEND flags -Xclang -target-feature -Xclang +avx @@ -342,7 +346,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) # simultaneously, we choose declare the builtins using the private space, # which will also work for the generic address space. set( supports_generic_addrspace FALSE ) - elseif( ARCH STREQUAL x86_64) + elseif( ARCH STREQUAL "amdgcn" ) + set( build_flags ) + set( opt_flags -O3 --amdgpu-oclc-reflect-enable=false ) + elseif( ARCH STREQUAL "x86_64") set( opt_flags ) set( supports_generic_addrspace FALSE ) else() diff --git a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_and.cl b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_and.cl index 8d0734d7f4d2d..88c3031046c23 100644 --- a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_and.cl +++ b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_and.cl @@ -10,13 +10,8 @@ #include #include -AMDGPU_ATOMIC(_Z17__spirv_AtomicAnd, int, i, __hip_atomic_fetch_and) -AMDGPU_ATOMIC(_Z17__spirv_AtomicAnd, unsigned int, j, __hip_atomic_fetch_and) -AMDGPU_ATOMIC(_Z17__spirv_AtomicAnd, long, l, __hip_atomic_fetch_and) -AMDGPU_ATOMIC(_Z17__spirv_AtomicAnd, unsigned long, m, __hip_atomic_fetch_and) +#define __CLC_OP & +#define __SPIRV_BUILTIN _Z17__spirv_AtomicAnd +#define __HIP_BUILTIN __hip_atomic_fetch_and -#undef AMDGPU_ATOMIC -#undef AMDGPU_ATOMIC_IMPL -#undef AMDGPU_ARCH_GEQ -#undef AMDGPU_ARCH_BETWEEN -#undef GET_ATOMIC_SCOPE_AND_ORDER +#include "atomic_safe.def" diff --git a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h index ea4e90c0ae3aa..6d33d98b8810a 100644 --- a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h +++ b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h @@ -9,6 +9,8 @@ #include #include +extern int __oclc_amdgpu_reflect(__constant char *); + #define AMDGPU_ARCH_GEQ(LOWER) __oclc_ISA_version >= LOWER #define AMDGPU_ARCH_BETWEEN(LOWER, UPPER) \ __oclc_ISA_version >= LOWER &&__oclc_ISA_version < UPPER @@ -72,14 +74,22 @@ AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, local, U3AS3, 1, BUILTIN) \ AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, , , 0, BUILTIN) -#define AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, AS, AS_MANGLED, \ - SUB1, OP) \ +// Safe atomics will either choose a slow CAS atomic impl (default) or a fast +// native atomic if --amdgpu-unsafe-int-atomics is passed to LLVM. +// +// Safe atomics using CAS may be necessary if PCIe does not support atomic +// operations such as and, or, xor +#define AMDGPU_SAFE_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, AS, AS_MANGLED, \ + SUB1, OP, USE_BUILTIN_COND, BUILTIN) \ _CLC_DEF TYPE \ FUNC_NAME##P##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ volatile AS TYPE *p, enum Scope scope, \ enum MemorySemanticsMask semantics, TYPE val) { \ int atomic_scope = 0, memory_order = 0; \ GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \ + if (USE_BUILTIN_COND) \ + return BUILTIN(p, val, memory_order, atomic_scope); \ + /* CAS atomics*/ \ TYPE oldval = __hip_atomic_load(p, memory_order, atomic_scope); \ TYPE newval = 0; \ do { \ @@ -89,7 +99,13 @@ return oldval; \ } -#define AMDGPU_CAS_ATOMIC(FUNC_NAME, TYPE, TYPE_MANGLED, OP) \ - AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, global, U3AS1, 1, OP) \ - AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, local, U3AS3, 1, OP) \ - AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, , , 0, OP) +#define AMDGPU_SAFE_ATOMIC(FUNC_NAME, TYPE, TYPE_MANGLED, OP, BUILTIN) \ + AMDGPU_SAFE_ATOMIC_IMPL( \ + FUNC_NAME, TYPE, TYPE_MANGLED, global, U3AS1, 1, OP, \ + __oclc_amdgpu_reflect("AMDGPU_OCLC_UNSAFE_INT_ATOMICS"), BUILTIN) \ + AMDGPU_SAFE_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, local, U3AS3, 1, OP, \ + true /* local AS should always use builtin*/, \ + BUILTIN) \ + AMDGPU_SAFE_ATOMIC_IMPL( \ + FUNC_NAME, TYPE, TYPE_MANGLED, , , 0, OP, \ + __oclc_amdgpu_reflect("AMDGPU_OCLC_UNSAFE_INT_ATOMICS"), BUILTIN) diff --git a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_or.cl b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_or.cl index 829eb5deba6e5..283802601bc84 100644 --- a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_or.cl +++ b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_or.cl @@ -10,13 +10,8 @@ #include #include -AMDGPU_ATOMIC(_Z16__spirv_AtomicOr, int, i, __hip_atomic_fetch_or) -AMDGPU_ATOMIC(_Z16__spirv_AtomicOr, unsigned int, j, __hip_atomic_fetch_or) -AMDGPU_ATOMIC(_Z16__spirv_AtomicOr, long, l, __hip_atomic_fetch_or) -AMDGPU_ATOMIC(_Z16__spirv_AtomicOr, unsigned long, m, __hip_atomic_fetch_or) +#define __CLC_OP | +#define __SPIRV_BUILTIN _Z16__spirv_AtomicOr +#define __HIP_BUILTIN __hip_atomic_fetch_or -#undef AMDGPU_ATOMIC -#undef AMDGPU_ATOMIC_IMPL -#undef AMDGPU_ARCH_GEQ -#undef AMDGPU_ARCH_BETWEEN -#undef GET_ATOMIC_SCOPE_AND_ORDER +#include "atomic_safe.def" diff --git a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_safe.def b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_safe.def new file mode 100644 index 0000000000000..fb2024869615b --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_safe.def @@ -0,0 +1,9 @@ +// Before including, define: __SPIRV_BUILTIN, __CLC_OP, __HIP_BUILTIN +// and include atomic_helpers.h to get AMDGPU_SAFE_ATOMIC + +AMDGPU_SAFE_ATOMIC(__SPIRV_BUILTIN, int, i, __CLC_OP, __HIP_BUILTIN) +AMDGPU_SAFE_ATOMIC(__SPIRV_BUILTIN, unsigned int, j, __CLC_OP, + __HIP_BUILTIN) +AMDGPU_SAFE_ATOMIC(__SPIRV_BUILTIN, long, l, __CLC_OP, __HIP_BUILTIN) +AMDGPU_SAFE_ATOMIC(__SPIRV_BUILTIN, unsigned long, m, __CLC_OP, + __HIP_BUILTIN) diff --git a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xor.cl b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xor.cl index 869164f16f55b..5c8ebc8d90629 100644 --- a/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xor.cl +++ b/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xor.cl @@ -10,16 +10,8 @@ #include #include -#define __CLC_XOR ^ +#define __CLC_OP ^ +#define __SPIRV_BUILTIN _Z17__spirv_AtomicXor +#define __HIP_BUILTIN __hip_atomic_fetch_xor -AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, int, i, __CLC_XOR) -AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, unsigned int, j, __CLC_XOR) -AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, long, l, __CLC_XOR) -AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, unsigned long, m, __CLC_XOR) - -#undef __CLC_XOR -#undef AMDGPU_ATOMIC -#undef AMDGPU_ATOMIC_IMPL -#undef AMDGPU_ARCH_GEQ -#undef AMDGPU_ARCH_BETWEEN -#undef GET_ATOMIC_SCOPE_AND_ORDER +#include "atomic_safe.def" diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index 15de429ca6095..f7962737bcae1 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -130,6 +130,11 @@ struct AMDGPULowerKernelAttributesPass PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM); }; +struct AMDGPUOclcReflectPass : public PassInfoMixin { +public: + PreservedAnalyses run(Function &M, FunctionAnalysisManager &AM); +}; + void initializeAMDGPULowerModuleLDSLegacyPass(PassRegistry &); extern char &AMDGPULowerModuleLDSLegacyPassID; diff --git a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp new file mode 100644 index 0000000000000..b0675b0a0919d --- /dev/null +++ b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp @@ -0,0 +1,91 @@ +//===- AMDGPUReflect.cpp --------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This pass searches for occurences of the AMDGPU_OCLC_REFLECT function, and +// replaces the calls with some val dependent on the operand of the func. This +// can be used to reflect across different implementations of functions at +// compile time based on a compiler flag or some other means. The first use case +// is to choose a safe or unsafe version of atomic_xor at compile time, which +// can be chosen at compile time by setting the flag +// --amdgpu-oclc-unsafe-int-atomics=true. +// +// This pass is similar to the NVPTX pass NVVMReflect. +// +//===----------------------------------------------------------------------===// + +#include "AMDGPU.h" +#include "llvm/IR/Constants.h" +#include "llvm/IR/Dominators.h" +#include "llvm/IR/InstIterator.h" +#include "llvm/IR/Instructions.h" +#include "llvm/Pass.h" +#include "llvm/Support/CommandLine.h" + +using namespace llvm; + +#define AMDGPU_OCLC_REFLECT "__oclc_amdgpu_reflect" + +static cl::opt + AMDGPUReflectEnabled("amdgpu-oclc-reflect-enable", cl::init(true), + cl::Hidden, + cl::desc("AMDGPU reflection, enabled by default")); +static cl::opt AMDGPUUnsafeIntAtomicsEnable( + "amdgpu-oclc-unsafe-int-atomics", cl::init(false), cl::Hidden, + cl::desc("Should unsafe int atomics be chosen. Disabled by default.")); + +PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, + FunctionAnalysisManager &AM) { + if (!AMDGPUReflectEnabled) + return PreservedAnalyses::all(); + + if (F.getName() == AMDGPU_OCLC_REFLECT) { + assert(F.isDeclaration() && + "__oclc_amdgpu_reflect function should not have a body"); + return PreservedAnalyses::all(); + } + + SmallVector ToRemove; + + for (Instruction &I : instructions(F)) { + CallInst *Call = dyn_cast(&I); + if (!Call) + continue; + if (Function *Callee = Call->getCalledFunction(); + !Callee || Callee->getName() != AMDGPU_OCLC_REFLECT) + continue; + + assert(Call->getNumOperands() == 2 && + "Wrong number of operands to __oclc_amdgpu_reflect function"); + + ToRemove.push_back(Call); + } + + for (Instruction *I : ToRemove) { + CallInst *Call = dyn_cast(I); + const Value *Str = Call->getArgOperand(0); + const Value *Operand = cast(Str)->getOperand(0); + StringRef ReflectArg = cast(Operand)->getAsString(); + ReflectArg = ReflectArg.substr(0, ReflectArg.size() - 1); + + if (ReflectArg == "AMDGPU_OCLC_UNSAFE_INT_ATOMICS") { + int ReflectVal = AMDGPUUnsafeIntAtomicsEnable ? 1 : 0; + Call->replaceAllUsesWith(ConstantInt::get(Call->getType(), ReflectVal)); + } else { + assert(false && "Invalid arg passed to __oclc_amdgpu_reflect"); + } + I->eraseFromParent(); + } + + if (!ToRemove.size()) + return PreservedAnalyses::all(); + + PreservedAnalyses PA; + PA.preserveSet(); + PA.preserve(); + return PA; +} diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp index 7b5da69069cde..90ac45634c824 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp @@ -663,6 +663,7 @@ void AMDGPUTargetMachine::registerPassBuilderCallbacks( PB.registerPipelineStartEPCallback( [](ModulePassManager &PM, OptimizationLevel Level) { FunctionPassManager FPM; + FPM.addPass(AMDGPUOclcReflectPass()); FPM.addPass(AMDGPUUseNativeCallsPass()); if (EnableLibCallSimplify && Level != OptimizationLevel::O0) FPM.addPass(AMDGPUSimplifyLibCallsPass()); diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt index b6f01b78d32d4..a0714f97594a1 100644 --- a/llvm/lib/Target/AMDGPU/CMakeLists.txt +++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt @@ -89,6 +89,7 @@ add_llvm_target(AMDGPUCodeGen AMDGPUPrintfRuntimeBinding.cpp AMDGPUPromoteAlloca.cpp AMDGPUPromoteKernelArguments.cpp + AMDGPUOclcReflect.cpp AMDGPURegBankCombiner.cpp AMDGPURegBankSelect.cpp AMDGPURegisterBankInfo.cpp diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll new file mode 100644 index 0000000000000..b95cb7d64de7c --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll @@ -0,0 +1,17 @@ +; RUN: opt -S -p amdgpu-oclc-reflect %s | FileCheck %s -check-prefixes=CHECK,CHECK-SAFE-ATOMICS +; RUN: opt -S -p amdgpu-oclc-reflect -amdgpu-oclc-unsafe-int-atomics=true %s | FileCheck %s -check-prefixes=CHECK,CHECK-UNSAFE-ATOMICS + +target triple = "amdgcn-amd-amdhsa" + +@.str = private unnamed_addr addrspace(4) constant [31 x i8] c"AMDGPU_OCLC_UNSAFE_INT_ATOMICS\00", align 1 + +declare hidden i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef) local_unnamed_addr + +define i32 @foo() { +; CHECK-NOT: call i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef @.str) +; CHECK-SAFE-ATOMICS: ret i32 0 +; CHECK-UNSAFE-ATOMICS: ret i32 1 + %call = tail call i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef @.str) + ret i32 %call +} + From 917704e1940379294d908f1a5acf7e1ab9fb854f Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 1 Apr 2024 13:28:44 +0100 Subject: [PATCH 02/13] Add SYCL test and fix filename in header Add a test to make sure that the correct safe/unsafe int atomics are being generated from SYCL tomic ref member functions. --- llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp | 2 +- .../hip/atomic/amdgpu_unsafe_atomics.cpp | 24 +++++++++++++++++++ 2 files changed, 25 insertions(+), 1 deletion(-) create mode 100644 sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp diff --git a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp index b0675b0a0919d..9567a2e84ed48 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp @@ -1,4 +1,4 @@ -//===- AMDGPUReflect.cpp --------------------------------------------------===// +//===- AMDGPUOclcReflect.cpp ----------------------------------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. diff --git a/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp b/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp new file mode 100644 index 0000000000000..0412412fa727d --- /dev/null +++ b/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp @@ -0,0 +1,24 @@ +// REQUIRES: hip_be +// RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-SAFE +// RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -mllvm --amdgpu-oclc-unsafe-int-atomics=true -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-UNSAFE + +#include + +int main() { + sycl::queue{}.single_task([=] { + int a; + sycl::atomic_ref + atomicInt(a); + atomicInt.fetch_xor(1); + atomicInt.fetch_and(1); + atomicInt.fetch_or(1); + // CHECK: __CLANG_OFFLOAD_BUNDLE____START__ sycl-amdgcn-amd-amdhsa- + // CHECK-SAFE: cmpxchg volatile + // CHECK-SAFE-NOT: atomicrmw + // CHECK-UNSAFE: atomicrmw volatile xor + // CHECK-UNSAFE: atomicrmw volatile and + // CHECK-UNSAFE: atomicrmw volatile or + // CHECK-UNSAFE-NOT: cmpxchg + // CHECK: __CLANG_OFFLOAD_BUNDLE____END__ sycl-amdgcn-amd-amdhsa- + }); +} From 7e27d852da0e675f87b710360d5c340199b85f62 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 16:02:03 +0100 Subject: [PATCH 03/13] List passname alphabetically --- llvm/lib/Target/AMDGPU/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt index a0714f97594a1..4f75da401bf7f 100644 --- a/llvm/lib/Target/AMDGPU/CMakeLists.txt +++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt @@ -82,6 +82,7 @@ add_llvm_target(AMDGPUCodeGen AMDGPUInsertSingleUseVDST.cpp AMDGPUMarkLastScratchLoad.cpp AMDGPUMIRFormatter.cpp + AMDGPUOclcReflect.cpp AMDGPUOpenCLEnqueuedBlockLowering.cpp AMDGPUPerfHintAnalysis.cpp AMDGPUPostLegalizerCombiner.cpp @@ -89,7 +90,6 @@ add_llvm_target(AMDGPUCodeGen AMDGPUPrintfRuntimeBinding.cpp AMDGPUPromoteAlloca.cpp AMDGPUPromoteKernelArguments.cpp - AMDGPUOclcReflect.cpp AMDGPURegBankCombiner.cpp AMDGPURegBankSelect.cpp AMDGPURegisterBankInfo.cpp From 7310737f748ad94b70f867b6f96c0bf5b1fba2c8 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 16:17:58 +0100 Subject: [PATCH 04/13] Remove blankline --- llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll | 1 - 1 file changed, 1 deletion(-) diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll index b95cb7d64de7c..3090a95792ffb 100644 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll @@ -14,4 +14,3 @@ define i32 @foo() { %call = tail call i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef @.str) ret i32 %call } - From 819407953ba9497e4d131defc4b03bf21ad5e48e Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 16:24:50 +0100 Subject: [PATCH 05/13] Register Pass With New Pass Registry .inc --- llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 1 + 1 file changed, 1 insertion(+) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def index 90f36fadf3590..a6c8fe4087048 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def +++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def @@ -36,6 +36,7 @@ FUNCTION_PASS("amdgpu-lower-kernel-arguments", AMDGPULowerKernelArgumentsPass(*this)) FUNCTION_PASS("amdgpu-lower-kernel-attributes", AMDGPULowerKernelAttributesPass()) +FUNCTION_PASS("amdgpu-oclc-reflect", AMDGPUOclcReflect()) FUNCTION_PASS("amdgpu-simplifylib", AMDGPUSimplifyLibCallsPass()) FUNCTION_PASS("amdgpu-promote-alloca", AMDGPUPromoteAllocaPass(*this)) FUNCTION_PASS("amdgpu-promote-alloca-to-vector", From 4d17f4f5f2b6e4cbc68df03b67122b3ca0e62edb Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 16:26:04 +0100 Subject: [PATCH 06/13] Typo --- libclc/CMakeLists.txt | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 468b03d7403cf..5d8d8e98d0ec5 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -347,9 +347,8 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) # which will also work for the generic address space. set( supports_generic_addrspace FALSE ) elseif( ARCH STREQUAL "amdgcn" ) - set( build_flags ) set( opt_flags -O3 --amdgpu-oclc-reflect-enable=false ) - elseif( ARCH STREQUAL "x86_64") + elseif( ARCH STREQUAL x86_64) set( opt_flags ) set( supports_generic_addrspace FALSE ) else() From 7f2771ba7c408b5e25242c2b6c6a8f740c378d09 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 16:31:03 +0100 Subject: [PATCH 07/13] Require the use of AMDGPU reflect --- llvm/lib/Target/AMDGPU/AMDGPU.h | 1 + llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index f7962737bcae1..e2cea6c30f139 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -133,6 +133,7 @@ struct AMDGPULowerKernelAttributesPass struct AMDGPUOclcReflectPass : public PassInfoMixin { public: PreservedAnalyses run(Function &M, FunctionAnalysisManager &AM); + static bool isRequired() { return true; } }; void initializeAMDGPULowerModuleLDSLegacyPass(PassRegistry &); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def index a6c8fe4087048..6f1e296fb09c0 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def +++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def @@ -36,7 +36,7 @@ FUNCTION_PASS("amdgpu-lower-kernel-arguments", AMDGPULowerKernelArgumentsPass(*this)) FUNCTION_PASS("amdgpu-lower-kernel-attributes", AMDGPULowerKernelAttributesPass()) -FUNCTION_PASS("amdgpu-oclc-reflect", AMDGPUOclcReflect()) +FUNCTION_PASS("amdgpu-oclc-reflect", AMDGPUOclcReflectPass()) FUNCTION_PASS("amdgpu-simplifylib", AMDGPUSimplifyLibCallsPass()) FUNCTION_PASS("amdgpu-promote-alloca", AMDGPUPromoteAllocaPass(*this)) FUNCTION_PASS("amdgpu-promote-alloca-to-vector", From 695fcfa7af7b68fc851a6e7952a2d2b8c10721ce Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 16:52:57 +0100 Subject: [PATCH 08/13] Restructure comment --- llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp index 9567a2e84ed48..e3982d5015746 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp @@ -9,10 +9,12 @@ // This pass searches for occurences of the AMDGPU_OCLC_REFLECT function, and // replaces the calls with some val dependent on the operand of the func. This // can be used to reflect across different implementations of functions at -// compile time based on a compiler flag or some other means. The first use case -// is to choose a safe or unsafe version of atomic_xor at compile time, which -// can be chosen at compile time by setting the flag -// --amdgpu-oclc-unsafe-int-atomics=true. +// compile time based on a compiler flag or some other means. This pass +// currently supports use cases: +// +// 1. Choose a safe or unsafe version of atomic_xor at compile time, which can +// be chosen at compile time by setting the flag +// --amdgpu-oclc-unsafe-int-atomics=true. // // This pass is similar to the NVPTX pass NVVMReflect. // From 6d247722dde345e89963e5aebecb45898443c360 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Mon, 22 Apr 2024 17:23:05 +0100 Subject: [PATCH 09/13] Respond to comments - Change getNumOperands to arg size - Move size check above the for loop --- llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp index e3982d5015746..d5f3d65dd61e0 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp @@ -61,12 +61,15 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, !Callee || Callee->getName() != AMDGPU_OCLC_REFLECT) continue; - assert(Call->getNumOperands() == 2 && + assert(Call->arg_size() == 1 && "Wrong number of operands to __oclc_amdgpu_reflect function"); ToRemove.push_back(Call); } + if (!ToRemove.size()) + return PreservedAnalyses::all(); + for (Instruction *I : ToRemove) { CallInst *Call = dyn_cast(I); const Value *Str = Call->getArgOperand(0); @@ -83,9 +86,6 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, I->eraseFromParent(); } - if (!ToRemove.size()) - return PreservedAnalyses::all(); - PreservedAnalyses PA; PA.preserveSet(); PA.preserve(); From 8eedac83effbb32c4439000eeace0152a603ea5c Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 23 Apr 2024 15:54:35 +0100 Subject: [PATCH 10/13] Respond to comments - Use a vector of CallInsts instead of Instructions. - Change assert(fasle) to report_fatal_error. --- llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp index d5f3d65dd61e0..8b62c632ef316 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp @@ -51,7 +51,7 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, return PreservedAnalyses::all(); } - SmallVector ToRemove; + SmallVector ToRemove; for (Instruction &I : instructions(F)) { CallInst *Call = dyn_cast(&I); @@ -70,8 +70,7 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, if (!ToRemove.size()) return PreservedAnalyses::all(); - for (Instruction *I : ToRemove) { - CallInst *Call = dyn_cast(I); + for (CallInst *Call : ToRemove) { const Value *Str = Call->getArgOperand(0); const Value *Operand = cast(Str)->getOperand(0); StringRef ReflectArg = cast(Operand)->getAsString(); @@ -81,9 +80,9 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, int ReflectVal = AMDGPUUnsafeIntAtomicsEnable ? 1 : 0; Call->replaceAllUsesWith(ConstantInt::get(Call->getType(), ReflectVal)); } else { - assert(false && "Invalid arg passed to __oclc_amdgpu_reflect"); + report_fatal_error("Invalid arg passed to __oclc_amdgpu_reflect"); } - I->eraseFromParent(); + Call->eraseFromParent(); } PreservedAnalyses PA; From c4ab1ed91e5662430380f4e0c7cb335e7cc9be71 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 23 Apr 2024 15:56:30 +0100 Subject: [PATCH 11/13] Typo --- libclc/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 5d8d8e98d0ec5..4570ac7a68000 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -315,7 +315,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) elseif( ARCH STREQUAL amdgcn ) # AMDGCN needs libclc to be compiled to high bc version since all atomic # clang builtins need to be accessible - list( APPEND flags -mcpu=gfx940 -mllvm --amdgpu-oclc-reflect-enable=false) + list( APPEND flags -mcpu=gfx940 -mllvm --amdgpu-oclc-reflect-enable=false ) elseif( ARCH STREQUAL x86_64) # TODO: This is used by SYCL Native Cpu, we should define an option to set this flags list( APPEND flags -Xclang -target-feature -Xclang +avx @@ -346,7 +346,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} ) # simultaneously, we choose declare the builtins using the private space, # which will also work for the generic address space. set( supports_generic_addrspace FALSE ) - elseif( ARCH STREQUAL "amdgcn" ) + elseif( ARCH STREQUAL amdgcn ) set( opt_flags -O3 --amdgpu-oclc-reflect-enable=false ) elseif( ARCH STREQUAL x86_64) set( opt_flags ) From bf5d2d0501c816852af1155c46be3a0227f7d19c Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Wed, 24 Apr 2024 10:20:55 +0100 Subject: [PATCH 12/13] Respond to comments - Use auto - Use drop_back to remove null byte - Replace hip_be with hip --- llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp | 4 ++-- .../check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp index 8b62c632ef316..47430cadf5afb 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUOclcReflect.cpp @@ -54,7 +54,7 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, SmallVector ToRemove; for (Instruction &I : instructions(F)) { - CallInst *Call = dyn_cast(&I); + auto *Call = dyn_cast(&I); if (!Call) continue; if (Function *Callee = Call->getCalledFunction(); @@ -74,7 +74,7 @@ PreservedAnalyses AMDGPUOclcReflectPass::run(Function &F, const Value *Str = Call->getArgOperand(0); const Value *Operand = cast(Str)->getOperand(0); StringRef ReflectArg = cast(Operand)->getAsString(); - ReflectArg = ReflectArg.substr(0, ReflectArg.size() - 1); + ReflectArg = ReflectArg.drop_back(1); if (ReflectArg == "AMDGPU_OCLC_UNSAFE_INT_ATOMICS") { int ReflectVal = AMDGPUUnsafeIntAtomicsEnable ? 1 : 0; diff --git a/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp b/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp index 0412412fa727d..18a76a6d4fec0 100644 --- a/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp +++ b/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp @@ -1,4 +1,4 @@ -// REQUIRES: hip_be +// REQUIRES: hip // RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-SAFE // RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -mllvm --amdgpu-oclc-unsafe-int-atomics=true -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-UNSAFE From bbce2f9f1c0efa415225a6800fdb15dbdb543af1 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Wed, 24 Apr 2024 10:33:03 +0100 Subject: [PATCH 13/13] Use update_test_checks.py Change opt test to use update_test_checks.py --- llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll index 3090a95792ffb..27557aa8a5c00 100644 --- a/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll +++ b/llvm/test/CodeGen/AMDGPU/amdgpu-oclc-reflect.ll @@ -1,3 +1,4 @@ +; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 4 ; RUN: opt -S -p amdgpu-oclc-reflect %s | FileCheck %s -check-prefixes=CHECK,CHECK-SAFE-ATOMICS ; RUN: opt -S -p amdgpu-oclc-reflect -amdgpu-oclc-unsafe-int-atomics=true %s | FileCheck %s -check-prefixes=CHECK,CHECK-UNSAFE-ATOMICS @@ -8,9 +9,14 @@ target triple = "amdgcn-amd-amdhsa" declare hidden i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef) local_unnamed_addr define i32 @foo() { -; CHECK-NOT: call i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef @.str) -; CHECK-SAFE-ATOMICS: ret i32 0 -; CHECK-UNSAFE-ATOMICS: ret i32 1 +; CHECK-SAFE-ATOMICS-LABEL: define i32 @foo() { +; CHECK-SAFE-ATOMICS-NEXT: ret i32 0 +; +; CHECK-UNSAFE-ATOMICS-LABEL: define i32 @foo() { +; CHECK-UNSAFE-ATOMICS-NEXT: ret i32 1 +; %call = tail call i32 @__oclc_amdgpu_reflect(ptr addrspace(4) noundef @.str) ret i32 %call } +;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: +; CHECK: {{.*}}