1f85ae058SJohannes Doerfert /*===---- openmp_wrapper/math.h -------- OpenMP math.h intercept ------ c++ -=== 2e62c693cSGheorghe-Teodor Bercea * 3e62c693cSGheorghe-Teodor Bercea * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4e62c693cSGheorghe-Teodor Bercea * See https://llvm.org/LICENSE.txt for license information. 5e62c693cSGheorghe-Teodor Bercea * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6e62c693cSGheorghe-Teodor Bercea * 7e62c693cSGheorghe-Teodor Bercea *===-----------------------------------------------------------------------=== 8e62c693cSGheorghe-Teodor Bercea */ 9e62c693cSGheorghe-Teodor Bercea 1017d83342SJohannes Doerfert // If we are in C++ mode and include <math.h> (not <cmath>) first, we still need 1117d83342SJohannes Doerfert // to make sure <cmath> is read first. The problem otherwise is that we haven't 1217d83342SJohannes Doerfert // seen the declarations of the math.h functions when the system math.h includes 1317d83342SJohannes Doerfert // our cmath overlay. However, our cmath overlay, or better the underlying 1417d83342SJohannes Doerfert // overlay, e.g. CUDA, uses the math.h functions. Since we haven't declared them 1517d83342SJohannes Doerfert // yet we get errors. CUDA avoids this by eagerly declaring all math functions 1617d83342SJohannes Doerfert // (in the __device__ space) but we cannot do this. Instead we break the 1717d83342SJohannes Doerfert // dependence by forcing cmath to go first. While our cmath will in turn include 1817d83342SJohannes Doerfert // this file, the cmath guards will prevent recursion. 1917d83342SJohannes Doerfert #ifdef __cplusplus 2017d83342SJohannes Doerfert #include <cmath> 2117d83342SJohannes Doerfert #endif 2217d83342SJohannes Doerfert 23f85ae058SJohannes Doerfert #ifndef __CLANG_OPENMP_MATH_H__ 24f85ae058SJohannes Doerfert #define __CLANG_OPENMP_MATH_H__ 25e62c693cSGheorghe-Teodor Bercea 26f85ae058SJohannes Doerfert #ifndef _OPENMP 27f85ae058SJohannes Doerfert #error "This file is for OpenMP compilation only." 28e62c693cSGheorghe-Teodor Bercea #endif 29e62c693cSGheorghe-Teodor Bercea 30f85ae058SJohannes Doerfert #include_next <math.h> 31f85ae058SJohannes Doerfert 32f85ae058SJohannes Doerfert // We need limits.h for __clang_cuda_math.h below and because it should not hurt 33f85ae058SJohannes Doerfert // we include it eagerly here. 34f85ae058SJohannes Doerfert #include <limits.h> 35f85ae058SJohannes Doerfert 36f85ae058SJohannes Doerfert // We need stdlib.h because (for now) __clang_cuda_math.h below declares `abs` 37f85ae058SJohannes Doerfert // which should live in stdlib.h. 38f85ae058SJohannes Doerfert #include <stdlib.h> 39f85ae058SJohannes Doerfert 40f85ae058SJohannes Doerfert #pragma omp begin declare variant match( \ 41f85ae058SJohannes Doerfert device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) 42f85ae058SJohannes Doerfert 43f85ae058SJohannes Doerfert #define __CUDA__ 447f1e6fcfSJohannes Doerfert #define __OPENMP_NVPTX__ 45f85ae058SJohannes Doerfert #include <__clang_cuda_math.h> 467f1e6fcfSJohannes Doerfert #undef __OPENMP_NVPTX__ 47f85ae058SJohannes Doerfert #undef __CUDA__ 48f85ae058SJohannes Doerfert 49f85ae058SJohannes Doerfert #pragma omp end declare variant 50f85ae058SJohannes Doerfert 51*713a5d12SPushpinder Singh #ifdef __AMDGCN__ 52*713a5d12SPushpinder Singh #pragma omp begin declare variant match(device = {arch(amdgcn)}) 53*713a5d12SPushpinder Singh 54*713a5d12SPushpinder Singh #define __OPENMP_AMDGCN__ 55*713a5d12SPushpinder Singh #include <__clang_hip_math.h> 56*713a5d12SPushpinder Singh #undef __OPENMP_AMDGCN__ 57*713a5d12SPushpinder Singh 58*713a5d12SPushpinder Singh #pragma omp end declare variant 59*713a5d12SPushpinder Singh #endif 60*713a5d12SPushpinder Singh 61f85ae058SJohannes Doerfert #endif 62