共享内存的分配以及清零操作

加速计算 CUDA
1
  1. 如果我需要在共享内存中开辟两个区域__global__ void ker()
    {
    shared float data[SIZE_DATA];
    shared igned int index[SIZE_INDEX];

}

可以使用动态的方式将根据程序的需要动态变化SIZE_DATA和SIZE_INDEX么？
主要是两种数据类型不同，不能一起声明。
2.如何在程序中清零共享内存，每个线程清0一部分共享内存是最合理高效的方式么？或者是否可以使用Memset类似的函数清0？

谢谢！

2

楼主您好，

（1）CUDA C目前还不支持VLA(variable length array), 请慢慢等待它支持此C99特性。（可能下一版就支持了，也可能需要下下个）

（2）最高效的方式是大家一起写入0. 在启用了8B/bank宽度的3.x卡上，最有效的方式是每人写入8B的整数0（或者double的0.0也行，0.0和整数0的编码格式一样）。在其他卡上，最有效的方式是写入4B的0.
使用memset()在设备端非常低效，不建议使用。

感谢您的来访。

3

谢版主，
对于第一个问题，我声明时的将SIZE_DATA，SIZE_INDEX做为宏定义是否就可以了？
另外由于我要使用别的线程的共享内存，是需要加上volatile么？
以及，对共享内存使用 extern 是什么含义？

对于第二个问题，写入0的操作就是执行for循环语句，将数组的每个元素 赋值为0么？

问题有点多，先谢版主，另外最近cuda论坛有时候无法访问，是什么原因引起的。

4

LZ您好：

第一个问题：
1：使用宏的问题不属于CUDA的问题，请您自行思考。
2：shared memory本身就是一个block的线程共享的，而不是某个线程私有的，所以您的说法其前提不成立。
3：这表示声明的shared memory的大小在调用kernel的时候由<<<>>>中的参数指定，详情请参考programming guide。

第二个问题：
请参照2#横扫斑竹的叙述，选择最有效的方式清零。

最近论坛因为某些不可抗力因素，运作不太正常，已经积极联系维护团队处理了，今天的情况已经有所改善。

祝您好运~