template <typename T>
struct CudaSupportedTypeImpl {
using type = T;
};
template <>
struct CudaSupportedTypeImpl<long long> {
using type = unsigned long long;
};
template <>
struct CudaSupportedTypeImpl<unsigned long> {
using type =
typename std::conditional<sizeof(unsigned long) == sizeof(unsigned int),
unsigned int, unsigned long long>::type;
};
template <>
struct CudaSupportedTypeImpl<long> {
using type = typename CudaSupportedTypeImpl<unsigned long>::type;
};
template <typename T>
using CudaSupportedType = typename CudaSupportedTypeImpl<T>::type;
template <typename T>
__device__ CudaSupportedType<T>* ToCudaSupportedPtr(T* ptr) {
return reinterpret_cast<CudaSupportedType<T>*>(ptr);
}
使用:
template <typename T, typename U>
__device__ detail::ToTypeIfConvertible<U, T> GpuAtomicAdd(T* ptr, U value) {
return atomicAdd(detail::ToCudaSupportedPtr(ptr), value);
}
|