关于CUDA中的const指针和参数传递

About const pointer and parameter-passing in CUDA

本文关键字：参数传递 指针 const CUDA 中的 关于      更新时间：2023-10-16

nvcc如何处理内核中的const指针?

根据nvidia的说法，在参数传递期间为指针添加const和restrict使NVCC能够进行积极的优化，这是否严格遵循C/c++的方式?

假设A是指向数据缓冲区的指针，该缓冲区可能会被其他线程/流频繁更新，但在此测试内核调用期间内容不会被修改:

test<<<blocks, threads>>>(const int *__restrict__ A, int *__restrict__ B);

那么NVCC是否可以保持这种正确性:在每次内核调用时加载A中的更新数据，而不是加载一些预缓存的过时数据?

const像c++一样工作。const变量不能更改，编译器会在编译时检查这一点。编译器只检查给定范围的const正确性，因为const可以通过C风格强制转换改变。

restrict以C方式工作。将指针标记为restrict时，编译器假定这些指针没有别名。这是你给定的事实，编译器不会检查这个事实是否为真。

回答你的问题，NVCC不会确保内核启动之间全局内存写和读的正确性。由于内核启动在CUDA中是异步的，您必须确保修改这些内存空间的内核不会同时执行。您可以通过同步内存拷贝和/或cudaDeviceSynchronize()来实现这一点。如果同时启动这些内核，则无法确保不同内核的所有更改都在其他内核访问之前提交到全局内存。

相关文章：