我有以下简单的转置核（忽略输入矩阵初始化）：
#include <cuda_runtime.h>
const int rows = 8192;
const int columns = 8192;
const int tile_dim = 32;
__global__ void transpose(float* in, float* out)
{
__shared__ float tile[tile_dim][tile_dim + 1];
int x = blockIdx.x * tile_dim + threadIdx.x;
int y = blockIdx.y * tile_dim + threadIdx.y;
tile[threadIdx.y][threadIdx.x] = in[y * columns + x];
__syncthreads();
x = blockIdx.y * tile_dim + threadIdx.x;
y = blockIdx.x * tile_dim + threadIdx.y;
out[y * columns + x] = tile[threadIdx.x][threadIdx.y]; // Good
// out[y * columns + x] = ((float*)tile)[threadIdx.x * tile_dim + threadIdx.y]; // Bad
}
int main()
{
float *in, *out;
size_t size = rows * columns * sizeof(float);
cudaMalloc(&in, size);
cudaMalloc(&out, size);
dim3 grid(rows / tile_dim, columns / tile_dim);
dim3 block(tile_dim, tile_dim);
transpose<<<grid, block>>>(in, out);
cudaDeviceSynchronize();
return 0;
}标记为“Good”和“Bad”的行是导致问题的原因。切换Good -〉Bad会导致运行时速度降低近2倍，并在共享内存负载上显示更多的事务（68 M vs 200 K）。
这是怎么回事？
1条答案
按热度按时间iszxjhcz1#
内核使用the "trick" with shared memory将列填充1：
以允许无存储体冲突的行式和列式存取。
这行代码利用了这一点：
这并不：
要在“1D”情况下获得等效的性能/索引行为，您需要乘以共享内存分配的*实际宽度：
顺便说一句，我怀疑从这个转换：
（仅）：
会给予数值上正确的结果（需要在内核代码中进行其他更改），但这似乎不是问题的关键。