Author: fineg74 Date: 2026-02-17T20:53:17Z New Revision: 5c518583d0b65228a88f41ea649bc867a33f8222
URL: https://github.com/llvm/llvm-project/commit/5c518583d0b65228a88f41ea649bc867a33f8222 DIFF: https://github.com/llvm/llvm-project/commit/5c518583d0b65228a88f41ea649bc867a33f8222.diff LOG: [OFFLOAD] Add headers to support complex math for spirv (#179846) This is to add support for complex operations for OpenMP kernels when compiled with spirv. This is the first PR to add support for spirv in OpenMP wrapper headers to reduce the size of PRs Added: Modified: clang/lib/Headers/__clang_cuda_complex_builtins.h clang/lib/Headers/openmp_wrappers/complex clang/lib/Headers/openmp_wrappers/complex.h Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_cuda_complex_builtins.h b/clang/lib/Headers/__clang_cuda_complex_builtins.h index e3038dcb8fd36..d265b000b7ef2 100644 --- a/clang/lib/Headers/__clang_cuda_complex_builtins.h +++ b/clang/lib/Headers/__clang_cuda_complex_builtins.h @@ -16,19 +16,25 @@ // to work with CUDA and OpenMP target offloading [in C and C++ mode].) #pragma push_macro("__DEVICE__") -#if defined(__OPENMP_NVPTX__) || defined(__OPENMP_AMDGCN__) +#if defined(__OPENMP_NVPTX__) || defined(__OPENMP_AMDGCN__) || \ + defined(__OPENMP_SPIRV__) #pragma omp declare target #define __DEVICE__ __attribute__((noinline, nothrow, cold, weak)) #else #define __DEVICE__ __device__ inline #endif -#ifdef __NVPTX__ +#if defined(__NVPTX__) // FIXME: NVPTX should use generic builtins. #define _SCALBNd __nv_scalbn #define _SCALBNf __nv_scalbnf #define _LOGBd __nv_logb #define _LOGBf __nv_logbf +#elif defined(__OPENMP_SPIRV__) +#define _SCALBNd __spirv_ocl_ldexp +#define _SCALBNf __spirv_ocl_ldexp +#define _LOGBd __spirv_ocl_logb +#define _LOGBf __spirv_ocl_logb #else #define _SCALBNd __builtin_scalbn #define _SCALBNf __builtin_scalbnf @@ -220,7 +226,8 @@ __DEVICE__ float _Complex __divsc3(float __a, float __b, float __c, float __d) { #undef _LOGBd #undef _LOGBf -#if defined(__OPENMP_NVPTX__) || defined(__OPENMP_AMDGCN__) +#if defined(__OPENMP_NVPTX__) || defined(__OPENMP_AMDGCN__) || \ + defined(__OPENMP_SPIRV__) #pragma omp end declare target #endif diff --git a/clang/lib/Headers/openmp_wrappers/complex b/clang/lib/Headers/openmp_wrappers/complex index 1ceecc1af8aec..afa8e3eb3d835 100644 --- a/clang/lib/Headers/openmp_wrappers/complex +++ b/clang/lib/Headers/openmp_wrappers/complex @@ -29,6 +29,12 @@ #undef __OPENMP_AMDGCN__ #endif // __AMDGCN__ +#ifdef __SPIRV__ +#define __OPENMP_SPIRV__ +#include <__clang_cuda_complex_builtins.h> +#undef __OPENMP_SPIRV__ +#endif // __SPIRV__ + #endif // Grab the host header too. @@ -45,7 +51,7 @@ #ifndef _LIBCPP_STD_VER #pragma omp begin declare variant match( \ - device = {arch(amdgcn, nvptx, nvptx64)}, \ + device = {arch(amdgcn, nvptx, nvptx64, spirv64)}, \ implementation = {extension(match_any, allow_templates)}) #include <complex_cmath.h> diff --git a/clang/lib/Headers/openmp_wrappers/complex.h b/clang/lib/Headers/openmp_wrappers/complex.h index 7e7c0866426bc..92839088b47d5 100644 --- a/clang/lib/Headers/openmp_wrappers/complex.h +++ b/clang/lib/Headers/openmp_wrappers/complex.h @@ -29,6 +29,12 @@ #undef __OPENMP_AMDGCN__ #endif +#ifdef __SPIRV__ +#define __OPENMP_SPIRV__ +#include <__clang_cuda_complex_builtins.h> +#undef __OPENMP_SPIRV__ +#endif // __SPIRV__ + #endif // Grab the host header too. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
