2.2CUDA-Memory(存储)和bank-conflict

在CUDA基本概念介绍有简单介绍CUDA memory。这里详细介绍：

每一个线程拥有自己的私有存储器，每一个线程块拥有一块共享存储器(Shared memory)；最后，grid中所有的线程都可以访问同一块全局存储器(global memory)。除此之外，还有两种可以被所有线程访问的只读存储器：常数存储器（constant memory）和纹理存储器(Texture memory),它们分别为不同的应用进行了优化。全局存储器、常数存储器和纹理存储器中的值在一个内核函数执行完成后将被继续保持，可以被同一程序中其也内核函数调用。

下表给出了这8种存储器的位置、缓存情况，访问权限及生存域

存储器	位置	拥有缓存	访问权限	变量生存周期
register	GPU片内	N/A	Device可读/写	与thread相同
Local memory	板载显存	无	Device可读/写	与thread相同
Shared memory	GPU片内	N/A	Device可读/写	与block相同
Constant memory	板载显存	有	Device可读，host可读写	可在程序中保持
Texture memory	板载显存	有	Device可读，host可读写	可在程序中保持
Global memory	板载显存	无	Device可读/写, host可读/写	可在程序中保持
Host memory	Host内存	无	host可读/写	可在程序中保持
Pinned memory	Host内存	无	host可读/写	可在程序中保持

kernel变量定义，使用范围和生命周期。

2.2CUDA-Memory(存储)和bank-conflict

其中__shared__和__constant__前面的__device__声明是可以省略的

Global memory，如果有一个thread修改啦global memory的值，其他的thread不能立即看到这个值的变化。需要终止这个kernel，然后lanuch一个新的kernel，这样新的kernel能看到global memory的变化。

由于访问速度比Global快的多，比如向量加法，每次从in指针里面取内容都是从global里面取。比如矩阵乘法：

__global__ void MatrixMulKernel(int m, int n, int k, float* A, float*
B, float* C)
{
    int Row = blockIdx.y*blockDim.y+threadIdx.y;
    int Col= blockIdx.x*blockDim.x+threadIdx.x;
    if ((Row < m) && (Col < k)) {
        float Cvalue = 0.0;
        for (int i = 0; i < n; ++i)
            /* A[Row, i] and B[i, Col] */
            Cvalue += A[Row*n+i] * B[Col+i*k];
        C[Row*k+Col] = Cvalue;
    }
}

View Code