PyCUDA 上的共享内存入门

Question

我正在尝试通过使用以下代码来理解共享内存：

import pycuda.driver as drv
import pycuda.tools
import pycuda.autoinit
import numpy
from pycuda.compiler import SourceModule

src='''
__global__ void reduce0(float *g_idata, float *g_odata) {
extern __shared__ float sdata[];
// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

sdata[tid] = g_idata[i];
__syncthreads();
// do reduction in shared mem
for(unsigned int s=1; s < blockDim.x; s *= 2) {
   if (tid % (2*s) == 0) {
      sdata[tid] += sdata[tid + s];
   }
__syncthreads();
}
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
'''

mod = SourceModule(src)
reduce0=mod.get_function('reduce0')

a = numpy.random.randn(400).astype(numpy.float32)

dest = numpy.zeros_like(a)
reduce0(drv.In(a),drv.Out(dest),block=(400,1,1))

我看不出这有什么明显的问题，但我不断收到同步错误并且它没有运行。

非常感谢任何帮助。

Answer 1

当你指定时

extern __shared__ float sdata[];

您是在告诉编译器调用者将提供共享内存。在 PyCUDA 中，这是通过在调用 CUDA 函数的行上指定 shared=nnnn 来完成的。在您的情况下，类似于：

reduce0(drv.In(a),drv.Out(dest),block=(400,1,1),shared=4*400)

或者，您可以删除 extern 关键字，并直接指定共享内存：

__shared__ float sdata[400];