[libc-commits] [clang] [compiler-rt] [libc] [libcxx] [llvm] [AMDGPU] Fix code object verion not being set to 'none' (PR #135036)
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Wed Apr 9 09:31:41 PDT 2025
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/135036
>From e41985970c254f3eda71cb5ef3a1dc321c8e6f56 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Wed, 9 Apr 2025 09:41:38 -0500
Subject: [PATCH 1/2] [AMDGPU] Fix code object verion not being set to 'none'
Summary:
Previously, we removed the special handling for the code object version
global. I erroneously thought that this meant we cold get rid of this
weird `-Xclang` option. However, this also emits an LLVM IR module flag,
which will then cause linking issues.
---
compiler-rt/cmake/builtin-config-ix.cmake | 1 +
compiler-rt/lib/builtins/CMakeLists.txt | 6 ++++++
.../modules/LLVMLibCCompileOptionRules.cmake | 2 ++
libcxx/cmake/caches/AMDGPU.cmake | 6 ++++--
offload/DeviceRTL/CMakeLists.txt | 2 +-
offload/test/api/amdgpu_code_object.c | 16 ++++++++++++++++
6 files changed, 30 insertions(+), 3 deletions(-)
create mode 100644 offload/test/api/amdgpu_code_object.c
diff --git a/compiler-rt/cmake/builtin-config-ix.cmake b/compiler-rt/cmake/builtin-config-ix.cmake
index e1945ba2b2230..7bd3269bd999d 100644
--- a/compiler-rt/cmake/builtin-config-ix.cmake
+++ b/compiler-rt/cmake/builtin-config-ix.cmake
@@ -22,6 +22,7 @@ builtin_check_c_compiler_flag(-Wno-pedantic COMPILER_RT_HAS_WNO_PEDANTIC
builtin_check_c_compiler_flag(-nogpulib COMPILER_RT_HAS_NOGPULIB_FLAG)
builtin_check_c_compiler_flag(-flto COMPILER_RT_HAS_FLTO_FLAG)
builtin_check_c_compiler_flag(-fconvergent-functions COMPILER_RT_HAS_FCONVERGENT_FUNCTIONS_FLAG)
+builtin_check_c_compiler_flag("-Xclang -mcode-object-version=none" COMPILER_RT_HAS_CODE_OBJECT_VERSION_FLAG)
builtin_check_c_compiler_flag(-Wbuiltin-declaration-mismatch COMPILER_RT_HAS_WBUILTIN_DECLARATION_MISMATCH_FLAG)
builtin_check_c_compiler_flag(/Zl COMPILER_RT_HAS_ZL_FLAG)
builtin_check_c_compiler_flag(-fcf-protection=full COMPILER_RT_HAS_FCF_PROTECTION_FLAG)
diff --git a/compiler-rt/lib/builtins/CMakeLists.txt b/compiler-rt/lib/builtins/CMakeLists.txt
index 5d78b5a780428..3cdbf21ed403d 100644
--- a/compiler-rt/lib/builtins/CMakeLists.txt
+++ b/compiler-rt/lib/builtins/CMakeLists.txt
@@ -833,6 +833,12 @@ else ()
append_list_if(COMPILER_RT_HAS_FLTO_FLAG -flto BUILTIN_CFLAGS)
append_list_if(COMPILER_RT_HAS_FCONVERGENT_FUNCTIONS_FLAG
-fconvergent-functions BUILTIN_CFLAGS)
+
+ # AMDGPU targets want to use a generic ABI.
+ if("${COMPILER_RT_DEFAULT_TARGET_ARCH}" MATCHES "amdgcn")
+ append_list_if(COMPILER_RT_HAS_CODE_OBJECT_VERSION_FLAG
+ "SHELL:-Xclang -mcode-object-version=none" BUILTIN_CFLAGS)
+ endif()
endif()
set(BUILTIN_DEFS "")
diff --git a/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake b/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake
index ddd18ef293c8d..0facb0b9be0c1 100644
--- a/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake
+++ b/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake
@@ -215,6 +215,8 @@ function(_get_common_compile_options output_var flags)
if(LIBC_CUDA_ROOT)
list(APPEND compile_options "--cuda-path=${LIBC_CUDA_ROOT}")
endif()
+ elseif(LIBC_TARGET_ARCHITECTURE_IS_AMDGPU)
+ list(APPEND compile_options "SHELL:-Xclang -mcode-object-version=none")
endif()
endif()
set(${output_var} ${compile_options} PARENT_SCOPE)
diff --git a/libcxx/cmake/caches/AMDGPU.cmake b/libcxx/cmake/caches/AMDGPU.cmake
index d4aa28b4134ea..e7bf3f53891f0 100644
--- a/libcxx/cmake/caches/AMDGPU.cmake
+++ b/libcxx/cmake/caches/AMDGPU.cmake
@@ -32,6 +32,8 @@ set(LIBCXX_TEST_CONFIG "amdgpu-libc++-shared.cfg.in" CACHE STRING "")
set(LIBCXX_TEST_PARAMS "optimization=none;long_tests=False;executor=amdhsa-loader" CACHE STRING "")
# Necessary compile flags for AMDGPU.
-set(LIBCXX_ADDITIONAL_COMPILE_FLAGS "-nogpulib;-flto;-fconvergent-functions" CACHE STRING "")
-set(LIBCXXABI_ADDITIONAL_COMPILE_FLAGS "-nogpulib;-flto;-fconvergent-functions" CACHE STRING "")
+set(LIBCXX_ADDITIONAL_COMPILE_FLAGS
+ "-nogpulib;-flto;-fconvergent-functions;SHELL:-Xclang -mcode-object-version=none" CACHE STRING "")
+set(LIBCXXABI_ADDITIONAL_COMPILE_FLAGS
+ "-nogpulib;-flto;-fconvergent-functions;SHELL:-Xclang -mcode-object-version=none" CACHE STRING "")
set(CMAKE_REQUIRED_FLAGS "-nogpulib" CACHE STRING "")
diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt
index 07888217b6c68..8f2a1fd01fabc 100644
--- a/offload/DeviceRTL/CMakeLists.txt
+++ b/offload/DeviceRTL/CMakeLists.txt
@@ -255,7 +255,7 @@ function(compileDeviceRTLLibrary target_name target_triple)
endfunction()
add_custom_target(omptarget.devicertl.amdgpu)
-compileDeviceRTLLibrary(amdgpu amdgcn-amd-amdhsa)
+compileDeviceRTLLibrary(amdgpu amdgcn-amd-amdhsa -Xclang -mcode-object-version=none)
add_custom_target(omptarget.devicertl.nvptx)
compileDeviceRTLLibrary(nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx63)
diff --git a/offload/test/api/amdgpu_code_object.c b/offload/test/api/amdgpu_code_object.c
new file mode 100644
index 0000000000000..95d14f6772e77
--- /dev/null
+++ b/offload/test/api/amdgpu_code_object.c
@@ -0,0 +1,16 @@
+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -Xclang \
+// RUN: -mcode-object-version=5
+// RUN: %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+#include <stdio.h>
+
+// Test to make sure we can build and run with the previous COV.
+int main() {
+#pragma omp target
+ ;
+
+ // CHECK: PASS
+ printf("PASS\n");
+}
>From b87322363d7f8b3fbf59eeaef46ac96e7b46c481 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Wed, 9 Apr 2025 11:31:30 -0500
Subject: [PATCH 2/2] Stop using undefined global now
---
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 57 ++++-----------------
clang/test/CodeGen/amdgpu-abi-version.c | 16 ++----
2 files changed, 14 insertions(+), 59 deletions(-)
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index b56b739094ff3..51b8db3c3a7d8 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -54,63 +54,26 @@ Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) {
/// Emit code based on Code Object ABI version.
/// COV_4 : Emit code to use dispatch ptr
/// COV_5+ : Emit code to use implicitarg ptr
-/// COV_NONE : Emit code to load a global variable "__oclc_ABI_version"
-/// and use its value for COV_4 or COV_5+ approach. It is used for
-/// compiling device libraries in an ABI-agnostic way.
Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
llvm::LoadInst *LD;
auto Cov = CGF.getTarget().getTargetOpts().CodeObjectVersion;
-
- if (Cov == CodeObjectVersionKind::COV_None) {
- StringRef Name = "__oclc_ABI_version";
- auto *ABIVersionC = CGF.CGM.getModule().getNamedGlobal(Name);
- if (!ABIVersionC)
- ABIVersionC = new llvm::GlobalVariable(
- CGF.CGM.getModule(), CGF.Int32Ty, false,
- llvm::GlobalValue::ExternalLinkage, nullptr, Name, nullptr,
- llvm::GlobalVariable::NotThreadLocal,
- CGF.CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant));
-
- // This load will be eliminated by the IPSCCP because it is constant
- // weak_odr without externally_initialized. Either changing it to weak or
- // adding externally_initialized will keep the load.
- Value *ABIVersion = CGF.Builder.CreateAlignedLoad(CGF.Int32Ty, ABIVersionC,
- CGF.CGM.getIntAlign());
-
- Value *IsCOV5 = CGF.Builder.CreateICmpSGE(
- ABIVersion,
- llvm::ConstantInt::get(CGF.Int32Ty, CodeObjectVersionKind::COV_5));
-
+ Value *GEP = nullptr;
+ if (Cov >= CodeObjectVersionKind::COV_5) {
// Indexing the implicit kernarg segment.
- Value *ImplicitGEP = CGF.Builder.CreateConstGEP1_32(
+ GEP = CGF.Builder.CreateConstGEP1_32(
CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);
-
- // Indexing the HSA kernel_dispatch_packet struct.
- Value *DispatchGEP = CGF.Builder.CreateConstGEP1_32(
- CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2);
-
- auto Result = CGF.Builder.CreateSelect(IsCOV5, ImplicitGEP, DispatchGEP);
- LD = CGF.Builder.CreateLoad(
- Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2)));
} else {
- Value *GEP = nullptr;
- if (Cov >= CodeObjectVersionKind::COV_5) {
- // Indexing the implicit kernarg segment.
- GEP = CGF.Builder.CreateConstGEP1_32(
- CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);
- } else {
- // Indexing the HSA kernel_dispatch_packet struct.
- GEP = CGF.Builder.CreateConstGEP1_32(
- CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2);
- }
- LD = CGF.Builder.CreateLoad(
- Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2)));
+ // Indexing the HSA kernel_dispatch_packet struct.
+ GEP = CGF.Builder.CreateConstGEP1_32(CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF),
+ 4 + Index * 2);
}
+ LD = CGF.Builder.CreateLoad(
+ Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2)));
llvm::MDBuilder MDHelper(CGF.getLLVMContext());
- llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1),
- APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));
+ llvm::MDNode *RNode = MDHelper.createRange(
+ APInt(16, 1), APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));
LD->setMetadata(llvm::LLVMContext::MD_range, RNode);
LD->setMetadata(llvm::LLVMContext::MD_noundef,
llvm::MDNode::get(CGF.getLLVMContext(), {}));
diff --git a/clang/test/CodeGen/amdgpu-abi-version.c b/clang/test/CodeGen/amdgpu-abi-version.c
index b9c1de0521b95..2bfb15c787da9 100644
--- a/clang/test/CodeGen/amdgpu-abi-version.c
+++ b/clang/test/CodeGen/amdgpu-abi-version.c
@@ -1,23 +1,15 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 5
// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm -mcode-object-version=none %s -o - | FileCheck %s
-//.
-// CHECK: @__oclc_ABI_version = external addrspace(4) global i32
-//.
// CHECK-LABEL: define dso_local i32 @foo(
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
-// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 4
-// CHECK-NEXT: [[TMP1:%.*]] = icmp sge i32 [[TMP0]], 500
-// CHECK-NEXT: [[TMP2:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// CHECK-NEXT: [[TMP3:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP2]], i32 12
-// CHECK-NEXT: [[TMP4:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
-// CHECK-NEXT: [[TMP5:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP4]], i32 4
-// CHECK-NEXT: [[TMP6:%.*]] = select i1 [[TMP1]], ptr addrspace(4) [[TMP3]], ptr addrspace(4) [[TMP5]]
-// CHECK-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG2:![0-9]+]], !invariant.load [[META3:![0-9]+]], !noundef [[META3]]
-// CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP7]] to i32
+// CHECK-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// CHECK-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 4
+// CHECK-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG2:![0-9]+]], !invariant.load [[META3:![0-9]+]], !noundef [[META3]]
+// CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP2]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
//
int foo() { return __builtin_amdgcn_workgroup_size_x(); }
More information about the libc-commits
mailing list