[AMDGPU] Fix code object version not being set to 'none' #135036

jhuber6 · 2025-04-09T15:00:24Z

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.

llvmbot · 2025-04-09T15:01:05Z

@llvm/pr-subscribers-clang
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-offload
@llvm/pr-subscribers-libcxx

@llvm/pr-subscribers-libc

Author: Joseph Huber (jhuber6)

Changes

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.

Full diff: https://github.com/llvm/llvm-project/pull/135036.diff

6 Files Affected:

(modified) compiler-rt/cmake/builtin-config-ix.cmake (+1)
(modified) compiler-rt/lib/builtins/CMakeLists.txt (+6)
(modified) libc/cmake/modules/LLVMLibCCompileOptionRules.cmake (+2)
(modified) libcxx/cmake/caches/AMDGPU.cmake (+4-2)
(modified) offload/DeviceRTL/CMakeLists.txt (+1-1)
(added) offload/test/api/amdgpu_code_object.c (+16)

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");
+}

llvmbot · 2025-04-09T15:01:06Z

@llvm/pr-subscribers-backend-amdgpu

Author: Joseph Huber (jhuber6)

Changes

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.

Full diff: https://github.com/llvm/llvm-project/pull/135036.diff

6 Files Affected:

(modified) compiler-rt/cmake/builtin-config-ix.cmake (+1)
(modified) compiler-rt/lib/builtins/CMakeLists.txt (+6)
(modified) libc/cmake/modules/LLVMLibCCompileOptionRules.cmake (+2)
(modified) libcxx/cmake/caches/AMDGPU.cmake (+4-2)
(modified) offload/DeviceRTL/CMakeLists.txt (+1-1)
(added) offload/test/api/amdgpu_code_object.c (+16)

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");
+}

ronlieb

works for me on downstream

arsenm · 2025-04-09T16:59:45Z

clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

@@ -62,7 +62,7 @@ Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {

  auto Cov = CGF.getTarget().getTargetOpts().CodeObjectVersion;

-  if (Cov == CodeObjectVersionKind::COV_None) {
+  if (Cov == CodeObjectVersionKind::COV_None && !CGF.getLangOpts().OpenMP) {


Shouldn't be language dependent? Does this mean the IR module flag?

Yeah, I'm realizing that now. This just needs to turn off the module flag but I don't want to have this lowering at all. I think it's still used in the ROCm Device Libs? So probably can't just delete it like I did before. Can I just add a flag to suppress the module metadata instead?

It's added in

llvm-project/clang/lib/CodeGen/CodeGenModule.cpp

Line 924 in 1aed6ad

if (getTarget().getTargetOpts().CodeObjectVersion !=

. Could I maybe just omit this if the user is compiling with something like -nogpulib? Otherwise, we just need another flag for it because None is pulling double duty here by doing the reduction.

What is the linking issue exactly? This is here to prevent you from linking incompatible builds

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.

…m#135036)" This reverts commit 2f41fa3.

…ne' (llvm#135036)"" This reverts commit 1a25f16.

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.

jhuber6 requested a review from a team as a code owner April 9, 2025 15:00

jhuber6 requested review from arsenm, jplehr, ronlieb, saiislam and shiltian April 9, 2025 15:00

llvmbot added compiler-rt libc++ libc++ C++ Standard Library. Not GNU libstdc++. Not libc++abi. backend:AMDGPU compiler-rt:builtins libc offload labels Apr 9, 2025

ronlieb approved these changes Apr 9, 2025

View reviewed changes

saiislam approved these changes Apr 9, 2025

View reviewed changes

lntue approved these changes Apr 9, 2025

View reviewed changes

llvmbot added clang Clang issues not falling into any other category clang:codegen IR generation bugs: mangling, exceptions, etc. labels Apr 9, 2025

ronlieb approved these changes Apr 9, 2025

View reviewed changes

jayfoad changed the title ~~[AMDGPU] Fix code object verion not being set to 'none'~~ [AMDGPU] Fix code object version not being set to 'none' Apr 9, 2025

jhuber6 force-pushed the FixCOV branch from b873223 to dee557e Compare April 9, 2025 16:57

arsenm reviewed Apr 9, 2025

View reviewed changes

jhuber6 force-pushed the FixCOV branch from dee557e to e9d477a Compare April 10, 2025 16:12

jhuber6 merged commit 2f41fa3 into llvm:main Apr 10, 2025
16 of 34 checks passed

searlmc1 pushed a commit to ROCm/llvm-project that referenced this pull request Apr 11, 2025

Revert "[AMDGPU] Fix code object version not being set to 'none' (llv…

1a25f16

…m#135036)" This reverts commit 2f41fa3.

searlmc1 pushed a commit to ROCm/llvm-project that referenced this pull request Apr 14, 2025

Revert "Revert "[AMDGPU] Fix code object version not being set to 'no…

faecb9e

…ne' (llvm#135036)"" This reverts commit 1a25f16.

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[AMDGPU] Fix code object version not being set to 'none' #135036

[AMDGPU] Fix code object version not being set to 'none' #135036

Uh oh!

jhuber6 commented Apr 9, 2025

Uh oh!

llvmbot commented Apr 9, 2025 •

edited

Loading

Uh oh!

llvmbot commented Apr 9, 2025

Uh oh!

ronlieb left a comment

Uh oh!

arsenm Apr 9, 2025

Uh oh!

jhuber6 Apr 9, 2025

Uh oh!

jhuber6 Apr 9, 2025

Uh oh!

arsenm Apr 10, 2025

Uh oh!

Uh oh!

Uh oh!

[AMDGPU] Fix code object version not being set to 'none' #135036

[AMDGPU] Fix code object version not being set to 'none' #135036

Uh oh!

Conversation

jhuber6 commented Apr 9, 2025

Uh oh!

llvmbot commented Apr 9, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Apr 9, 2025

Uh oh!

ronlieb left a comment

Choose a reason for hiding this comment

Uh oh!

arsenm Apr 9, 2025

Choose a reason for hiding this comment

Uh oh!

jhuber6 Apr 9, 2025

Choose a reason for hiding this comment

Uh oh!

jhuber6 Apr 9, 2025

Choose a reason for hiding this comment

Uh oh!

arsenm Apr 10, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

llvmbot commented Apr 9, 2025 •

edited

Loading