nVidia Thrust: 设备指针的const正确性

Question

在我的项目中，我广泛使用nVidia CUDA技术，有时候也会使用Thrust库来处理某些特定的任务。其中一个在这个库里实现得非常好的算法就是“Reduce”。其中一种使用情况是通过将向量中的每个非负元素除以所有元素的总和来对向量进行规范化操作。

template <typename T>
void normalise(T const* const d_input, const unsigned int size, T* d_output)
{
    const thrust::device_ptr<T> X = thrust::device_pointer_cast(const_cast<T*>(d_input));
    T sum = thrust::reduce(X, X + size);

    thrust::constant_iterator<T> denominator(sum);
    thrust::device_ptr<T> Y = thrust::device_pointer_cast(d_output);
    thrust::transform(X, X + size, denominator, Y, thrust::divides<T>());
}

(T通常是float或double)

总的来说，我不想在整个代码库中完全依赖于Thrust，因此我尽量确保像上面示例中的函数仅接受原始的CUDA设备指针。这意味着一旦它们被NVCC编译，我可以将它们静态链接到其他代码中而无需NVCC。

然而，这段代码让我担心。我希望该函数是const正确的，但我似乎找不到thrust::device_pointer_cast(...)的const版本 - 是否存在这样的东西？在这个代码版本中，我采用了const_cast，以便在函数签名中使用const，这让我感到难过。

顺便说一下，将reduce的结果复制到主机，然后再将其发送回设备进行下一步操作感觉有些奇怪。有更好的方法吗？

- Xharlie

---

你认为为什么要使用 non const？你为什么要进行类型转换？ - Ivan Aksamentov - Drop

我进行类型转换是因为没有进行类型转换会导致编译错误。"no instance of constructor "thrust::pointer<Element, Tag, Reference, Derived>::pointer [with Element=float, Tag=thrust::device_system_tag, Reference=thrust::device_reference<float>, Derived=thrust::device_ptr<float>]" matches the argument list"（你提供的链接是复制构造函数）。 - Xharlie

@Xharlie，你尝试过使用带有const修饰的设备指针吗？即thrust::device_ptr<const T>？ - Angew is no longer proud of SO

@Angew - 这个可以。thrust::device_ptr<const T> X = thrust::device_pointer_cast<const T>(d_input) 很好。 - Xharlie

@Angew - 感谢你的提示。 - Xharlie

Answer 1

如果你想要常量正确性(const-correctness)，那么你需要在任何地方都保持常量正确性。 input 是一个指向 const T 的指针，因此 X 也应该是这样的常量指针。

const thrust::device_ptr<const T> X = thrust::device_pointer_cast(d_input);

- Angew is no longer proud of SO

---

@Drop 是的（尽管这样写会是 const auto X = ...，以完全匹配原帖中的代码）。 - Angew is no longer proud of SO

如果@Angew在他的答案中使用了auto，我将从这个问题中学到很少的知识。他给出完整类型说明的事实才使它有用。 - Xharlie

确保T与指针中包含的类型相匹配。例如，如果d_input是uint64并且您尝试：const thrust :: device_ptr <float> X = thrust :: device_pointer_cast（d_input）;您将获得该错误。 - Alexander Higgins