[clang] [Clang] Fix deleted C++20 destructor call in GPU header (PR #214837)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-x86

Author: Joseph Huber (jhuber6)

<details>
<summary>Changes</summary>

Summary:
The -llvm header used this incorrectly so it wouldn't work with newer
C++.


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


2 Files Affected:

- (modified) clang/lib/Headers/__clang_gpu_builtin_vars.h (+10-5) 
- (added) clang/test/Headers/gpu-builtin-vars.cpp (+54) 


``````````diff
diff --git a/clang/lib/Headers/__clang_gpu_builtin_vars.h b/clang/lib/Headers/__clang_gpu_builtin_vars.h
index b80248dcd2be3..21bc80d7c29b9 100644
--- a/clang/lib/Headers/__clang_gpu_builtin_vars.h
+++ b/clang/lib/Headers/__clang_gpu_builtin_vars.h
@@ -61,11 +61,16 @@ __GPU_COORD_BUILTIN(__gpu_builtin_gridDim_t, __gpu_num_blocks_x(),
 #pragma pop_macro("__GPU_COORD_BUILTIN")
 #pragma pop_macro("__GPU_DISALLOW_BUILTINVAR_ACCESS")
 
-static inline const
-    __attribute__((device)) __gpu_builtin_threadIdx_t threadIdx{};
-static inline const __attribute__((device)) __gpu_builtin_blockIdx_t blockIdx{};
-static inline const __attribute__((device)) __gpu_builtin_blockDim_t blockDim{};
-static inline const __attribute__((device)) __gpu_builtin_gridDim_t gridDim{};
+#pragma push_macro("__GPU_BUILTIN_VAR")
+#define __GPU_BUILTIN_VAR                                                      \
+  extern const __attribute__((device)) __attribute__((weak))
+
+__GPU_BUILTIN_VAR __gpu_builtin_threadIdx_t threadIdx;
+__GPU_BUILTIN_VAR __gpu_builtin_blockIdx_t blockIdx;
+__GPU_BUILTIN_VAR __gpu_builtin_blockDim_t blockDim;
+__GPU_BUILTIN_VAR __gpu_builtin_gridDim_t gridDim;
+
+#pragma pop_macro("__GPU_BUILTIN_VAR")
 
 #endif // device compile
 #endif // __CLANG_GPU_BUILTIN_VARS_H__
diff --git a/clang/test/Headers/gpu-builtin-vars.cpp b/clang/test/Headers/gpu-builtin-vars.cpp
new file mode 100644
index 0000000000000..5ed68a2f5957d
--- /dev/null
+++ b/clang/test/Headers/gpu-builtin-vars.cpp
@@ -0,0 +1,54 @@
+// HIP on AMDGPU.
+// RUN: %clang_cc1 -std=c++17 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple amdgpu9.0a-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -x hip -fcuda-is-device -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+// RUN: %clang_cc1 -std=c++20 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple amdgpu9.0a-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -x hip -fcuda-is-device -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+
+// HIP on SPIR-V.
+// RUN: %clang_cc1 -std=c++20 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple spirv64-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -x hip -fcuda-is-device -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+
+// CUDA on NVPTX.
+// RUN: %clang_cc1 -std=c++17 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple nvptx64-nvidia-cuda -aux-triple x86_64-unknown-unknown \
+// RUN:   -x cuda -fcuda-is-device -target-cpu sm_70 -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+// RUN: %clang_cc1 -std=c++20 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple nvptx64-nvidia-cuda -aux-triple x86_64-unknown-unknown \
+// RUN:   -x cuda -fcuda-is-device -target-cpu sm_70 -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+
+// HIP host compilation.
+// RUN: %clang_cc1 -std=c++20 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple x86_64-unknown-unknown -aux-triple amdgpu9.0a-amd-amdhsa \
+// RUN:   -x hip -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+
+// CUDA host compilation.
+// RUN: %clang_cc1 -std=c++20 -internal-isystem %S/Inputs/include \
+// RUN:   -internal-isystem %S/../../lib/Headers \
+// RUN:   -triple x86_64-unknown-unknown -aux-triple nvptx64-nvidia-cuda \
+// RUN:   -aux-target-cpu sm_70 -x cuda -fsyntax-only -verify %s \
+// RUN:   -include __clang_gpu_builtin_vars.h
+
+// expected-no-diagnostics
+
+__attribute__((global)) void test_kernel(unsigned *out) {
+  unsigned i = threadIdx.x + threadIdx.y + threadIdx.z;
+  i += blockIdx.x + blockIdx.y + blockIdx.z;
+  i += blockDim.x + blockDim.y + blockDim.z;
+  i += gridDim.x + gridDim.y + gridDim.z;
+  *out = i;
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/214837