Начала работы с общей памятью на 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 это делается путем указания shared=nnnn на линии, вызывающей функцию CUDA. В вашем случае, что-то вроде:

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

В качестве альтернативы, вы можете отказаться от Экстерна ключевого слова, и указать общую память непосредственно:

__shared__ float sdata[400];