3. Memory Access and Shared Memory

本篇总结一下之前提到但没有进一步展开的内存访问问题和共享内存问题。

__device__关键字与内存访问

#include <iostream>

__device__ int cnt1 = 1, cnt2 = 1; // GPU上的全局变量

__global__ void kernal(int type) {
    if (type == 0) {
        printf("int parameter <<<2, 1>>>, grid.shape=(2,), block.shape(1,)\tcnt1=%d\n", atomicAdd(&cnt1, 1));
    }
    else {
        printf("dim3 parameter <<<dim3(2, 2), dim3(3, 3)>>>, grid.shape=(2, 2),"
                                "block.shape=(3, 3)\tcnt2=%d\n", atomicAdd(&cnt2, 1));
    }
}

/*
__global__ void kernal(int type， int& cnt, const char* str) {
    printf("%s\tcnt%d=%d\n", str, type + 1, ++cnt);
}
*/

int main() {
    kernal<<<2, 1>>>(0);
    dim3 gridshape(2, 2);
    dim3 blockshape(3, 3);
    kernal<<<gridshape, blockshape>>>(1);
    cudaDeviceSynchronize();
    return 0;
}

再以该程序为例，首先可以看到__device__关键字可以用于定义函数，也可以定义在GPU上的全局变量，__host__同理，但__global__不能用于定义变量。
当然在设备端函数上的局部变量就是在GPU上的局部变量了。

然后看程序中注释掉的部分，这个核函数尝试使用引用和指针来访问CPU端的内存，这显然是不允许的。类似地，CPU端不能直接访问在GPU端的内存。
因此上一篇的第一个思考题答案就是cudaMalloc需要在CPU端进行调用，但是CPU端不能直接访问A[0]，所以需要用A2来分配内存并在GPU端完成对A[0]的赋值。B同理。
注意指针变量所在的位置和其所指向的内存无关，即CPU端的指针变量可以指向GPU端的内存，GPU端的指针变量也可以指向CPU端的内存。

共享内存

共享内存本质上是GPU上的一种高速缓存，访问速度极快，甚至可以和寄存器相比较。与之相对的其容量比较有限。
每个Block有一个共享内存，正如之前提到的，Block内资源共享，这也是一个体现。
定义共享内存变量与定义局部静态变量有些相似，前者在Block内共享，而后者以及全局静态变量都是在进程内共享。

#include <iostream>

__global__ void f1() {
    // 静态分配
    __shared__ char s[11];
    s[threadIdx.x] = '1';
    if (threadIdx.x == 9) {
        s[10] = '\0';
        printf("f1(): %s\n", s);
    }
} 

__global__ void f2() {
    // 动态分配
    extern __shared__ char s[];
    s[threadIdx.x] = '2';
    if (threadIdx.x == 9) {
        s[10] = '\0';
        printf("f2(): %s\n", s);
    }
}

int main() {
    f1<<<1, 10>>>();
    f2<<<1, 10, 10 * sizeof(char)>>>();
    cudaDeviceSynchronize();

    return 0;
}

共享内存的使用有静态和动态两种。静态分配在核函数中指定，定义时和普通变量一样需要预先确定好大小；动态分配则是在调用核函数时才确定，而extern __shared__ ...的语句只是变量声明，没有定义。
注意extern在这里是和__shared__共同表示动态共享内存，正常情况没有__shared__的话，extern不能用于声明局部变量，而且这样也没有意义。