在NVIDIA的2.x架构中,每个warp都有64kb的内存,默认情况下分为48kb的共享内存和16kb的L1缓存(服务global和constant内存).
我们都知道访问共享内存的银行冲突 - 内存分为32个大小为32位的存储区,允许所有32个线程同时独立访问.另一方面,全局内存虽然慢得多,但不会遇到银行冲突,因为内存请求在整个warp中合并.
问题: 假设来自全局或常量内存的一些数据缓存在L1缓存中以用于给定的warp.访问此数据是否存在银行冲突,例如共享内存(因为L1缓存和共享内存实际上是相同的硬件),还是以全局/常量内存的方式无冲突？
我是学习CUDA并行编程的新手.现在我对设备的全局内存访问感到困惑.这是关于扭曲模型和合并.
有一些要点:
据说一个块中的线程被分成了warp.在每个warp中最多有32个线程.这意味着同一warp的所有这些线程将与同一处理器同时执行.那么半翘曲的感觉是什么？
当谈到一个块的共享内存时,它将被分成16个库.为了避免银行冲突,多个线程可以同时读取一个银行而不是写入同一个银行.这是正确的解释吗？
提前致谢!
我有以下代码，它使用共享内存执行平铺矩阵转置以提高性能。共享内存用 1 列填充，以避免 32x32 线程块的存储体冲突。
__global__ void transpose_tiled_padded(float *A, float *B, int n)
{
int i_in = blockDim.x*blockIdx.x + threadIdx.x;
int j_in = blockDim.y*blockIdx.y + threadIdx.y;
int i_out = blockDim.x*blockIdx.y + threadIdx.x;
int j_out = blockDim.y*blockIdx.x + threadIdx.y;
extern __shared__ float tile[];
// coalesced read of A rows to (padded) shared tile column (transpose)
tile[threadIdx.y + threadIdx.x*(blockDim.y+1)] = A[i_in + j_in*n];
__syncthreads();
// coalesced write from (padded) shared tile column to B rows
B[i_out + j_out*n] = tile[threadIdx.x + threadIdx.y*(blockDim.x+1)]; …我正在编写一些 N 体模拟代码，在 CUDA 中针对 Volta 和图灵系列卡进行短程交互。我计划使用共享内存，但我不太清楚这样做时如何避免银行冲突。由于我的交互是本地的，我计划将我的粒子数据分类到本地组中，我可以将这些数据发送到每个 SM 的共享内存（还没有担心粒子的邻居正在从另一个 SM 工作。为了变得更好性能（避免库冲突），仅每个线程从/向共享内存的不同地址读取/写入就足够了，但每个线程可以无序访问该内存而不会受到惩罚？
我看到的所有信息似乎只提到内存被合并以从全局内存到共享内存的复制，但我没有看到任何关于扭曲（或整个 SM）中的线程是否关心共享内存中的合并。