c++ - 当使用标签调度存在自定义 CUDA 分配器时重载类新运算符

Question

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

为了解决这个问题，我认为一个简单的标签分派(dispatch)是合适的。如果分配器是 cudaAllocator，new 应该在统一内存中创建类，如果不是，它应该只返回常规的 new 输出。不幸的是，我认为我遗漏了一些关于 tag-dispatch 的工作原理。这是类(class)的相关部分:

#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

对我做错了什么有什么想法吗？

Answer 1

这里有几个问题:

1. new 缺少要创建的对象 的标识符。假设您的容器名为 myContainer 它应该是:

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

2. 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__