关于threadfence、volatile、threadfence_block、threadfence_system、syncthreads的一些见解和问题

  • 4 replies
  • 1886 views
volatile:
读写数据都会bypass掉L1 cache;

syncthreads:
1. (threadblock-wide) execution barrier,当同一个block中的线程都执行到该处时才会继续往下执行;
2. (device-wide) memory fence, 保证该语句前的write的数据对同一个block中的其他线程可见,这个可见意味着数据被写入L2 cache,同时因为L1 cache是write through cache,所以L1 cache中的数据也会被更新;

threadfence:
1. 有了这个指令后, 线程则会暂停, 等待写入的过程完成到一定程度,保证线程自己的写入操作能被整个grid的线程"看"到, 才会继续执行,但是不会导致其他线程的停止运行,这里的看到也是指数据被写入L2 cache,同时因为L1 cache是write through cache,所以L1 cache中的数据也会被更新,但是仅限于已经执行该指令的SM,对于其他的SM该数据的缓存则不会改变或者失效,其他的线程仍然有可能会从L1 cache中读取到老旧的数据;
2. 将threadfence函数前后的访存隔离开, 编译器不会跨越threadfence的边界, 不会重新调整访存语句的顺序


但是threadfence_block和threadfence_system我是不太清楚的,这里的可见性具体是写入到什么程度呢?
threadfence_block也是写入到L2 cache才算对同一个block内的线程可见? threadfence_system是写入到global memory才算对多个显卡的线程可见?

volatile:
读写数据都会bypass掉L1 cache;

syncthreads:
1. (threadblock-wide) execution barrier,当同一个block中的线程都执行到该处时才会继续往下执行;
2. (device-wide) memory fence, 保证该语句前的write的数据对同一个block中的其他线程可见,这个可见意味着数据被写入L2 cache,同时因为L1 cache是write through cache,所以L1 cache中的数据也会被更新;

threadfence:
1. 有了这个指令后, 线程则会暂停, 等待写入的过程完成到一定程度,保证线程自己的写入操作能被整个grid的线程"看"到, 才会继续执行,但是不会导致其他线程的停止运行,这里的看到也是指数据被写入L2 cache,同时因为L1 cache是write through cache,所以L1 cache中的数据也会被更新,但是仅限于已经执行该指令的SM,对于其他的SM该数据的缓存则不会改变或者失效,其他的线程仍然有可能会从L1 cache中读取到老旧的数据;
2. 将threadfence函数前后的访存隔离开, 编译器不会跨越threadfence的边界, 不会重新调整访存语句的顺序


但是threadfence_block和threadfence_system我是不太清楚的,这里的可见性具体是写入到什么程度呢?
threadfence_block也是写入到L2 cache才算对同一个block内的线程可见? threadfence_system是写入到global memory才算对多个显卡的线程可见?

(1)如同我们用的CPU编译器(例如gcc好了)在特定的CPU硬件上的编译。任何对threadfence的指定都有2个效果,一个是软件效果(编译器效果),一个是硬件效果(CPU)。GPU也同理。

前者可能会导致编译器对threadfence前后的读写指令的妥当安排(例如可能不会跨越fence将后面的读取指令调到前面,等于说你用fence限制了编译器的对你的原始代码行的自由调整的发挥)。后者可能在特定的硬件上生成特定的指令,要求硬件确定特定的指令生效顺序(例如可能要求某些写入指令在某些单元上,至少执行到某些特定的结果,从而使得其他单元能够看到写入效果)。

(2)__syncthreads()隐含了fence效果。等于fence调用一次+线程块同步一次(+可选的块内数据迭代)(看后缀)。

(3)volatile我不懂。但有的时候是可以越过L2的。例如这个例子:
__global__ void labandlib(...., volatile int *flag)
{
    while(....)
    {
         ....//long running kernel code
         if (*flag == 0x12345678) return; //early return
         ....//long running kernel code   
    }
}
其中flag是通过cudaMallocHost分配的,并随时通过CPU设定值来打断长期运行的kernel的执行的。

这种代码可以正常执行。所以它在某些情况下至少得bypass L2. 更多的得你教晓我。

(1)如同我们用的CPU编译器(例如gcc好了)在特定的CPU硬件上的编译。任何对threadfence的指定都有2个效果,一个是软件效果(编译器效果),一个是硬件效果(CPU)。GPU也同理。

前者可能会导致编译器对threadfence前后的读写指令的妥当安排(例如可能不会跨越fence将后面的读取指令调到前面,等于说你用fence限制了编译器的对你的原始代码行的自由调整的发挥)。后者可能在特定的硬件上生成特定的指令,要求硬件确定特定的指令生效顺序(例如可能要求某些写入指令在某些单元上,至少执行到某些特定的结果,从而使得其他单元能够看到写入效果)。

(2)__syncthreads()隐含了fence效果。等于fence调用一次+线程块同步一次(+可选的块内数据迭代)(看后缀)。

(3)volatile我不懂。但有的时候是可以越过L2的。例如这个例子:
__global__ void labandlib(...., volatile int *flag)
{
    while(....)
    {
         ....//long running kernel code
         if (*flag == 0x12345678) return; //early return
         ....//long running kernel code   
    }
}
其中flag是通过cudaMallocHost分配的,并随时通过CPU设定值来打断长期运行的kernel的执行的。

这种代码可以正常执行。所以它在某些情况下至少得bypass L2. 更多的得你教晓我。
https://developer.download.nvidia.com/video/gputechconf/gtc/2020/presentations/s21819-optimizing-applications-for-nvidia-ampere-gpu-architecture.pdf这篇文章第七页中图示对于system memory不会在L2 上cache,上述的例子会不会是这个原因


https://developer.download.nvidia.com/video/gputechconf/gtc/2020/presentations/s21819-optimizing-applications-for-nvidia-ampere-gpu-architecture.pdf这篇文章第七页中图示对于system memory不会在L2 上cache,上述的例子会不会是这个原因

有可能是这个原因。但是去掉了volatile后永远不会检测到正确的flag了。

实际上有没有volatile是两种指令,
一种是:
LDG.STRONG.SYS
另外一种:LDG
所以可能L2也会缓冲系统内存(因为后一种多次读取都检测不到,可能是kernel运行期间将旧值缓冲了)。

有可能是这个原因。但是去掉了volatile后永远不会检测到正确的flag了。

实际上有没有volatile是两种指令,
一种是:
LDG.STRONG.SYS
另外一种:LDG
所以可能L2也会缓冲系统内存(因为后一种多次读取都检测不到,可能是kernel运行期间将旧值缓冲了)。
嗯嗯从结果来看是会缓冲的