https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/214331
>From 9950531e024449e96d7da227217d0b738449da0e Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Tue, 4 Aug 2026 22:47:33 -0500 Subject: [PATCH] [Clang] Rework GPU wrapper headers for the `-llvm` environment Summary; This environment is intended to be the hermetic LLVM interface so we can compile GPU things in a standalone environment. For this to work, we resolve the standard include directory and refine the API split. Future work will add missing functions, but this should let us compile HIP device code on godbolt without needing a full ROCm interface. --- clang/lib/Driver/ToolChains/AMDGPU.cpp | 8 ++- clang/lib/Headers/CMakeLists.txt | 21 ++++++- clang/lib/Headers/__clang_gpu_builtin_vars.h | 9 --- .../Headers/__clang_gpu_device_functions.h | 61 ++----------------- clang/lib/Headers/__clang_gpu_intrinsics.h | 42 +++++-------- .../lib/Headers/__clang_gpu_runtime_wrapper.h | 28 +++++++++ .../hip_wrappers/hip/device_functions.h | 17 ++++++ .../Headers/hip_wrappers/hip/hip_runtime.h | 23 +++++++ .../test/Driver/hip-device-libs-llvm-env.hip | 4 +- clang/test/Headers/gpu-device-functions.cpp | 35 ++++++++--- 10 files changed, 142 insertions(+), 106 deletions(-) create mode 100644 clang/lib/Headers/__clang_gpu_runtime_wrapper.h create mode 100644 clang/lib/Headers/hip_wrappers/hip/device_functions.h create mode 100644 clang/lib/Headers/hip_wrappers/hip/hip_runtime.h diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp index 7bce060de0596..d1477e84e65b6 100644 --- a/clang/lib/Driver/ToolChains/AMDGPU.cpp +++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp @@ -1275,8 +1275,12 @@ void AMDGPUToolChain::AddHIPIncludeArgs(const ArgList &DriverArgs, if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc, true) && !DriverArgs.hasArg(options::OPT_nohipwrapperinc) && - !DriverArgs.hasArg(options::OPT_nobuiltininc)) - CC1Args.append({"-include", "__clang_gpu_device_functions.h"}); + !DriverArgs.hasArg(options::OPT_nobuiltininc)) { + SmallString<128> P(getDriver().ResourceDir); + llvm::sys::path::append(P, "include", "hip_wrappers"); + CC1Args.push_back("-internal-isystem"); + CC1Args.push_back(DriverArgs.MakeArgString(P)); + } return; } diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index b1d1832a598d9..21b6eb5e38052 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -310,6 +310,12 @@ set(gpu_files __clang_gpu_builtin_vars.h __clang_gpu_device_functions.h __clang_gpu_intrinsics.h + __clang_gpu_runtime_wrapper.h + ) + +set(gpu_hip_wrapper_files + hip_wrappers/hip/hip_runtime.h + hip_wrappers/hip/device_functions.h ) set(windows_only_files @@ -471,7 +477,7 @@ endfunction(clang_generate_header) foreach( f ${files} ${cuda_wrapper_files} ${cuda_wrapper_bits_files} ${cuda_wrapper_utility_files} ${ppc_wrapper_files} ${openmp_wrapper_files} ${zos_wrapper_files} ${hlsl_files} ${llvm_libc_wrapper_files} - ${llvm_offload_wrapper_files}) + ${llvm_offload_wrapper_files} ${gpu_hip_wrapper_files}) copy_header_to_output_dir(${CMAKE_CURRENT_SOURCE_DIR} ${f}) endforeach( f ) @@ -602,7 +608,7 @@ add_header_target("systemz-resource-headers" "${systemz_files};${zos_wrapper_fil add_header_target("ve-resource-headers" "${ve_files}") add_header_target("webassembly-resource-headers" "${webassembly_files}") add_header_target("x86-resource-headers" "${x86_files}") -add_header_target("gpu-resource-headers" "${gpu_files}") +add_header_target("gpu-resource-headers" "${gpu_files};${gpu_hip_wrapper_files}") # Other header groupings add_header_target("hlsl-resource-headers" "${hlsl_files};${hlsl_generated_files}") @@ -662,6 +668,11 @@ install( DESTINATION ${header_install_dir}/llvm_offload_wrappers COMPONENT clang-resource-headers) +install( + FILES ${gpu_hip_wrapper_files} + DESTINATION ${header_install_dir}/hip_wrappers/hip + COMPONENT clang-resource-headers) + install( FILES ${zos_wrapper_files} DESTINATION ${header_install_dir}/zos_wrappers @@ -807,6 +818,12 @@ install( EXCLUDE_FROM_ALL COMPONENT gpu-resource-headers) +install( + FILES ${gpu_hip_wrapper_files} + DESTINATION ${header_install_dir}/hip_wrappers/hip + EXCLUDE_FROM_ALL + COMPONENT gpu-resource-headers) + if(NOT CLANG_ENABLE_HLSL) set(EXCLUDE_HLSL EXCLUDE_FROM_ALL) endif() diff --git a/clang/lib/Headers/__clang_gpu_builtin_vars.h b/clang/lib/Headers/__clang_gpu_builtin_vars.h index b80248dcd2be3..096bef90b26b6 100644 --- a/clang/lib/Headers/__clang_gpu_builtin_vars.h +++ b/clang/lib/Headers/__clang_gpu_builtin_vars.h @@ -13,15 +13,6 @@ #include <gpuintrin.h> -// The warpSize is a runtime value rather than a compile-time constant. -static inline __attribute__((device)) const struct { - __attribute__((device, always_inline, const)) operator int() const noexcept { - return __gpu_num_lanes(); - } -} warpSize{}; - -// Make sure nobody can create instances of the coordinate types, take their -// address, copy, or assign them. #pragma push_macro("__GPU_DISALLOW_BUILTINVAR_ACCESS") #define __GPU_DISALLOW_BUILTINVAR_ACCESS(__tag) \ __attribute__((device)) __tag() = delete; \ diff --git a/clang/lib/Headers/__clang_gpu_device_functions.h b/clang/lib/Headers/__clang_gpu_device_functions.h index a5fb1253926f3..96674a76b399e 100644 --- a/clang/lib/Headers/__clang_gpu_device_functions.h +++ b/clang/lib/Headers/__clang_gpu_device_functions.h @@ -11,30 +11,11 @@ #if defined(__HIP__) || defined(__CUDA__) -#ifndef __device__ -#define __host__ __attribute__((host)) -#define __device__ __attribute__((device)) -#define __global__ __attribute__((global)) -#define __shared__ __attribute__((shared)) -#define __constant__ __attribute__((constant)) -#define __managed__ __attribute__((managed)) -#endif - #include <gpuintrin.h> #pragma push_macro("__GPU_DEVICE__") #define __GPU_DEVICE__ static __inline__ __attribute__((device, always_inline)) -#pragma push_macro("MAYBE_UNDEF") -#define MAYBE_UNDEF __attribute__((maybe_undef)) - -// warpSize and the threadIdx/blockIdx/blockDim/gridDim coordinate variables. -#include <__clang_gpu_builtin_vars.h> - -//===----------------------------------------------------------------------===// -// Integer intrinsics. -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ unsigned int __popc(unsigned int __x) { return __builtin_popcountg(__x); } @@ -128,10 +109,6 @@ __GPU_DEVICE__ unsigned int __byte_perm(unsigned int __x, unsigned int __y, return __result; } -//===----------------------------------------------------------------------===// -// Bitfield operations. -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ unsigned int __lastbit_u32_u64(unsigned long long __x) { return (unsigned int)__builtin_ctzg(__x, -1); } @@ -168,10 +145,6 @@ __GPU_DEVICE__ unsigned long long __bitinsert_u64(unsigned long long __dst, return (__dst & ~(__mask << __o)) | ((__src & __mask) << __o); } -//===----------------------------------------------------------------------===// -// Type punning. -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ int __float_as_int(float __x) { return __builtin_bit_cast(int, __x); } @@ -202,12 +175,6 @@ __GPU_DEVICE__ double __hiloint2double(int __hi, int __lo) { (unsigned long long)(unsigned int)__lo); } -//===----------------------------------------------------------------------===// -// Numeric conversions with explicit rounding. -//===----------------------------------------------------------------------===// - -// Floating-point to integer conversions map directly onto the rounding -// builtins. #define __GPU_CVT_FP2INT(__name, __res, __arg) \ __GPU_DEVICE__ __res __name##_rd(__arg __x) { \ return (__res)__builtin_elementwise_floor(__x); \ @@ -231,11 +198,9 @@ __GPU_CVT_FP2INT(__float2ull, unsigned long long, float) #undef __GPU_CVT_FP2INT -// Round-to-nearest is a plain conversion, so the '_rn' variants are exact. -// -// TODO: Directed rounding (rd/ru/rz) for integer-to-float and the narrowing -// double-to-float conversions has no portable builtin yet (need pragma STDC -// FENV_ROUND pragma), so these are stubbed. +// TODO: Directed rounding (rd/ru/rz) on the widening conversions needs a +// FENV_ROUND pragma that is not yet wired up, so only round-to-nearest is +// exact. #define __GPU_CVT_TO_F(__name, __res, __arg) \ __GPU_DEVICE__ __res __name##_rd(__arg __x) { __builtin_trap(); } \ __GPU_DEVICE__ __res __name##_rn(__arg __x) { return (__res)__x; } \ @@ -252,15 +217,9 @@ __GPU_CVT_TO_F(__double2float, float, double) #undef __GPU_CVT_TO_F -// Integer to double conversions are always exact, so only round-to-nearest is -// defined by CUDA and HIP. __GPU_DEVICE__ double __int2double_rn(int __x) { return (double)__x; } __GPU_DEVICE__ double __uint2double_rn(unsigned int __x) { return (double)__x; } -//===----------------------------------------------------------------------===// -// Wavefront vote and lane identity. -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ unsigned int __lane_id(void) { return __gpu_lane_id(); } __GPU_DEVICE__ unsigned long long __ballot(int __pred) { @@ -324,11 +283,7 @@ __GPU_DEVICE__ int __fns(unsigned int __mask, unsigned int __base, return __fns32(__mask, __base, __offset); } -//===----------------------------------------------------------------------===// -// Synchronization and fences -//===----------------------------------------------------------------------===// - -// CUDA provides __syncthreads as an NVPTX compiler builtin directly. +// CUDA lowers __syncthreads to an NVPTX builtin directly. #if !defined(__NVPTX__) __GPU_DEVICE__ void __syncthreads(void) { __gpu_sync_threads(); } #endif @@ -387,10 +342,6 @@ __GPU_DEVICE__ void __threadfence_system(void) { __scoped_atomic_thread_fence(__ATOMIC_SEQ_CST, __MEMORY_SCOPE_SYSTEM); } -//===----------------------------------------------------------------------===// -// Timers -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ long long __clock64(void) { return (long long)__builtin_readcyclecounter(); } @@ -401,10 +352,6 @@ __GPU_DEVICE__ long long wall_clock64(void) { return (long long)__builtin_readsteadycounter(); } -// Warp shuffle / synchronization / reduction intrinsics. -#include <__clang_gpu_intrinsics.h> - -#pragma pop_macro("MAYBE_UNDEF") #pragma pop_macro("__GPU_DEVICE__") #endif // device compile diff --git a/clang/lib/Headers/__clang_gpu_intrinsics.h b/clang/lib/Headers/__clang_gpu_intrinsics.h index 4781bb8bb6f90..ae4644413c3f2 100644 --- a/clang/lib/Headers/__clang_gpu_intrinsics.h +++ b/clang/lib/Headers/__clang_gpu_intrinsics.h @@ -11,14 +11,19 @@ #if defined(__HIP__) || defined(__CUDA__) -#ifndef __GPU_DEVICE__ -#error \ - "__clang_gpu_intrinsics.h must be included via __clang_gpu_device_functions.h" -#endif +#include <__clang_gpu_device_functions.h> +#include <gpuintrin.h> -//===----------------------------------------------------------------------===// -// Wavefront shuffles -//===----------------------------------------------------------------------===// +static inline __attribute__((device)) const struct { + __attribute__((device, always_inline, const)) operator int() const noexcept { + return __gpu_num_lanes(); + } +} warpSize{}; + +#pragma push_macro("__GPU_DEVICE__") +#define __GPU_DEVICE__ static __inline__ __attribute__((device, always_inline)) +#pragma push_macro("MAYBE_UNDEF") +#define MAYBE_UNDEF __attribute__((maybe_undef)) template <typename __T> __GPU_DEVICE__ __T __gpu_shuffle_idx_impl(__T __v, unsigned int __idx, @@ -67,20 +72,12 @@ __GPU_DEVICE__ __T __shfl_xor(MAYBE_UNDEF __T __var, int __lane_mask, __var, (unsigned int)(__tgt >= __width ? __rel : __tgt), __width); } -//===----------------------------------------------------------------------===// -// Warp synchronization -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ void __syncwarp(unsigned long long __mask = -1) { __scoped_atomic_thread_fence(__ATOMIC_RELEASE, __MEMORY_SCOPE_WVFRNT); __gpu_sync_lane(__mask); __scoped_atomic_thread_fence(__ATOMIC_ACQUIRE, __MEMORY_SCOPE_WVFRNT); } -//===----------------------------------------------------------------------===// -// Wave syncrhonization sync aliases. -//===----------------------------------------------------------------------===// - template <typename __MaskT> __GPU_DEVICE__ unsigned long long __ballot_sync(__MaskT __mask, int __pred) { return __ballot(__pred) & (unsigned long long)__mask; @@ -121,10 +118,6 @@ __GPU_DEVICE__ __T __shfl_xor_sync(__MaskT __mask, MAYBE_UNDEF __T __var, return __shfl_xor(__var, __lane_mask, __width); } -//===----------------------------------------------------------------------===// -// Match primitives. -//===----------------------------------------------------------------------===// - template <typename __T> __GPU_DEVICE__ unsigned long long __match_any(__T __value) { if constexpr (sizeof(__T) == sizeof(unsigned long long)) { @@ -161,10 +154,6 @@ __GPU_DEVICE__ unsigned long long __match_all_sync(__MaskT __mask, __T __value, return __match_all(__value, __pred); } -//===----------------------------------------------------------------------===// -// Wave reductions. -//===----------------------------------------------------------------------===// - template <typename __MaskT> __GPU_DEVICE__ unsigned int __reduce_add_sync(__MaskT __mask, unsigned int __val) { @@ -213,10 +202,6 @@ __GPU_DEVICE__ unsigned int __reduce_xor_sync(__MaskT __mask, return __gpu_lane_xor_u32((unsigned long long)__mask, __val); } -//===----------------------------------------------------------------------===// -// Funnel shifts. -//===----------------------------------------------------------------------===// - __GPU_DEVICE__ unsigned int __funnelshift_l(unsigned int __lo, unsigned int __hi, unsigned int __shift) { unsigned int __s = __shift & 31u; @@ -238,5 +223,8 @@ __funnelshift_rc(unsigned int __lo, unsigned int __hi, unsigned int __shift) { return (unsigned int)(((unsigned long long)__hi << 32 | __lo) >> __s); } +#pragma pop_macro("MAYBE_UNDEF") +#pragma pop_macro("__GPU_DEVICE__") + #endif // device compile #endif // __CLANG_GPU_INTRINSICS_H__ diff --git a/clang/lib/Headers/__clang_gpu_runtime_wrapper.h b/clang/lib/Headers/__clang_gpu_runtime_wrapper.h new file mode 100644 index 0000000000000..7ccd71eab96a5 --- /dev/null +++ b/clang/lib/Headers/__clang_gpu_runtime_wrapper.h @@ -0,0 +1,28 @@ +//===---- __clang_gpu_runtime_wrapper.h - Hermetic HIP/CUDA bootstrap ------=== +// +// 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 +// +//===-----------------------------------------------------------------------=== + +// Dialect-neutral bootstrap shared by the hermetic HIP and CUDA offload paths. +// It establishes only the device attribute macros the vendor headers expect and +// pulls in no device code. + +#ifndef __CLANG_GPU_RUNTIME_WRAPPER_H__ +#define __CLANG_GPU_RUNTIME_WRAPPER_H__ + +#if defined(__HIP__) || defined(__CUDA__) + +#ifndef __device__ +#define __host__ __attribute__((host)) +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) +#define __constant__ __attribute__((constant)) +#define __managed__ __attribute__((managed)) +#endif + +#endif // defined(__HIP__) || defined(__CUDA__) +#endif // __CLANG_GPU_RUNTIME_WRAPPER_H__ diff --git a/clang/lib/Headers/hip_wrappers/hip/device_functions.h b/clang/lib/Headers/hip_wrappers/hip/device_functions.h new file mode 100644 index 0000000000000..da7daec29b667 --- /dev/null +++ b/clang/lib/Headers/hip_wrappers/hip/device_functions.h @@ -0,0 +1,17 @@ +//===---- device_functions.h - Hermetic HIP device functions shadow --------=== +// +// 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 +// +//===-----------------------------------------------------------------------=== + +#ifndef __CLANG_HIP_DEVICE_FUNCTIONS_SHADOW_H__ +#define __CLANG_HIP_DEVICE_FUNCTIONS_SHADOW_H__ + +#include <__clang_gpu_builtin_vars.h> +#include <__clang_gpu_device_functions.h> +#include <__clang_gpu_intrinsics.h> +#include <__clang_gpu_runtime_wrapper.h> + +#endif // __CLANG_HIP_DEVICE_FUNCTIONS_SHADOW_H__ diff --git a/clang/lib/Headers/hip_wrappers/hip/hip_runtime.h b/clang/lib/Headers/hip_wrappers/hip/hip_runtime.h new file mode 100644 index 0000000000000..5d9414495ad83 --- /dev/null +++ b/clang/lib/Headers/hip_wrappers/hip/hip_runtime.h @@ -0,0 +1,23 @@ +//===---- hip_runtime.h - Hermetic HIP runtime shadow ---------------------=== +// +// 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 +// +//===-----------------------------------------------------------------------=== + +#ifndef __CLANG_HIP_RUNTIME_SHADOW_H__ +#define __CLANG_HIP_RUNTIME_SHADOW_H__ + +#include <__clang_gpu_runtime_wrapper.h> + +#if __has_include_next(<hip/hip_runtime.h>) +#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_RUNTIME_H +#include_next <hip/hip_runtime.h> +#endif + +#include <__clang_gpu_builtin_vars.h> +#include <__clang_gpu_device_functions.h> +#include <__clang_gpu_intrinsics.h> + +#endif // __CLANG_HIP_RUNTIME_SHADOW_H__ diff --git a/clang/test/Driver/hip-device-libs-llvm-env.hip b/clang/test/Driver/hip-device-libs-llvm-env.hip index 7232285237550..86ea18140e4c7 100644 --- a/clang/test/Driver/hip-device-libs-llvm-env.hip +++ b/clang/test/Driver/hip-device-libs-llvm-env.hip @@ -24,7 +24,7 @@ // RUN: 2>&1 | FileCheck --check-prefix=INC %s // INC: "-cc1" "-triple" "amdgcn-amd-amdhsa-llvm" -// INC-SAME: "-include" "__clang_gpu_device_functions.h" +// INC-SAME: "-internal-isystem" "{{.*}}hip_wrappers" // RUN: %clang -### --target=x86_64-linux-gnu -nogpuinc --rocm-path=%S/Inputs/rocm \ // RUN: --offload-targets=amdgcn-amd-amdhsa-llvm --offload-arch=gfx90a \ @@ -38,4 +38,4 @@ // RUN: %S/Inputs/hip_multiple_inputs/b.hip \ // RUN: 2>&1 | FileCheck --check-prefix=NOINC %s -// NOINC-NOT: __clang_gpu_device_functions.h +// NOINC-NOT: hip_wrappers diff --git a/clang/test/Headers/gpu-device-functions.cpp b/clang/test/Headers/gpu-device-functions.cpp index 09c9b467a293c..373b031cdaa8d 100644 --- a/clang/test/Headers/gpu-device-functions.cpp +++ b/clang/test/Headers/gpu-device-functions.cpp @@ -3,47 +3,68 @@ // 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_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // RUN: %clang_cc1 -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/../../lib/Headers \ // RUN: -triple amdgpu11.00-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -x hip -fcuda-is-device -fsyntax-only -verify %s \ -// RUN: -include __clang_gpu_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // HIP on SPIR-V. // RUN: %clang_cc1 -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_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // CUDA on NVPTX. // RUN: %clang_cc1 -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_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // HIP host compilation. // RUN: %clang_cc1 -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_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // // HIP host compilation with a SPIR-V device. // RUN: %clang_cc1 -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/../../lib/Headers \ // RUN: -triple x86_64-unknown-unknown -aux-triple spirv64-amd-amdhsa \ // RUN: -x hip -fsyntax-only -verify %s \ -// RUN: -include __clang_gpu_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // // CUDA host compilation. // RUN: %clang_cc1 -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_device_functions.h +// RUN: -include __clang_gpu_runtime_wrapper.h \ +// RUN: -include __clang_gpu_builtin_vars.h \ +// RUN: -include __clang_gpu_device_functions.h \ +// RUN: -include __clang_gpu_intrinsics.h // expected-no-diagnostics _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
