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

Question

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

根据nvidia的说法,在参数传递过程中为指针添加const和restrict会使NVCC进行积极优化,这是否严格遵循C/C++方式？

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

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

然后NVCC可以保持这种正确性:在每次内核调用时加载A中的更新数据,而不是加载一些预先缓存的过时数据？

Answer 1

const像C++一样工作.甲const变量不能改变,并且这是在由编译器编译时进行检查.编译器仅针对给定范围检查const-correctness,因为可以使用C样式转换更改const.

restrict以C方式工作.当您将指针标记为限制编译器时,假定这些指针没有别名.这是你的事实,编译器不会检查这个事实是否属实.

提出您的问题,NVCC将无法确保全局内存写入和内核启动之间读取的正确性.由于内核启动在CUDA中是异步的,因此必须确保修改这些内存空间的内核不会同时执行.您可以通过同步内存副本和/或实现此目的cudaDeviceSynchronize().如果同时启动这些内核,则在从其他内核访问之前,无法确保不同内核的所有更改都提交到全局内存.