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 <stdint.h> 44 #endif 45 46 #define __OPENMP_AMDGCN__ 47 #pragma push_macro("__device__") 48 #define __device__ 49 50 /// Include declarations for libdevice functions. 51 #include <__clang_hip_libdevice_declares.h> 52 53 #pragma pop_macro("__device__") 54 #undef __OPENMP_AMDGCN__ 55 56 #pragma omp end declare variant 57 #endif 58 59 #ifdef __cplusplus 60 } // extern "C" 61 #endif 62 63 // Ensure we make `_ZdlPv`, aka. `operator delete(void*)` available without the 64 // need to `include <new>` in C++ mode. 65 #ifdef __cplusplus 66 67 // We require malloc/free. 68 #include <cstdlib> 69 70 #pragma push_macro("OPENMP_NOEXCEPT") 71 #if __cplusplus >= 201103L 72 #define OPENMP_NOEXCEPT noexcept 73 #else 74 #define OPENMP_NOEXCEPT 75 #endif 76 77 // Device overrides for non-placement new and delete. new(__SIZE_TYPE__ size)78inline void *operator new(__SIZE_TYPE__ size) { 79 if (size == 0) 80 size = 1; 81 return ::malloc(size); 82 } 83 84 inline void *operator new[](__SIZE_TYPE__ size) { return ::operator new(size); } 85 delete(void * ptr)86inline void operator delete(void *ptr)OPENMP_NOEXCEPT { ::free(ptr); } 87 88 inline void operator delete[](void *ptr) OPENMP_NOEXCEPT { 89 ::operator delete(ptr); 90 } 91 92 // Sized delete, C++14 only. 93 #if __cplusplus >= 201402L delete(void * ptr,__SIZE_TYPE__ size)94inline void operator delete(void *ptr, __SIZE_TYPE__ size)OPENMP_NOEXCEPT { 95 ::operator delete(ptr); 96 } 97 inline void operator delete[](void *ptr, __SIZE_TYPE__ size) OPENMP_NOEXCEPT { 98 ::operator delete(ptr); 99 } 100 #endif 101 102 #pragma pop_macro("OPENMP_NOEXCEPT") 103 #endif 104 105 #endif 106