Skip to content

Commit a866ff4

Browse files
committed
[ROCm] Guard placement new/delete to fix build on ROCm 7.12+
ROCm 7.12 (clang 22) provides __device__ placement new/delete via cuda_wrappers/new, causing redefinition errors. Guard with __CLANG_CUDA_WRAPPERS_NEW so the manual definitions are only compiled on older ROCm versions that lack them.
1 parent ed8c2aa commit a866ff4

1 file changed

Lines changed: 3 additions & 0 deletions

File tree

mlx/backend/rocm/sort.hip

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,12 +10,15 @@
1010

1111
// Workaround: rocprim headers use placement new in __device__ code,
1212
// which requires __device__ overloads of operator new/delete.
13+
// ROCm 7.12+ (clang 22+) already provides these via cuda_wrappers/new.
1314
#ifdef __HIP_DEVICE_COMPILE__
15+
#ifndef __CLANG_CUDA_WRAPPERS_NEW
1416
__device__ inline void* operator new(size_t, void* p) noexcept { return p; }
1517
__device__ inline void* operator new[](size_t, void* p) noexcept { return p; }
1618
__device__ inline void operator delete(void*, void*) noexcept {}
1719
__device__ inline void operator delete[](void*, void*) noexcept {}
1820
#endif
21+
#endif
1922

2023
#include <rocprim/rocprim.hpp>
2124
#include <cassert>

0 commit comments

Comments
 (0)