109 lines · plain
1/*===---- new - CUDA wrapper for <new> -------------------------------------===2 *3 * Permission is hereby granted, free of charge, to any person obtaining a copy4 * of this software and associated documentation files (the "Software"), to deal5 * in the Software without restriction, including without limitation the rights6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell7 * copies of the Software, and to permit persons to whom the Software is8 * furnished to do so, subject to the following conditions:9 *10 * The above copyright notice and this permission notice shall be included in11 * all copies or substantial portions of the Software.12 *13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER17 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,18 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN19 * THE SOFTWARE.20 *21 *===-----------------------------------------------------------------------===22 */23 24#ifndef __CLANG_CUDA_WRAPPERS_NEW25#define __CLANG_CUDA_WRAPPERS_NEW26 27#include_next <new>28 29#if !defined(__device__)30// The header has been included too early from the standard C++ library31// and CUDA-specific macros are not available yet.32// Undo the include guard and try again later.33#undef __CLANG_CUDA_WRAPPERS_NEW34#else35 36#pragma push_macro("CUDA_NOEXCEPT")37#if __cplusplus >= 201103L38#define CUDA_NOEXCEPT noexcept39#else40#define CUDA_NOEXCEPT41#endif42 43// Device overrides for non-placement new and delete.44__device__ inline void *operator new(__SIZE_TYPE__ size) {45 if (size == 0) {46 size = 1;47 }48 return ::malloc(size);49}50__device__ inline void *operator new(__SIZE_TYPE__ size,51 const std::nothrow_t &) CUDA_NOEXCEPT {52 return ::operator new(size);53}54 55__device__ inline void *operator new[](__SIZE_TYPE__ size) {56 return ::operator new(size);57}58__device__ inline void *operator new[](__SIZE_TYPE__ size,59 const std::nothrow_t &) {60 return ::operator new(size);61}62 63__device__ inline void operator delete(void* ptr) CUDA_NOEXCEPT {64 if (ptr) {65 ::free(ptr);66 }67}68__device__ inline void operator delete(void *ptr,69 const std::nothrow_t &) CUDA_NOEXCEPT {70 ::operator delete(ptr);71}72 73__device__ inline void operator delete[](void* ptr) CUDA_NOEXCEPT {74 ::operator delete(ptr);75}76__device__ inline void operator delete[](void *ptr,77 const std::nothrow_t &) CUDA_NOEXCEPT {78 ::operator delete(ptr);79}80 81// Sized delete, C++14 only.82#if __cplusplus >= 201402L83__device__ inline void operator delete(void *ptr,84 __SIZE_TYPE__ size) CUDA_NOEXCEPT {85 ::operator delete(ptr);86}87__device__ inline void operator delete[](void *ptr,88 __SIZE_TYPE__ size) CUDA_NOEXCEPT {89 ::operator delete(ptr);90}91#endif92 93// Device overrides for placement new and delete.94#if !(_LIBCPP_STD_VER >= 26 || __cpp_lib_constexpr_new >= 202406L)95__device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT {96 return __ptr;97}98__device__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT {99 return __ptr;100}101#endif102__device__ inline void operator delete(void *, void *) CUDA_NOEXCEPT {}103__device__ inline void operator delete[](void *, void *) CUDA_NOEXCEPT {}104 105#pragma pop_macro("CUDA_NOEXCEPT")106 107#endif // __device__108#endif // include guard109