Numba CUDA 在运行时共享内存大小?
Numba CUDA shared memory size at runtime?
在 CUDA C 中,可以直接定义在运行时指定大小的共享内存。我如何使用 Numba/NumbaPro CUDA 执行此操作?
到目前为止我所做的只是导致消息 "Argument 'shape' must be a constant".
出现错误
编辑:澄清一下,我想要的是 C CUDA 中的以下内容(示例取自 here:
__global__ void dynamicReverse(int *d, int n)
{
extern __shared__ int s[];
// some work in the kernel with the shared memory
}
int main(void)
{
const int n = 64;
int a[n];
// run dynamic shared memory version
dynamicReverse<<<1,n,n*sizeof(int)>>>(a, n);
}
我找到了解决方案(通过非常有用的 Continuum Analytics 用户支持)。我们所做的是像通常那样定义共享内存,但将形状设置为 0。然后,要定义共享数组的大小,我们必须将其作为第四个参数(在流标识符之后)提供给内核。例如:
@cuda.autojit
def myKernel(a):
sm = cuda.shared.array(shape=0,dtype=numba.int32)
# do stuff
arga = np.arange(512)
grid = 1
block = 512
stream = 0
sm_size = arga.size * arga.dtype.itemsize
myKernel[grid,block,stream,sm_size](arga)
在 CUDA C 中,可以直接定义在运行时指定大小的共享内存。我如何使用 Numba/NumbaPro CUDA 执行此操作?
到目前为止我所做的只是导致消息 "Argument 'shape' must be a constant".
出现错误编辑:澄清一下,我想要的是 C CUDA 中的以下内容(示例取自 here:
__global__ void dynamicReverse(int *d, int n)
{
extern __shared__ int s[];
// some work in the kernel with the shared memory
}
int main(void)
{
const int n = 64;
int a[n];
// run dynamic shared memory version
dynamicReverse<<<1,n,n*sizeof(int)>>>(a, n);
}
我找到了解决方案(通过非常有用的 Continuum Analytics 用户支持)。我们所做的是像通常那样定义共享内存,但将形状设置为 0。然后,要定义共享数组的大小,我们必须将其作为第四个参数(在流标识符之后)提供给内核。例如:
@cuda.autojit
def myKernel(a):
sm = cuda.shared.array(shape=0,dtype=numba.int32)
# do stuff
arga = np.arange(512)
grid = 1
block = 512
stream = 0
sm_size = arga.size * arga.dtype.itemsize
myKernel[grid,block,stream,sm_size](arga)