如何解决为什么 cuda 指针内存访问比全局设备内存访问慢?
#include <vector_functions.h>
#include <vector_types.h>
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <string>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
__device__ int foo[16];
__device__ int bar[16];
__global__ void go(const int* ptr) {
printf("device: tid = %d,foo = %p\n",blockIdx.x,foo);
printf("device: tid = %d,ptr = %p\n",ptr);
int val = threadIdx.x;
for (int i = 0; i < (1 << 20); i++) {
bar[blockIdx.x] = val;
val = (val * 19 + ptr[threadIdx.x]) % (int)(1e9 + 7); // change ptr to foo for experiment
}
}
int main() {
int* ptr = nullptr;
cudaGetSymbolAddress((void**)&ptr,foo);
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
go<<<16,16>>>(ptr);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaDeviceSynchronize();
float ms;
cudaEventelapsedtime(&ms,start,stop);
printf("%.6fms\n",ms);
return 0;
}
在我的 GeForce GTX 1080 上:
使用 ptr
需要 180 毫秒,但使用 foo
只需要 36 毫秒,尽管 ptr
和 foo
指向完全相同的地址。我认为它们应该以相同的速度执行,因为它们都是由 L2 缓存的全局内存。
我使用的是 Linux,我的编译命令是:
nvcc -gencode=arch=compute_61,code=compute_61 -Xptxas -O3 test.cu -o test
谁能解释一下为什么?
解决方法
这两种情况不同的原因是,当显式使用 foo
时,编译器(在这种情况下为 ptxas
)知道 foo
不会 {{3 }} bar
,等可以做具体的优化。当改用内核参数 ptr
时,编译器不知道此别名是否发生,并假设它可能发生。这对设备代码生成有重大影响。
作为证明点,使用以下内核原型重新编译您的测试用例:
__global__ void go(const int* __restrict__ ptr) {
你会看到时差消失了。这是 alias 编译器,ptr
不能为任何其他已知位置(例如 bar
)设置别名,因此这允许在两种情况下生成类似的代码。 (在现实世界中,当您准备与编译器签订此类合同时,您才会/应该只使用此类装饰。)
详情:
请务必记住,设备代码编译器是一个优化编译器。此外,从单线程的角度来看,设备代码编译器主要对正确性感兴趣。多线程访问同一个位置不是考虑到这个答案,也确实不是设备代码编译器考虑的。当多个线程访问同一位置时,确保正确性是程序员的责任。
有了那个序言,这里的主要区别似乎是优化之一。知道 foo
(或 ptr
)没有别名 bar
并且仅考虑单个执行线程,很明显您的内核循环代码可以重写为:
int val = threadIdx.x;
int ptrval = ptr[threadIdx.x]; // becomes a LDG instruction
for (int i = 0; i < ((1 << 20)-1); i++) {
val = (val * 19 + ptrval) % (int)(1e9 + 7);
}
bar[blockIdx.x] = val; // becomes a STG instruction
此优化的一个主要影响是我们从多次写入 bar
变为仅写入一次。通过这种优化,ptr
的读取也可以“优化到寄存器中”(因为我们现在知道它是循环不变的)。最终效果是消除了循环中的所有全局加载和存储。另一方面,如果 ptr
可能会或可能不会别名 bar
,那么我们必须考虑这种可能性,并且上述优化将不成立。
这似乎是编译器正在做的事情。在我们使用 foo
(或 __restrict__
)的情况下,编译器(在 sass 代码中)在开头安排了一个全局加载,在结尾安排了一个全局存储,以及部分展开的充满整数运算的循环。
然而,当我们将代码保持原样/发布时,编译器也部分展开了循环,但在部分展开的循环中散布了 LDG
和 STG
指令。>
您可以使用 informing 自行观察,例如:
cuobjdump -sass test
(对于每种情况)
设备代码 printf
语句不会实质性地改变这里的任何观察结果,因此为了分析的简单起见,我将删除它们。
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。