#include <stdio.h>
__global__ void staticReverse(int *d, int n)
{
__shared__ int s[64];
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
}
__global__ void dynamicReverse(int *d, int n)
{
extern __shared__ int s[];
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
}
int main(void)
{
const int n = 64;
int a[n], r[n], d[n];
for (int i = 0; i < n; i++) {
a[i] = i;
r[i] = n-i-1;
d[i] = 0;
}
int *d_d;
cudaMalloc(&d_d, n * sizeof(int));
// run version with static shared memory
cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
staticReverse<<<1,n>>>(d_d, n);
cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < n; i++)
if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
// run dynamic shared memory version
cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < n; i++)
if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
}
__global__ void dynamicReverse(int *d, int n)
{
extern __shared__ int s[]; //**** extern 获取核函数分配的动态内存
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
}
Bank Conflict
当然，shared memory 也不是完全完美的，其还有可能引发一种新的冲突 ———— bank conflict: bank conflict 发生在 同一个warp中不同的线程访问同一个bank不同的数据 上面四个条件缺一不可造成 bank conflict： 1. 同一个 warp，不同的 warp 一般不考虑 bank conflict，因为 warp 的分发是随机的我们是不可知的。 2. 不同的 thread，同一个 thread 一个 clock cycle 只能存取一个 data 所以这是一句废话，但是要保证在同一个 warp 里面所以不是一句废话 3. 同一个 bank，对于不同的 bank 我们管不着 4. 不同的 data，对于同一个 bank 中，如果 不同的 thread 都访存的是一个 data，cuda 提供了广播的方法，类似于 python 中的 numpy 广播，反而是对性能的提升而不是下降，只有访存不同的 data 的时候才会出现 bank conflict
Figure Explanation
在 Left 和 Right 中，每个 thread 都访问的是不同的 bank，所以不满足 bank conflict 的第三个条件，没有发生 bank conflict。 在 Middle 中，相当于是 stride=16 的访问导致了偶数编号的 bank 同时被两个线程所访问，而且访问的是不同的 data，所以出现了 two-way bank conflict。
Example
举矩阵乘法的例子：
$C=A*B$
global memory
__global__
void simpleMultiply(float *a, float* b, float *c, int N)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int i = 0; i < TILE_DIM; i++) {
sum += a[row*TILE_DIM+i] * b[i*N+col];
}
c[row*N+col] = sum;
}
partial shared memory
__global__
void coalescedMultiply(float *a, float* b, float *c, int N)
{
__shared__ float aTile[TILE_DIM][TILE_DIM];
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
__syncwarp(); //***
for (int i = 0; i < TILE_DIM; i++) {
sum += aTile[threadIdx.y][i]* b[i*N+col];
}
c[row*N+col] = sum;
}
__global__
void simpleMultiply(float *a, float *c, int M)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int i = 0; i < TILE_DIM; i++) {
sum += a[row*TILE_DIM+i] * a[col*TILE_DIM+i];
}
c[row*M+col] = sum;
}
all shared memory
__global__
void coalescedMultiply(float *a, float *c, int M)
{
__shared__ float aTile[TILE_DIM][TILE_DIM];
__shared__ float transposedTile[TILE_DIM][TILE_DIM];
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
transposedTile[threadIdx.x][threadIdx.y] =
a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
threadIdx.x];
__syncthreads();
for (int i = 0; i < TILE_DIM; i++) {
sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
}
c[row*M+col] = sum;
}
但是测试发现性能并没有预期的提升，是因为出现了 bank conflict： 分析 sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x] 的部分，每个 thread 自己在 transposedTile 做 TILE_DIM-stride 遍历，不同的 thread 不会碰在一起，所以是前面出了问题。 在拷贝的时候，同一个 warp 中的 threadIdx.y 是相同的，而在TransposedTile 拷贝的过程中出现了同一个 warp 中所有的 thread 都在同一列上复制，而我们的 TransposedTile 的 col=TILE_DIM=32，刚好就是每个 bank 管一个 col，那所有的 col 都在访问同一个 bank 的不同数据，造成了最严重的 32-way bank conflict !!! 所以性能必然下降，解决方法也很简单。
__host__
cudaError_t cudaFuncSetCacheConfig (const void *func, cudaFuncCache cacheConfig)
Effect: Sets the preferred cache configuration for **a device function**.
Params:
1. const void *func: 必须是由 __global__ 声明的核函数
2. cudaFuncCache cacheConfig:
1. cudaFuncCachePreferNone: no preference for shared memory or L1 (default)
2. cudaFuncCachePreferShared: prefer larger shared memory and smaller L1 cache
3. cudaFuncCachePreferL1: prefer larger L1 cache and smaller shared memory
4. cudaFuncCachePreferEqual: prefer equal size L1 cache and shared memory
__host__
cudaError_t cudaDeviceSetCacheConfig (cudaFuncCache cacheConfig)
Effect: Sets the preferred cache configuration for the **current device**.
上面的是针对某个指定的核函数 func 进行cache的分配
这个就是直接对当前的 Device 直接设置，只要在 Device 上面运行的核函数都有 cache 的 Prefer
注意： 官方文档也说了：
This is only a preference. The runtime will use the requested configuration if possible, but it is free to choose a different configuration if required to execute func.
上面的三种访存，左侧数字对应的同一个 warp 中 thread 的编号（在 CC 2.0 以上每个 warp 中都是 32个线程），右侧的黄色是 bank 的编号（在 CC 2.0 以上shared memory 中的 bank 的个数也是32），注意黄色的 bank 内部的橙色小矩形代表的是 data。
上面的左中右都没有出现 bank conflict，因为即使有两个线程访问同一个 bank，他们内部访问的都是同一个 data，不满足 bank conflict 的第四个条件，所以都没有发生 bank conflict。
