使用标记调度存在自定义 CUDA 分配器时重载类新运算符

Posted 2023-02-22

【中文标题】使用标记调度存在自定义 CUDA 分配器时重载类新运算符

【英文标题】：overload class new operator when custom CUDA allocator is present using tag-dispatch

【发布时间】：2015-07-20 11:14:47

【问题描述】：

创建了一个 STL 风格的分配器感知类，我正在尝试将其与自定义 CUDA 分配器一起使用。 CUDA 分配器可以很好地在统一内存中分配数据存储，但是为了在主机和设备上都可以访问 this 指针，我需要确保每当数据在统一内存中分配时，该类也是。

为了解决这个问题，我认为一个简单的标签调度是合适的。如果分配器是cudaAllocator，new 应该在统一内存中创建类，如果不是，它应该只返回常规的 new 输出。不幸的是，我认为我错过了标签调度的工作原理。这是课程的相关部分：

#ifdef __CUDACC__
public:

    using cudaAllocator_tag = std::true_type;
    using hostAllocator_tag = std::false_type;

    void *operator new(size_t len)
    
        return new(len, std::is_same<Alloc, cudaAllocator<value_type>>());
    

    void operator delete(void *ptr) 
    
        return delete(ptr, std::is_same<Alloc, cudaAllocator<value_type>>());
    

    void *operator new(size_t len, cudaAllocator_tag)
    
        void *ptr;
        cudaMallocManaged(&ptr, len);
        cudaDeviceSynchronize();
        return ptr;
    

    void operator delete(void *ptr, cudaAllocator_tag) 
    
        cudaDeviceSynchronize();
        cudaFree(ptr);
    

    void *operator new(size_t len, hostAllocator_tag)
    
        return ::new(len);
    

    void operator delete(void *ptr, hostAllocator_tag)
    
        ::delete(ptr);
    

#endif // __CUDACC__

但是 (NVCC) 编译器抛出以下错误：

2> error : expected a type specifier
2>  
2>            detected during instantiation of "void *CircularQueue<T, Alloc>::operator new(size_t) [with T=float, Alloc=cudaAllocator<float>]" 
2>  
2>  main.cu(21): here
2>  
2>  
2>  
2> error : no instance of overloaded "operator new" matches the argument list
2>  
2>              argument types are: (unsigned long long, size_t, std::is_same<cudaAllocator<float>, cudaAllocator<float>>)
2>  
2>            detected during instantiation of "void *CircularQueue<T, Alloc>::operator new(size_t) [with T=float, Alloc=cudaAllocator<float>]" 
2>  
2>  main.cu(21): here

关于我做错了什么有什么想法吗？

【参考方案1】：

这里有几个问题：

new 缺少要创建的对象 的标识符。假设你的容器叫myContainer 应该是：

static void *operator new(size_t len)

    return new myContainer(len, std::is_same<Alloc, cudaAllocator<value_type>>());

operator delete 的参数不能像 new 一样被重载。您可以通过让delete 调用自定义destroy 函数来解决此问题，使用内联和标记调度来避免任何运行时损失。

static void operator delete(void *ptr) 

    destroy(ptr, std::is_same<Alloc, cudaAllocator<value_type>>());

为了避免混淆/无限递归，最好也使用new。

完整解决方案：

#ifdef __CUDACC__
public:

    using cudaAllocator_tag = std::true_type;
    using hostAllocator_tag = std::false_type;
    using isCudaAllocator   = typename std::is_same<Alloc, cudaAllocator<value_type>>;

    static void *operator new(size_t len)
    
        return create(len, isCudaAllocator());
    

    static void operator delete(void *ptr) 
    
        destroy(ptr, isCudaAllocator());
    

protected:

    static inline void *create(size_t len, cudaAllocator_tag)
    
        void *ptr;
        cudaMallocManaged(&ptr, len);
        cudaDeviceSynchronize();
        return ptr;
    

    static inline void destroy(void *ptr, cudaAllocator_tag)
    
        cudaDeviceSynchronize();
        cudaFree(ptr);
    

    static inline void *create(size_t len, hostAllocator_tag)
    
        return ::new CircularQueue(len);
    

    static inline void destroy(void *ptr, hostAllocator_tag)
    
        ::delete(static_cast<CircularQueue*>(ptr));
    
#endif // __CUDACC__