入門用下面的代碼打開始與共享內存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];