虽然可以在设备代码中使用 Package 器，如thrust::device_ptr和thrust::device_reference，但它们并不设计为将指针作为值类型保存。解引用device_ptr给出device_reference，其具有实现的典型值类型的许多运算符，但不具有指针类型的运算符，如进一步解引用，->或通常使用成员方法（在下面的示例中为.get()）。因此，必须强制（但安全地）将引用强制转换为它的底层类型，即static_cast。我不知道为什么device_reference::operator value_type (void)没有出现在documentation中，但是从device_reference<T>到T的转换/转换肯定是一个预期的特性。这创建了一个副本，而不是一个引用T&，这对T = device_ptr<...>来说很好，因为指针是轻量级的。
因此，虽然可以使用device_vector<device_ptr<T>>（见下面的示例），但这可能不是一个好主意。这些 Package 器的主要目的是避免在Thrust算法中使用显式执行策略。当用于设备代码时，无法向主机调度，因此设备 Package 器最多只能携带不必要的信息。
使用嵌套device_ptr的示例（即使它确实有效，也要避免这种情况）：

#include <thrust/device_vector.h>
#include <thrust/for_each.h>
#include <thrust/device_ptr.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/execution_policy.h>

class Foo {
    int i{};

    public:
    Foo() = default;

    __host__ __device__
    Foo(int j) noexcept : i{j} {}

    __host__ __device__
    int get() const noexcept { return i; }
};

int main() {
    thrust::device_vector<Foo> foo(1, 42);
    thrust::device_vector<thrust::device_ptr<Foo>> bar(1, nullptr);
    
    auto foo_ptr = foo.data();     // gives device_ptr<Foo>
    auto foo_ptr_ptr = bar.data(); // gives device_ptr<device_ptr<Foo>>

    thrust::for_each(
        thrust::device, // this is optional here as device is the default when the given iterators allow for it, but I like to be explicit about it
        thrust::make_counting_iterator(0),
        thrust::make_counting_iterator(1),
        [foo_ptr, foo_ptr_ptr] __host__ __device__ (int idx) {
            foo_ptr_ptr[idx] = &foo_ptr[idx]; // device_reference<T>::operator=(T&) (with T = device_ptr<Foo>) is implemented
            printf(
                "%d\n", 
                static_cast<thrust::device_ptr<Foo>>(
                    foo_ptr_ptr[idx] // gives device_reference<device_ptr<Foo>>
                )->get() // device_ptr<T>::operator->() is implemented
            );
        });
    
    return 0;
}

如果要使用*解引用，则需要两次强制转换，即

static_cast<Foo>(*static_cast<device_ptr<Foo>>(foo_ptr_ptr[idx])).get()

因为X1 M15 N1 X不具有X1 M16 N1 X成员函数。
不使用 Package （这样做是为了提高可读性）：

// ...same includes and Foo class as in above snippet...

int main() {
    thrust::device_vector<Foo> foo(1, 42);
    thrust::device_vector<Foo*> bar(1, nullptr);
    
    auto foo_ptr = thrust::raw_pointer_cast(foo.data());     // gives Foo*
    auto foo_ptr_ptr = thrust::raw_pointer_cast(bar.data()); // gives Foo**

    thrust::for_each(
        thrust::device, // this is optional here as device is the default when the given iterators allow for it, but I like to be explicit about it
        thrust::make_counting_iterator(0),
        thrust::make_counting_iterator(1),
        [foo_ptr, foo_ptr_ptr] __host__ __device__ (int idx) {
            foo_ptr_ptr[idx] = &foo_ptr[idx];
            printf("%d\n", foo_ptr_ptr[idx]->get());
        });
    
    return 0;
}

根据迭代器和静态调度，您应该为此使用 raw_pointer_cast：

hittable *pHittable = thrust::raw_pointer_cast(object);
if (pHittable->hit(r, t_min, t_max, temp_rec)) {