1 /*===- __clang_openmp_device_functions.h - OpenMP device function declares -=== 2 * 3 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 * See https://llvm.org/LICENSE.txt for license information. 5 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 * 7 *===-----------------------------------------------------------------------=== 8 */ 9 10 #ifndef __CLANG_OPENMP_DEVICE_FUNCTIONS_H__ 11 #define __CLANG_OPENMP_DEVICE_FUNCTIONS_H__ 12 13 #ifndef _OPENMP 14 #error "This file is for OpenMP compilation only." 15 #endif 16 17 #ifdef __cplusplus 18 extern "C" { 19 #endif 20 21 #pragma omp begin declare variant match( \ 22 device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) 23 24 #define __CUDA__ 25 #define __OPENMP_NVPTX__ 26 27 /// Include declarations for libdevice functions. 28 #include <__clang_cuda_libdevice_declares.h> 29 30 /// Provide definitions for these functions. 31 #include <__clang_cuda_device_functions.h> 32 33 #undef __OPENMP_NVPTX__ 34 #undef __CUDA__ 35 36 #pragma omp end declare variant 37 38 #ifdef __AMDGCN__ 39 #pragma omp begin declare variant match(device = {arch(amdgcn)}) 40 41 // Import types which will be used by __clang_hip_libdevice_declares.h 42 #ifndef __cplusplus 43 #include <stdbool.h> 44 #include <stdint.h> 45 #endif 46 47 #define __OPENMP_AMDGCN__ 48 #pragma push_macro("__device__") 49 #define __device__ 50 51 /// Include declarations for libdevice functions. 52 #include <__clang_hip_libdevice_declares.h> 53 54 #pragma pop_macro("__device__") 55 #undef __OPENMP_AMDGCN__ 56 57 #pragma omp end declare variant 58 #endif 59 60 #ifdef __cplusplus 61 } // extern "C" 62 #endif 63 64 // Ensure we make `_ZdlPv`, aka. `operator delete(void*)` available without the 65 // need to `include <new>` in C++ mode. 66 #ifdef __cplusplus 67 68 // We require malloc/free. 69 #include <cstdlib> 70 71 #pragma push_macro("OPENMP_NOEXCEPT") 72 #if __cplusplus >= 201103L 73 #define OPENMP_NOEXCEPT noexcept 74 #else 75 #define OPENMP_NOEXCEPT 76 #endif 77 78 // Device overrides for non-placement new and delete. new(__SIZE_TYPE__ size)79inline void *operator new(__SIZE_TYPE__ size) { 80 if (size == 0) 81 size = 1; 82 return ::malloc(size); 83 } 84 85 inline void *operator new[](__SIZE_TYPE__ size) { return ::operator new(size); } 86 delete(void * ptr)87inline void operator delete(void *ptr)OPENMP_NOEXCEPT { ::free(ptr); } 88 89 inline void operator delete[](void *ptr) OPENMP_NOEXCEPT { 90 ::operator delete(ptr); 91 } 92 93 // Sized delete, C++14 only. 94 #if __cplusplus >= 201402L delete(void * ptr,__SIZE_TYPE__ size)95inline void operator delete(void *ptr, __SIZE_TYPE__ size)OPENMP_NOEXCEPT { 96 ::operator delete(ptr); 97 } 98 inline void operator delete[](void *ptr, __SIZE_TYPE__ size) OPENMP_NOEXCEPT { 99 ::operator delete(ptr); 100 } 101 #endif 102 103 #pragma pop_macro("OPENMP_NOEXCEPT") 104 #endif 105 106 #endif 107