cuda::ptr_rebind
template <typename U, typename T>
[[nodiscard]] __host__ __device__ inline
U* ptr_rebind(T* ptr) noexcept
template <typename U, typename T>
[[nodiscard]] __host__ __device__ inline
const U* ptr_rebind(const T* ptr) noexcept
template <typename U, typename T>
[[nodiscard]] __host__ __device__ inline
volatile U* ptr_rebind(volatile T* ptr) noexcept
template <typename U, typename T>
[[nodiscard]] __host__ __device__ inline
const volatile U* ptr_rebind(const volatile T* ptr) noexcept
The functions return the pointer ptr
cast to type U*
or const U*
. They are shorter and safer alternative to reinterpret_cast
.
Parameters
ptr
: The pointer.
Return value
The pointer cast to type
U*
orconst U*
.
Constraints
ptr
must be aligned toalignof(U)
andalignof(T)
.
Performance considerations
The returned pointer is decorated with
__builtin_assume_aligned
with thealignof(U)
value to help the compiler generate better code.
Example
#include <cuda/memory>
#include <cuda/std/cstdint>
__global__ void kernel(const int* ptr, volatile int* ptr2) {
const uint64_t* ptr_res1 = cuda::ptr_rebind<uint64_t>(ptr);
volatile uint64_t* ptr_res2 = cuda::ptr_rebind<uint64_t>(ptr2);
}
int main() {
int* ptr;
cudaMalloc(&ptr, 100 * sizeof(int));
kernel<<<1, 1>>>(ptr);
cudaDeviceSynchronize();
return 0;
}