CUDA内存栅栏（Memory Fence）理解
枝江的雨我的口水
2022年09月22日 11:12

CUDA中文教材的翻译太晦涩了，英文文档也全都是长难句，导致内存栅栏非常难理解。在本文中我会重新梳理一遍自己的理解，并用实验验证。

让我们在全局内存中初始化两个变量：

代码块
text/x-c++src
自动换行
复制代码
__device__ int X = 1, Y = 2;
复制成功

然后在线程 i 中修改这两个变量：

代码块
text/x-c++src
自动换行
复制代码
__device__ void writeXY()
{
    X = 10;
    Y = 20;
}
复制成功

从线程 i 的视角来看，变量X的写入在变量Y的写入之前，这看似是对的，但编译器会发现X和Y之间没有依赖关系，然后把它们优化成并行写入，所以此时如果用另一个并行的线程读取X和Y变量，X在Y之前写入的这个顺序就无法保证。读取操作也是如此，顺序无法得到保障。

读取变量：

代码块
text/x-c++src
自动换行
复制代码
__device__ void readXY()
{
    int B = Y;
    int A = X;
    printf("A=%d, B=%d", A, B);
}
复制成功

假设write和read函数分别由并行的线程 i 和线程 j 运行，此时打印出来可能有四种结果：

    A=1, B=2 ；   说明 j 在 i 之前读取了变量，可以接受

    A=10, B=20；说明 i 在 j 之前写入了变量，可以接受

    A=10, B=2；  说明 i 已经写入A，但没来得及写入B，j 已经读取了变量，可以接受

    A=1, B=20;     说明 i 写入顺序被并行化，而B写在了A前面，这是不能接受的情况，因为它可能会误导其他线程

为了避免第四种情况的出现，需要用内存栅栏函数，来保证线程 i 的读写顺序从线程 j 的视角来看不会颠倒。完整代码如下：

代码块
text/x-c++src
自动换行
复制代码
#ifndef __CUDACC__
#define __CUDACC__
#endif
#include<iostream>
#include<cuda_runtime.h>
#include<cooperative_groups.h>

namespace cg = cooperative_groups;

__device__ int X = 1, Y = 2;

__device__ void writeXY()
{
    X = 10;
    __threadfence();
    Y = 20;
}
__device__ void readXY()
{
    int B = Y;
    __threadfence();
    int A = X;
    printf("A=%d, B=%d\n", A, B);
}

__global__ void kernel(){
    cg::thread_block tb = cg::this_thread_block();

    if(tb.thread_rank() == 0)  writeXY();
    if(tb.thread_rank() == 1) readXY();
}

int main(){
    cudaLaunchCooperativeKernel(kernel, 1, 2, NULL, 0, 0);
    cudaDeviceSynchronize();
    return 0;
}

/* 运行结果: A=10, B=20 */
复制成功

我们改变一下线程配置，就能出现不同的结果：

代码块
text/x-c++src
自动换行
复制代码
__global__ void kernel(){
    cg::thread_block tb = cg::this_thread_block();

    if(tb.thread_rank() == 0)  writeXY();
    if(tb.thread_rank() == 32) readXY();
}

int main(){
    cudaLaunchCooperativeKernel(kernel, 1, 64, NULL, 0, 0);
    cudaDeviceSynchronize();
    return 0;
}

/* A=10, B=2 */
复制成功

代码块
text/x-c++src
自动换行
复制代码
__global__ void kernel(){
    cg::thread_block tb = cg::this_thread_block();

    if(tb.thread_rank() == 0)  readXY();
    if(tb.thread_rank() == 1) writeXY();
}

int main(){
    cudaLaunchCooperativeKernel(kernel, 1, 2, NULL, 0, 0);
    cudaDeviceSynchronize();
    return 0;
}

/* A=1, B=2 */
复制成功

无论线程配置如何改动，都只会出现前三种可以接受的情况，而第四种情况是不可能出现的，因为内存栅栏__threadfence保证了读写顺序。

新版的内存栅栏函数似乎改名了（cuda::atomic_thread_fence），旧版的函数名可能会在未来被弃用。

cv18722474