CUDA全局内存访问速度答案

【问题标题】：CUDA global memory access speedCUDA全局内存访问速度
【发布时间】：2014-01-10 20:46:51
【问题描述】：

这是简单的 cuda 代码。
我正在测试访问全局内存的时间。阅读并正确。

下面是核函数(test1())。

enter code here

__global__ void test1(int *direct_map)   
{  
    int index = 10;  
    int index2;  

    for(int j=0; j<1024; j++)  
    {  
        index2 = direct_map[index];  
        direct_map[index] = -1;  
        index = index2;  
    }  
}

direct_map 是 683*1024 的线性矩阵，每个像素都有一个偏移值来访问其他像素。
index 和 index2 不是续地址。

这个核函数大约需要 600 微秒。

但是，如果我删除代码，
direct_map[索引] = -1;

只需要 27 微秒。

我认为代码已经从全局内存中读取了direct_map[index]的值

index2 = direct_map[index];

那么，它应该位于二级缓存。
所以，在做“direct_map[index] = -1;”的时候，速度要快。

而且，我测试了对全局内存的随机写入 (test2())。

大约需要 120 微秒。

enter code here

__global__ void test2(int *direct_map)   
{  
    int index = 10;  

    for(int j=0; j<1024; j++)  
    {  
        direct_map[index] = -1;  
        index = j*683 + j/3 - 1;  
    }  
}

所以，我不知道为什么 test1() 需要超过 600 微秒。谢谢。

【问题讨论】：

标签： memory cuda global

【解决方案1】：

当你删除代码行时：

direct_map[index] = -1;

你的内核没有做任何有用的。编译器可以识别这一点并消除与内核启动相关的大部分代码。对内核代码的修改意味着内核不再影响任何全局状态，从编译器的角度来看，代码实际上是无用的。

您可以通过转储编译器在每种情况下生成的汇编代码来验证这一点，例如cuobjdump -sass myexecutable

任何时候您对代码进行了微小的更改并看到时序发生了很大的变化，您应该怀疑您所做的更改允许编译器做出不同的优化决策。

【讨论】：