CUDA优化实例（三）共享内存
CUDA优化实例（三）共享内存
经过前面的实验发现，共享内存是优化CUDA程序的核心方法。共享内存可以通过对全局内存数据进行合并访问，让kernel内交错的内存需求去访问共享内存。如：矩阵转置问题，将二维内存的行写入二维内存的列。对列的写入就是一个内存交错访问的例子。可以用合并的方式将块要操作的数据写入共享内存，让复杂的内存交错访问访问共享内存，然后将结果以合并的方式写入全局内存。
在使用共享内存时需要注意
共享内存是服务于块的，所以一般要注意同步即__syncthreads() 的使用。
共享内存对复杂的内存处理有优势是因为它被分为32个内存存储体，而32个内存存储体可同时被访问，即内存请求分布在不同的内存存储体上的话，会同时获得该数据，但请求在同一（行）内存存储体上的话，速度会受影响。填充是一种克服内存存储体冲突的方法，下面会介绍。
每个共享内存存储体的宽可以是32位或64位，即4字节和8字节，用于对付单精度和双精度的数据。一般情况下，我们认为我们的每个kernel只处理一个数据只处理4字节的单精度或8字节的双精度，这样一个线程束的32个线程纠正好与32个（32个说成，一行更准确）内存存储体的大小一样都是128b，这样如果不存在内存体冲突，只需要一个共享内存访问事务，反之则需要多个。共享内存资源处理单精度的数据默认存储体宽度的32位即可，若处理双精度，为了不浪费稀有的共享内存资源，显式的定义宽度为64位。
共享内存是稀缺的SM上的资源，所以它会影响块在SM上的分布，影响占用率，计算一下每个块共享的低下限。对于存储体32位情况：我的GTX1050Ti 每个SM的共享内存大小是48KB，如果保证占用率的话，每个SM最多可容纳16个块，也就是每个块可以最多分配48KB/16 = 3KB的空间，换算到768个单精度的浮点数据。也就是说一个块最多可以申请768个4字节的数据，如果超过这个空间的大小，那么不能保证所有SM中有16个块，不能保证占用率。（提一下寄存器资源2048个线程运行时，每个线程32字节）
此实验是在NVIDIA官方找的一个 。
#include &stdio.h&
#include &assert.h&
cudaError_t checkCuda(cudaError_t result)
#if defined(DEBUG) || defined(_DEBUG)
if (result != cudaSuccess) {
fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
assert(result == cudaSuccess);
const int TILE_DIM = 32;
const int BLOCK_ROWS = 8;
const int NUM_REPS = 100;
void postprocess(const float *ref, const float *res, int n, float ms)
bool passed = true;
for (int i = 0; i & i++)
if (res[i] != ref[i]) {
printf("%d %f %f\n", i, res[i], ref[i]);
printf("%25s\n", "*** FAILED ***");
passed = false;
if (passed)
printf("%20.2f\n", 2 * n * sizeof(float) * 1e-6 * NUM_REPS / ms );
__global__ void copy(float *odata, const float *idata)
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j & TILE_DIM; j+= BLOCK_ROWS)
odata[(y+j)*width + x] = idata[(y+j)*width + x];
__global__ void copySharedMem(float *odata, const float *idata)
__shared__ float tile[TILE_DIM * TILE_DIM];
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j & TILE_DIM; j += BLOCK_ROWS)
tile[(threadIdx.y+j)*TILE_DIM + threadIdx.x] = idata[(y+j)*width + x];
__syncthreads();
for (int j = 0; j & TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x] = tile[(threadIdx.y+j)*TILE_DIM + threadIdx.x];
__global__ void transposeNaive(float *odata, const float *idata)
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j & TILE_DIM; j+= BLOCK_ROWS)
odata[x*width + (y+j)] = idata[(y+j)*width + x];
__global__ void transposeCoalesced(float *odata, const float *idata)
__shared__ float tile[TILE_DIM][TILE_DIM];
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j & TILE_DIM; j += BLOCK_ROWS)
tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
__syncthreads();
x = blockIdx.y * TILE_DIM + threadIdx.x;
y = blockIdx.x * TILE_DIM + threadIdx.y;
for (int j = 0; j & TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
__global__ void transposeNoBankConflicts(float *odata, const float *idata)
__shared__ float tile[TILE_DIM][TILE_DIM+1];
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j & TILE_DIM; j += BLOCK_ROWS)
tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
__syncthreads();
x = blockIdx.y * TILE_DIM + threadIdx.x;
y = blockIdx.x * TILE_DIM + threadIdx.y;
for (int j = 0; j & TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
int main(int argc, char **argv)
const int nx = 1024;
const int ny = 1024;
const int mem_size = nx*ny*sizeof(float);
dim3 dimGrid(nx/TILE_DIM, ny/TILE_DIM, 1);
dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1);
int devId = 0;
if (argc & 1) devId = atoi(argv[1]);
cudaDeviceP
checkCuda( cudaGetDeviceProperties(&prop, devId));
printf("\nDevice : %s\n", prop.name);
printf("Matrix size: %d %d, Block size: %d %d, Tile size: %d %d\n",
nx, ny, TILE_DIM, BLOCK_ROWS, TILE_DIM, TILE_DIM);
printf("dimGrid: %d %d %d. dimBlock: %d %d %d\n",
dimGrid.x, dimGrid.y, dimGrid.z, dimBlock.x, dimBlock.y, dimBlock.z);
checkCuda( cudaSetDevice(devId) );
float *h_idata = (float*)malloc(mem_size);
float *h_cdata = (float*)malloc(mem_size);
float *h_tdata = (float*)malloc(mem_size);
float *gold
= (float*)malloc(mem_size);
float *d_idata, *d_cdata, *d_
checkCuda( cudaMalloc(&d_idata, mem_size) );
checkCuda( cudaMalloc(&d_cdata, mem_size) );
checkCuda( cudaMalloc(&d_tdata, mem_size) );
if (nx % TILE_DIM || ny % TILE_DIM) {
printf("nx and ny must be a multiple of TILE_DIM\n");
goto error_
if (TILE_DIM % BLOCK_ROWS) {
printf("TILE_DIM must be a multiple of BLOCK_ROWS\n");
goto error_
for (int j = 0; j & j++)
for (int i = 0; i & i++)
h_idata[j*nx + i] = j*nx +
for (int j = 0; j & j++)
for (int i = 0; i & i++)
gold[j*nx + i] = h_idata[i*nx + j];
checkCuda( cudaMemcpy(d_idata, h_idata, mem_size, cudaMemcpyHostToDevice) );
cudaEvent_t startEvent, stopE
checkCuda( cudaEventCreate(&startEvent) );
checkCuda( cudaEventCreate(&stopEvent) );
printf("%25s%25s\n", "Routine", "Bandwidth (GB/s)");
printf("%25s", "copy");
checkCuda( cudaMemset(d_cdata, 0, mem_size) );
copy&&&dimGrid, dimBlock&&&(d_cdata, d_idata);
checkCuda( cudaEventRecord(startEvent, 0) );
for (int i = 0; i & NUM_REPS; i++)
copy&&&dimGrid, dimBlock&&&(d_cdata, d_idata);
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
checkCuda( cudaMemcpy(h_cdata, d_cdata, mem_size, cudaMemcpyDeviceToHost) );
postprocess(h_idata, h_cdata, nx*ny, ms);
printf("%25s", "shared memory copy");
checkCuda( cudaMemset(d_cdata, 0, mem_size) );
copySharedMem&&&dimGrid, dimBlock&&&(d_cdata, d_idata);
checkCuda( cudaEventRecord(startEvent, 0) );
for (int i = 0; i & NUM_REPS; i++)
copySharedMem&&&dimGrid, dimBlock&&&(d_cdata, d_idata);
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
checkCuda( cudaMemcpy(h_cdata, d_cdata, mem_size, cudaMemcpyDeviceToHost) );
postprocess(h_idata, h_cdata, nx * ny, ms);
printf("%25s", "naive transpose");
checkCuda( cudaMemset(d_tdata, 0, mem_size) );
transposeNaive&&&dimGrid, dimBlock&&&(d_tdata, d_idata);
checkCuda( cudaEventRecord(startEvent, 0) );
for (int i = 0; i & NUM_REPS; i++)
transposeNaive&&&dimGrid, dimBlock&&&(d_tdata, d_idata);
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
checkCuda( cudaMemcpy(h_tdata, d_tdata, mem_size, cudaMemcpyDeviceToHost) );
postprocess(gold, h_tdata, nx * ny, ms);
printf("%25s", "coalesced transpose");
checkCuda( cudaMemset(d_tdata, 0, mem_size) );
transposeCoalesced&&&dimGrid, dimBlock&&&(d_tdata, d_idata);
checkCuda( cudaEventRecord(startEvent, 0) );
for (int i = 0; i & NUM_REPS; i++)
transposeCoalesced&&&dimGrid, dimBlock&&&(d_tdata, d_idata);
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
checkCuda( cudaMemcpy(h_tdata, d_tdata, mem_size, cudaMemcpyDeviceToHost) );
postprocess(gold, h_tdata, nx * ny, ms);
printf("%25s", "conflict-free transpose");
checkCuda( cudaMemset(d_tdata, 0, mem_size) );
transposeNoBankConflicts&&&dimGrid, dimBlock&&&(d_tdata, d_idata);
checkCuda( cudaEventRecord(startEvent, 0) );
for (int i = 0; i & NUM_REPS; i++)
transposeNoBankConflicts&&&dimGrid, dimBlock&&&(d_tdata, d_idata);
checkCuda( cudaEventRecord(stopEvent, 0) );
checkCuda( cudaEventSynchronize(stopEvent) );
checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
checkCuda( cudaMemcpy(h_tdata, d_tdata, mem_size, cudaMemcpyDeviceToHost) );
postprocess(gold, h_tdata, nx * ny, ms);
error_exit:
checkCuda( cudaEventDestroy(startEvent) );
checkCuda( cudaEventDestroy(stopEvent) );
checkCuda( cudaFree(d_tdata) );
checkCuda( cudaFree(d_cdata) );
checkCuda( cudaFree(d_idata) );
free(h_idata);
free(h_tdata);
free(h_cdata);
free(gold);
性能由好到坏是：
shared memory copy
conflict-free transpose
coalesced transpose
naive transpose
前三个基本一样。后面解释
const int TILE_DIM = 32;
const int BLOCK_ROWS = 8;
const int NUM_REPS = 100;
每个块只有32×8的大小，但其处理了32×32的数据，因为核函数中一下处理了4行，如下：
for (int j = 0; j & TILE_DIM; j+= BLOCK_ROWS)
odata[(y+j)*width + x] = idata[(y+j)*width + x];
同时NUM_REPS是为了运行每个核函数100次并取平均带宽值作为结果
各个核函数的含义：
naive transpose：
这是实现矩阵转置最朴素的方法，行访问列，其中行访问合并，但列访问不合并。它是所以性能的下限，即最差的情况。
for (int j = 0; j & TILE_DIM; j+= BLOCK_ROWS)
odata[x*width + (y+j)] = idata[(y+j)*width + x];
是理论带宽的最大值，它的读取和存储都是合并，是理论峰值。这问题是转置问题，存储合并不可能都是合并访问的，所以copy不是转置，而是简单的赋值或者叫复制，看check就看出来了，check的都是输入数组，它的作用是提供一个理论峰值作为参考。如下：都是原数组，没有转置。
postprocess(h_idata, h_cdata, nx*ny, ms);
shared memory copy:
与copy函数一样，这个函数也是copy，只不过是使用shared memory去copy，是shared memory 合并访问全局内存，全局内存合并访问shared memory，这与是无法转置的，也是为了提供理论峰值。同时还可以做到控制变量。
transposeCoalesced：
此例是共享内存合并访问全局内存，全局内存非合并访问共享内存，由于全局内存写的时候访问的是一列数据，正好都在同一个内存存储体中，所以这是个内存存储体冲突较为严重的例子。但其性能还是较不使用共享内存高了很多。
conflict-free transpose
此例是共享内存合并访问全局内存，全局内存非合并访问共享内存，由于全局内存写的时候访问的是一列数据，该函数在共享内存后加了一列，达到了填充的目的：
__shared__ float tile[TILE_DIM][TILE_DIM+1];
填充：因为要访问的数据都是在一列，填充一列之后，要访问的数据就正好全部错开了，完全消除了存储体冲突，性能和合并一样好。
前三个性能一样是因为它们都是合并加载/存储全局变量都是理论峰值
填充在此例中是完美的被运用。
共享内存对程序有很好的优化，但其运用也稍有些复杂，掌握好共享内存的最终目的是让全局内存被合并访问，达到最高的内存效率。
没有更多推荐了，
加入CSDN，享受更精准的内容推荐，与500万程序员共同成长！温馨提示！由于新浪微博认证机制调整，您的新浪微博帐号绑定已过期，请重新绑定！&&|&&
LOFTER精选
网易考拉推荐
用微信&&“扫一扫”
将文章分享到朋友圈。
用易信&&“扫一扫”
将文章分享到朋友圈。
我需要一个数组，动态分配，那么直接调用cudaMalloc来为a分配内存的话，是不行的。具体做法如下[cpp]&[cpp]&使用cudaMemcpyToSymbol来把一个动态分配的设备指针写入一个静态的符号。所以是sizeof(int *)，只是把指针写给a。2. 定义和使用常量变量1、赋值方法__constant__ char g_chFuncIcutilSafeCall( cudaMemcpyToSymbol( "g_chFuncIndex", i_chFuncIndex, sizeof( char ) ) ) ;&2、访问权限&& 只可以在cuda的__global__ 或者__device__函数中访问，不可以在外部C函数（extern "C" ）中使用3.定义和使用共享变量在__global__核函数内部定义__shared__float Mds[TILE_WIDTH][TILE_WIDTH];直接使用即可。4.定义和使用local类型变量默认在__global__内部直接定义的int a，即为local类型的变量
阅读(1742)|
用微信&&“扫一扫”
将文章分享到朋友圈。
用易信&&“扫一扫”
将文章分享到朋友圈。
历史上的今天
在LOFTER的更多文章
loftPermalink:'',
id:'fks_',
blogTitle:'cuda中怎么使用各类型变量的数据',
blogAbstract:'1 定义和使用全局变量[cpp]&
{if x.moveFrom=='wap'}
{elseif x.moveFrom=='iphone'}
{elseif x.moveFrom=='android'}
{elseif x.moveFrom=='mobile'}
${a.selfIntro|escape}{if great260}${suplement}{/if}
{list a as x}
推荐过这篇日志的人：
{list a as x}
{if !!b&&b.length>0}
他们还推荐了：
{list b as y}
转载记录：
{list d as x}
{list a as x}
{list a as x}
{list a as x}
{list a as x}
{if x_index>4}{break}{/if}
${fn2(x.publishTime,'yyyy-MM-dd HH:mm:ss')}
{list a as x}
{if !!(blogDetail.preBlogPermalink)}
{if !!(blogDetail.nextBlogPermalink)}
{list a as x}
{if defined('newslist')&&newslist.length>0}
{list newslist as x}
{if x_index>7}{break}{/if}
{list a as x}
{var first_option =}
{list x.voteDetailList as voteToOption}
{if voteToOption==1}
{if first_option==false},{/if}&&“${b[voteToOption_index]}”&&
{if (x.role!="-1") },“我是${c[x.role]}”&&{/if}
&&&&&&&&${fn1(x.voteTime)}
{if x.userName==''}{/if}
网易公司版权所有&&
{list x.l as y}
{if defined('wl')}
{list wl as x}{/list}2.3k 人阅读
标签：至少1个，最多5个
本文翻译自NVIDIA官方博客Parallel Forall，内容仅供参考，如有疑问请访问原网站:。
在以前发布的文章中，我们学习了被一组线程访问的全局内存如何被合并为一次事务以及对于不同的CUDA硬件，对齐和步长如何影响合并访问。对于最近的CUDA硬件，没有对齐的数据访问并不是什么大问题。然而不论是哪一代的CUDA硬件，跨越全局存储器都是个大问题，而且在很多情况下也是很难避免的，例如沿着第二和更高维度访问多维阵列中的元素时。但是，如果我们使用共享存储器的话，也是有可能进行合并访问的。在我向你说明如何避免直接跨越全局存储器之前，我首先需要详细地介绍一下共享存储器。
共享存储器
因为它是一个片上存储器，所以共享存储器比本地存储器和全局存储器要快得多。实际上共享存储器的延迟大约比没有缓存的全局存储器低100倍(假设线程之间没有bank冲突，在之后的文章中我们会介绍)。共享存储器被分配给每个线程块，所以块内的线程可以访问同一个共享存储器。线程可以访问共享内存中由同一线程块中的其他线程从全局内存加载的数据。这种能力（与线程同步相结合）具有许多用途，例如用户管理的数据高速缓存，高性能并行协作算法（例如并行归约），并且在其它情况不可能的情况下促进全局存储器的合并访问 。
当在线程之间共享数据时，我们需要小心以避免竞态条件(race conditions)，因为线程块中的线程之间虽然逻辑上是并行的，但是物理上并不是同时执行的。让我们假设线程A和线程B分别从全局存储器中加载了一个数据并且将它存到了共享存储器。然后，线程A想要从共享存储器中读取B的数据，反之亦然。我们还要假设线程A和B位于不同的warp。如果在A尝试读取B的数据时，B还未写入，这样就会导致未定义的行为和错误的结果。
为了保证在并行线程协作时得到正确的结果，我们必须对线程进行同步。CUDA提供了一个简单的栅栏同步原语，__syncthreads()。每个线程只能在块中所有的线程执行完__syncthreads()函数后，才能继续执行__syncthreads()的语句。因此我们可以在向共享存储器存数据后以及在向共享存储器加载数据前调用__syncthreads()，这样就避免了上面所描述的竞态条件(race conditions)。我们必须要牢记__syncthreads()被用在分支代码块中是未定义的行为，很可能会导致死锁——线程块中所有的线程必须在同一点调用__syncthreads()
共享内存的例子
在设备代码中声明共享内存要使用__shared__变量声明说明符。在核函数中有多种方式声明共享内存，这取决于你要申请的内存大小是在编译时确定还是在运行时确定。下面完整的代码(可以在Github上)展示了使用共享内存的两种方法。
#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 & i++) {
r[i] = n-i-1;
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 & 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 & i++)
if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
上面的代码使用共享存储器对大小为64的数组进行逆序处理。这两个核函数十分相似，不同之处在于共享内存数组的声明以及核函数的调用。
静态的共享内存
如果共享内存数组的大小在编译时就可以确定，就像在上节代码中staticReverse核函数中写的那样，我们就可以显式地声明固定大小的数组，下面是我们声明的s数组：
__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];
在这个核函数中，t和tr分别代表了原始和倒序之后数组的下标索引。每个线程使用语句s[t] = d[t]将全局内存的数据拷贝到共享内存，反向工作是通过语句d[t] = s[tr]来完成的。但是在执行线程访问共享内存中被线程写入的数据前，记住要使用__syncthreads()来确保所有的线程都已经完全将数据加载到共享内存。
在这个例子中，使用共享内存是用于促进全局内存合并访问(在旧的CUDA设备上，计算能力1.1或更低)。对于读取和写入都实现了最优的全局存储器合并，因为全局内存总是通过线性对齐的索引t来访问的。反向索引tr仅用于访问共享存储器，其不具有全局存储器的顺序访问限制，因此不能获得最佳性能。共享内存的唯一性能问题是bank冲突，我们之后会做讨论。
NOTE:注意在计算能力为1.2或更高版本的设备上，内存系统仍然可以完全地合并访问，即使是反向的保存在全局存储器中。这一技术在其他访问模式下也是很有用的，我会在下一篇博客中介绍。
动态的共享内存
另一个核函数使用了动态分配共享内存的方式，这主要用于共享内存的大小在编译时不能确定的情况。在这种情况下，每个线程块中共享内存的大小必须在核函数第三个执行配置参数中指定(以字节为单位)，如下所示：
dynamicReverse&&&1, n, n*sizeof(int)&&&(d_d, n);
该动态共享内存的核函数dynamicReverse()使用了未指定大小的extern数组语法(extern __shared__ int s[])来声明共享内存数组。
NOTE:注意中括号与extern说明符。
当核函数被启动时，数组大小从第三个执行配置参数被隐式地确定。该核函数其余部分的代码与staticReverse()核函数相同。
而如果你想在一个核函数中动态地申请多个数组时该怎么办呢？你必须在首先申请一个单独的未指定大小的extern数组，然后使用指针将它分为多个数组，如下所示：
extern __shared__ int s[];
int *integerData =
// nI ints
float *floatData = (float*)&integerData[nI]; // nF floats
char *charData = (char*)&floatData[nF];
// nC chars
这样的话，你需要在核函数中这样指定共享内存的大小：
myKernel&&&gridSize, blockSize, nI*sizeof(int)+nF*sizeof(float)+nC*sizeof(char)&&&(...);
共享内存的bank冲突
为了获得较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，可以被同时访问。因此任何跨越b个不同bank的n个地址的读写操作可以被同时进行，这样就大大提高了整体带宽 ——可达到单独一个bank带宽的b倍。
然而，如果多个线程请求的地址映射到相同的内存bank，那么访问就会被顺序执行。硬件会把冲突的内存请求分为尽可能多的单独的没有冲突的请求，这样就会减少一定的带宽，减少的因子与冲突的内存请求个数相等。当然，也有例外的情况：当一个warp中的所有线程访问同一个共享内存地址时，就会产生一次广播。计算能力为2.0及以上的设备还可以多次广播共享内存访问，这意味着一个warp中任意数量的线程对于同一位置的多次访问也可以同时进行。
译者注:这里关于warp的多播与bank冲突原文并未详细介绍，详细内容及例子可以参考CUDA programming guide。我在后续的博客中也会详细介绍这部分。
为了尽量减少bank冲突，理解共享内存地址如何映射到bank是非常重要的。共享内存的bank是这样组织的：连续的32-bits字被分配到连续的bank中，每个bank的带宽是每个时钟周期32bits。
译者注:这里不同计算能力的bank的带宽是不同的，原文提到的带宽大小是计算能力5.0的设备，对于计算能力2.0的设备每个bank的带宽是每两个时钟周期32bits；对于计算能力3.0的设备，每个bank的带宽是每个时钟周期64bits。详情请参考CUDA C programming guide。
对于计算能力1.x的设备，warp的大小是32而bank的数量是16。一个warp中线程对共享内存的请求被划分为两次请求：一个请求是前半个warp的另一个请求时后半个warp的。注意如果每个bank中只有一个内存地址是被半个warp中的线程访问的话，是不会有bank冲突的。
对于计算能力为2.x的设备，warp的大小是32而bank的数量也是32。一个warp中线程对共享内存的请求不会像计算能力1.x的设备那样被划分开，这就意味着同一个warp中的前半个warp中的线程与后半个warp中的线程会有可能产生bank冲突的。
计算能力为3.x的设备的bank大小是可以配置的，我们可以通过函数来设置，要么设置为4字节(默认为cudaSharedMemBankSizeFourByte)，要么设置为8字节(cudaSharedMemBankSizeEightByte)。当访问双精度的数据时，将bank大小设置为8字节可以帮助避免bank冲突。
配置共享内存的数量
在计算能力为2.x和3.x的设备上，每个多处理器有64KB的片上内存，它可以被划分为L1高速缓存和共享内存。对于计算能力为2.x的设备，总共有两种设置：48KB的共享内存/16KBL1高速缓存和16KB的共享内存/16KB的L1高速缓存。我们可以在运行时使用在主机端为所有的核函数配置或者使用为单个的核函数配置。它们有三个选项可以设置：cudaFuncCachePreferNone(在共享内存和L1中不设置首选项，即使用默认设置), cudaFuncCachePreferShared(共享内存大于L1), 和cudaFuncCachePreferL1(L1大于共享内存)。驱动程序将按照指定的首选项，除非核函数中每个线程块需要比指定配置中更多的共享内存。在计算能力3.x的设备上允许有第三种设置选项——32KB的共享内存/32KB的L1高速缓存，可以通过cudaFuncCachePreferEqual选项设置。
对于写出高性能的CUDA代码，共享内存的确是一个十分强大的特性。由于共享内存位于片上，所以访问共享内存比访问全局内存快很多。由于共享内存在线程块中可以被线程共享，所以才提供了相应的机制来保证线程的正常协作。使用共享内存来利用这种线程协作的一种方法是启用全局内存的合并访问，正如如本文中的数组逆序所演示的。在使用共享内存来使数组逆序的例子中，我们可以使用单位步长执行所有全局内存读取和写入，从而在任何CUDA GPU上实现完全地合并访问。
0 收藏&&|&&3
你可能感兴趣的文章
你可能感兴趣的文章
分享到微博？
我要该，理由是：
在 SegmentFault，学习技能、解决问题
每个月，我们帮助 1000 万的开发者解决各种各样的技术问题。并助力他们在技术能力、职业生涯、影响力上获得提升。