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