Начало работы с общей памятью в PyCUDA

Я пытаюсь понять общую память, играя со следующим кодом:

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))

Я не вижу в этом ничего явно плохого, но я продолжаю получать ошибки синхронизации, и он не запускается.

Любая помощь очень ценится.


python cuda pycuda gpgpu pyopencl
person reckoner    schedule 07.06.2015    source источник
comment
Вы не указываете размер общей памяти для запуска ядра.   -  person talonmies    schedule 07.06.2015
comment
Оказывается, внешний shared float sdata[]; это то, как это нужно компилятору nvcc.   -  person reckoner    schedule 08.06.2015
comment
Да, но когда ядро ​​использует динамически выделяемую разделяемую память, вы должны передать размер выделения разделяемой памяти в байтах в качестве аргумента запуска ядра. Код, который вы разместили, этого не делает.   -  person talonmies    schedule 08.06.2015

Ответы (1)


arrow_upward
6
arrow_downward

Когда вы указываете

extern __shared__ float sdata[];

вы сообщаете компилятору, что вызывающая сторона предоставит разделяемую память. В PyCUDA это делается путем указания shared=nnnn в строке, которая вызывает функцию CUDA. В вашем случае что-то вроде:

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

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

__shared__ float sdata[400];
person ime    schedule 09.06.2015
comment
На самом деле мне пришлось отказаться от ключевого слова extern и использовать аргумент shared=, чтобы заставить его работать. - person reckoner; 10.06.2015