+关注
已关注

分类  

暂无分类

标签  

暂无标签

日期归档  

暂无数据

在PyCUDA上共享内存入门

发布于2020-11-19 05:45     阅读(1018)     评论(0)     点赞(14)     收藏(4)


0

1

2

3

4

5

6

我正在尝试通过玩以下代码来了解共享内存:

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

我看不到任何明显与此有关的错误,但是我不断收到同步错误,并且它没有运行。

任何帮助,不胜感激。


解决方案


当您指定

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];

0

1

2

3

4

5

6

7



所属网站分类: 技术文章 > 问答

作者:黑洞官方问答小能手

链接: https://www.pythonheidong.com/blog/article/619624/beb69accf0f86250f604/

来源: python黑洞网

任何形式的转载都请注明出处,如有侵权 一经发现 必将追究其法律责任

14 0
收藏该文
已收藏

评论内容:(最多支持255个字符)