CUDA Thrust库中遍历device_vector<device_ptr>时的解引用错误

thrust::device_vector是Thrust库中管理设备端内存的容器，它的迭代器在设备上解引用时，返回的不是直接的元素值，而是thrust::device_reference<T>类型的代理对象。这个代理对象作为中间层，负责与设备内存交互——因为设备内存无法被主机端直接访问，但在设备端代码中，它用来模拟对设备内存元素的引用操作。

当你用范围for循环for (const auto object: objects)遍历device_vector<device_ptr<hittable>>时，循环变量object的实际类型是device_reference<const device_ptr<hittable>>，而非你期望的device_ptr<hittable>。这就是直接调用object->hit()报错的原因：device_reference不是指针类型，不支持->运算符。

针对这个问题，有几种可行的解决方式：

1. 显式转换为device_ptr #

利用官方文档提到的特性：对device_reference取地址会得到指向对应元素的device_ptr，解引用该device_ptr即可获取实际的device_ptr<hittable>对象：

__device__ bool hittable_list::hit(const ray& r, float t_min, float t_max, hit_record& rec) const {
    hit_record temp_rec;
    bool hit_anything = false;
    float total_prob = 1.0;

    for (const auto& ref : objects) {
        // 取device_reference的地址得到device_ptr<const device_ptr<hittable>>，再解引用得到目标对象
        const thrust::device_ptr<hittable> object = *(&ref);
        if (object->hit(r, t_min, t_max, temp_rec)) {
            hit_anything = true;
            total_prob *= temp_rec.trans_prob;
            temp_rec.trans_prob = total_prob;
            rec = temp_rec;
        }
    }
    return hit_anything;
}

2. 直接赋值转换 #

device_reference支持隐式转换为它所引用的元素类型，因此可以直接将循环中的代理对象赋值给device_ptr变量：

__device__ bool hittable_list::hit(const ray& r, float t_min, float t_max, hit_record& rec) const {
    hit_record temp_rec;
    bool hit_anything = false;
    float total_prob = 1.0;

    for (const auto& ref : objects) {
        thrust::device_ptr<hittable> object = ref;
        if (object->hit(r, t_min, t_max, temp_rec)) {
            hit_anything = true;
            total_prob *= temp_rec.trans_prob;
            temp_rec.trans_prob = total_prob;
            rec = temp_rec;
        }
    }
    return hit_anything;
}

3. 使用索引循环替代范围for #

如果你更习惯传统遍历方式，也可以通过device_vector的operator[]访问元素后转换为device_ptr：

__device__ bool hittable_list::hit(const ray& r, float t_min, float t_max, hit_record& rec) const {
    hit_record temp_rec;
    bool hit_anything = false;
    float total_prob = 1.0;

    for (size_t i = 0; i < objects.size(); ++i) {
        thrust::device_ptr<hittable> object = objects[i];
        if (object->hit(r, t_min, t_max, temp_rec)) {
            hit_anything = true;
            total_prob *= temp_rec.trans_prob;
            temp_rec.trans_prob = total_prob;
            rec = temp_rec;
        }
    }
    return hit_anything;
}

内容的提问来源于stack exchange，提问作者John Meneghini