10b57cec5SDimitry Andric/*===---- new - CUDA wrapper for <new> -------------------------------------=== 20b57cec5SDimitry Andric * 30b57cec5SDimitry Andric * Permission is hereby granted, free of charge, to any person obtaining a copy 40b57cec5SDimitry Andric * of this software and associated documentation files (the "Software"), to deal 50b57cec5SDimitry Andric * in the Software without restriction, including without limitation the rights 60b57cec5SDimitry Andric * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 70b57cec5SDimitry Andric * copies of the Software, and to permit persons to whom the Software is 80b57cec5SDimitry Andric * furnished to do so, subject to the following conditions: 90b57cec5SDimitry Andric * 100b57cec5SDimitry Andric * The above copyright notice and this permission notice shall be included in 110b57cec5SDimitry Andric * all copies or substantial portions of the Software. 120b57cec5SDimitry Andric * 130b57cec5SDimitry Andric * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 140b57cec5SDimitry Andric * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 150b57cec5SDimitry Andric * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 160b57cec5SDimitry Andric * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 170b57cec5SDimitry Andric * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 180b57cec5SDimitry Andric * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 190b57cec5SDimitry Andric * THE SOFTWARE. 200b57cec5SDimitry Andric * 210b57cec5SDimitry Andric *===-----------------------------------------------------------------------=== 220b57cec5SDimitry Andric */ 230b57cec5SDimitry Andric 240b57cec5SDimitry Andric#ifndef __CLANG_CUDA_WRAPPERS_NEW 250b57cec5SDimitry Andric#define __CLANG_CUDA_WRAPPERS_NEW 260b57cec5SDimitry Andric 270b57cec5SDimitry Andric#include_next <new> 280b57cec5SDimitry Andric 290b57cec5SDimitry Andric#if !defined(__device__) 300b57cec5SDimitry Andric// The header has been included too early from the standard C++ library 310b57cec5SDimitry Andric// and CUDA-specific macros are not available yet. 320b57cec5SDimitry Andric// Undo the include guard and try again later. 330b57cec5SDimitry Andric#undef __CLANG_CUDA_WRAPPERS_NEW 340b57cec5SDimitry Andric#else 350b57cec5SDimitry Andric 360b57cec5SDimitry Andric#pragma push_macro("CUDA_NOEXCEPT") 370b57cec5SDimitry Andric#if __cplusplus >= 201103L 380b57cec5SDimitry Andric#define CUDA_NOEXCEPT noexcept 390b57cec5SDimitry Andric#else 400b57cec5SDimitry Andric#define CUDA_NOEXCEPT 410b57cec5SDimitry Andric#endif 420b57cec5SDimitry Andric 430b57cec5SDimitry Andric// Device overrides for non-placement new and delete. 440b57cec5SDimitry Andric__device__ inline void *operator new(__SIZE_TYPE__ size) { 450b57cec5SDimitry Andric if (size == 0) { 460b57cec5SDimitry Andric size = 1; 470b57cec5SDimitry Andric } 480b57cec5SDimitry Andric return ::malloc(size); 490b57cec5SDimitry Andric} 500b57cec5SDimitry Andric__device__ inline void *operator new(__SIZE_TYPE__ size, 510b57cec5SDimitry Andric const std::nothrow_t &) CUDA_NOEXCEPT { 520b57cec5SDimitry Andric return ::operator new(size); 530b57cec5SDimitry Andric} 540b57cec5SDimitry Andric 550b57cec5SDimitry Andric__device__ inline void *operator new[](__SIZE_TYPE__ size) { 560b57cec5SDimitry Andric return ::operator new(size); 570b57cec5SDimitry Andric} 580b57cec5SDimitry Andric__device__ inline void *operator new[](__SIZE_TYPE__ size, 590b57cec5SDimitry Andric const std::nothrow_t &) { 600b57cec5SDimitry Andric return ::operator new(size); 610b57cec5SDimitry Andric} 620b57cec5SDimitry Andric 630b57cec5SDimitry Andric__device__ inline void operator delete(void* ptr) CUDA_NOEXCEPT { 640b57cec5SDimitry Andric if (ptr) { 650b57cec5SDimitry Andric ::free(ptr); 660b57cec5SDimitry Andric } 670b57cec5SDimitry Andric} 680b57cec5SDimitry Andric__device__ inline void operator delete(void *ptr, 690b57cec5SDimitry Andric const std::nothrow_t &) CUDA_NOEXCEPT { 700b57cec5SDimitry Andric ::operator delete(ptr); 710b57cec5SDimitry Andric} 720b57cec5SDimitry Andric 730b57cec5SDimitry Andric__device__ inline void operator delete[](void* ptr) CUDA_NOEXCEPT { 740b57cec5SDimitry Andric ::operator delete(ptr); 750b57cec5SDimitry Andric} 760b57cec5SDimitry Andric__device__ inline void operator delete[](void *ptr, 770b57cec5SDimitry Andric const std::nothrow_t &) CUDA_NOEXCEPT { 780b57cec5SDimitry Andric ::operator delete(ptr); 790b57cec5SDimitry Andric} 800b57cec5SDimitry Andric 810b57cec5SDimitry Andric// Sized delete, C++14 only. 820b57cec5SDimitry Andric#if __cplusplus >= 201402L 830b57cec5SDimitry Andric__device__ inline void operator delete(void *ptr, 840b57cec5SDimitry Andric __SIZE_TYPE__ size) CUDA_NOEXCEPT { 850b57cec5SDimitry Andric ::operator delete(ptr); 860b57cec5SDimitry Andric} 870b57cec5SDimitry Andric__device__ inline void operator delete[](void *ptr, 880b57cec5SDimitry Andric __SIZE_TYPE__ size) CUDA_NOEXCEPT { 890b57cec5SDimitry Andric ::operator delete(ptr); 900b57cec5SDimitry Andric} 910b57cec5SDimitry Andric#endif 920b57cec5SDimitry Andric 930b57cec5SDimitry Andric// Device overrides for placement new and delete. 940b57cec5SDimitry Andric__device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT { 950b57cec5SDimitry Andric return __ptr; 960b57cec5SDimitry Andric} 970b57cec5SDimitry Andric__device__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT { 980b57cec5SDimitry Andric return __ptr; 99} 100__device__ inline void operator delete(void *, void *) CUDA_NOEXCEPT {} 101__device__ inline void operator delete[](void *, void *) CUDA_NOEXCEPT {} 102 103#pragma pop_macro("CUDA_NOEXCEPT") 104 105#endif // __device__ 106#endif // include guard 107