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