CUDA中共享内存、常量内存和纹理内存的概念和应用（小白入门）

这篇写的很好

1 共享内存（一个block 48KB=49152B ）

常用于向量点积¹

■ 在考虑线程块大小的时候经常用到向上取整，这里使用了技巧 ceil( a / b ) == floor( (a-1) / b) + 1

■ 算法总体想法是在GPU中将很长的向量分段放入GPU的各线程块中，每个线程块利用共享内存和多线程分别计算乘法和加法。结果整理为每个线程块输出一个浮点数，置于全局内存中，这样就将待计算的元素数量降到了 gridDim.x 的水平，再返回CPU中完成剩下的加法。

■ 算法预先规定了每个线程块使用256个线程（blockDim.x == 256），那么使用的线程块数量应该满足 gridDim.x * blockDim.x ≥ N（待计算的向量长度），另外代码中规定线程块数量至少为32（“选择其他的只可能产生更高或更差的性能，这取决于CPU和GPU的相对速度”）

■ 在核函数中使用了既定大小的共享内存 shared float cache[threadsPerBlock]; ，并采用 __syncthreads(); 函数进行线程同步（因为接下来要进行规约运算，前提就是该线程块内所有的线程已经独立计算完毕）。

/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * NVIDIA Corporation and its licensors retain all intellectual property and 
 * proprietary rights in and to this software and related documentation. 
 * Any use, reproduction, disclosure, or distribution of this software 
 * and related documentation without an express license agreement from
 * NVIDIA Corporation is strictly prohibited.
 *
 * Please refer to the applicable NVIDIA end user license agreement (EULA) 
 * associated with this source code for terms and conditions that govern 
 * your use of this NVIDIA software.
 * 
 */


#include "../common/book.h"

#define imin(a,b) (a<b?a:b)

const int N = 33 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid =
            imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );


__global__ void dot( float *a, float *b, float *c ) {
   
   
    __shared__ float cache[threadsPerBlock]; //共享内存申请方式：__shared__  变量类型 变量名；
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < N) {
   
   
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    
    // set the cache values
    cache[cacheIndex] = temp;
    
    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
   
   
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}


int main( void ) {
   
   
    float   *a, *b, c, *partial_c;
    float   *dev_a, *dev_b, *dev_partial_c;

    // allocate memory on the cpu side
    a = (float*)malloc( N*sizeof(float) );
    b = (float*)malloc( N*sizeof(float) );
    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
                              N*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
                              N*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
                              blocksPerGrid*sizeof(float) ) );

    // fill in the host memory with data
    for (int i=0; i<N; i++) {
   
   
        a[i] = i;
        b[i] = i*2