const int buffer_size = 8192; // assume an 8k buffer
float *device_buffer = ; // assume you have a buffer already on the device with the data you want.
my_kernel<<<num_blocks, num_threads, buffer_size>>>(..., buffer_size, device_buffer);
__global__ void my_kernel(..., int buffer_size, const float *device_buffer) {
extern __shared__ float shmem_buffer[];
for (int idx = threadIdx.x; idx < buffer_sze; idx += blockDim.x) {
shmem_buffer[idx] = device_buffer[idx];
}
__syncthreads();
// rest of your kernel goes here. You can access data in shmem_buffer;
}
2条答案
按热度按时间webghufk1#
不可能的填充共享内存的唯一方法是使用CUDA内核中的线程。
如果您希望内核在启动时可以使用一组(只读)数据,那么当然可以使用
__constant__
内存。这样的存储器可以使用文档中指示的API在主机代码上/由主机代码设置,即cudaMemcpyToSymbol
。__constant__
内存实际上只有在每个线程都将在给定的访问周期中访问 * 相同的位置 * 时才有用,例如否则,使用普通的全局内存,静态或动态分配,使用适当的主机API进行初始化(动态:
cudaMemcpy
,静态:cudaMemcpyToSymbol
)。bf1o4zei2#
虽然您请求的特定行为无法自动实现,但这实际上是一个相当常见的CUDA范例:
首先,让所有线程将表复制到shmem中。
同步线程
访问内核中的数据。
如果您对数据进行相当随机的访问,并且您希望平均对每个条目进行多次访问,那么这将大大提高性能。实际上,您使用shmem作为托管缓存,并将来自DRAM的负载聚合到shmem一次,以便多次使用。此外,shmem对于未合并的负载没有惩罚。
例如,您可以编写如下代码:
换句话说,您只需显式地对副本进行编码。由于来自DRAM的所有负载将被完美地合并,因此这应该接近最佳效率。