如何在 cuda 程序中实现 [] 运算符重载？答案

【问题标题】：How to implement [] operator overloading in cuda program?如何在 cuda 程序中实现 [] 运算符重载？
【发布时间】：2016-02-14 12:17:19
【问题描述】：

我正在 Cuda 中实现 device_vector，我正在从著名的库 Thust 中获取想法。

现在要访问和修改 device_vector (v) 中的元素，我需要执行 v[N] = x。为此，我需要重载 [] 运算符。

这是用于重载 [] 运算符的代码：

T& operator[] (unsigned int index)
{
    if (index >= numEle)
        return ptr[0];
    else
        return ptr[index];
}

问题是：要修改设备内存中的任何内存位置，我们需要进行 Cuda 内核调用，而 Cuda 内核调用无法返回任何内容。

就 [] 重载而言，它返回对我们要修改的元素的引用。

我们如何为 Cuda 内核做到这一点？

注意：我知道 Thrust Library 会以某种方式做到这一点，但我无法理解。

【问题讨论】：

您可以返回一个临时对象，允许从unsigned int转换和赋值。
@Kos 我听不懂...你能不能给我更多的细节。
这个函数应该运行在主机还是设备上？
将从主机使用运算符。但是要修改设备中的内存，我们需要进行内核调用。
这不是启动的内核名称。我并不认为推力可能会使用内核启动来完成各种事情，但在您的测试用例中，nvprof 输出中的内核启动行 precedes [CUDA memcpy DtoH] 行。您正在见证的内核启动发生在 thrust::device_vector 的实例化时。 [CUDA memcpy DtoH] 是 nvprof 报告发生 cudaMemcpy 操作（不是内核启动）的方式，并且该操作的指定方向是 cudaMemcpyDeviceToHost（这是由于您的 v[N] “取消引用”而发生的）。

标签： c++ cuda

【解决方案1】：

cmets 具有非常好的指针，但作为示例，您可以创建一个对象，该对象将允许您使用 [] 运算符直接写入 CUDA 数组（或执行您选择的任何其他操作）：

struct CudaVector {

    unsigned int get(unsigned int index) {
        cout << "Get from device: " << index << endl;
        return 0; // TODO read actual value
    }
    void set(unsigned int index, unsigned int value) {
        cout << "Set in device: " << index << " " << value << endl;
        // TODO write actual value
    }

    struct Item {
        CudaVector& vector;
        unsigned int index;
        operator unsigned int() const {
            return vector.get(index);
        }       
        unsigned int operator=(unsigned int other) {
            vector.set(index, other);
            return other;
        }
        unsigned int operator=(const Item& other) {
            return (*this = static_cast<unsigned int>(other));
        }
    };

    Item operator[](unsigned int index) {
        return Item{*this, index};
    }
};

这就像：

CudaVector vector;
unsigned int foo = vector[8];
vector[5] = vector[6] = vector[7];

输出：

从设备获取：8
从设备获取：7
设置在设备中：6 0
在设备中设置：5 0

想法是您的operator[] 不返回引用，而是返回能够处理“读取”（使用转换运算符）和“写入”（使用赋值运算符）的临时对象。

（第二个重载允许链式分配，因为如果您不首先从unsigned int 分配，则不会自动拾取第一个。）

【讨论】：

这真的很有帮助。非常感谢先生。