我试图在 GPU 上实现 FDTD 方程。我最初实现了使用全局内存的内核。内存合并不是那么好。因此,我实现了另一个使用共享内存来加载值的内核。我正在研究1024x1024
.
代码如下
__global__ void update_Hx(float *Hx, float *Ez, float *coef1, float* coef2){
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;
__shared__ float Ez_shared[BLOCKSIZE_HX][BLOCKSIZE_HY + 1];
/*int top = offset + x_index_dim;*/
if(threadIdx.y == (blockDim.y - 1)){
Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
Ez_shared[threadIdx.x][threadIdx.y + 1] = Ez[offset + x_index_dim];
}
else{
Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
}
}
常数BLOCKSIZE_HX
=16
和BLOCKSIZE_HY
= 16
。
当我运行视觉分析器时,它仍然说内存没有合并。
编辑:我使用 GT 520 显卡,cuda 计算能力为 2.1。我的全局 L2 事务/访问 =7.5
即有245 760
用于
32768
执行该行的
L2 事务Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
Global memory load efficiency
是50%
。
Global memory load efficiency
=100 * gld_requested_throughput/ gld_throughput
尽管我的线程正在查看 16 个连续值,但我无法弄清楚为什么会有如此多的内存访问。有人可以指出我做错了什么吗?
编辑:感谢所有帮助。