将 16 字节和 256 字节复制到共享内存之间的区别答案

【问题标题】：difference between copying 16-bytes and 256-bytes to shared memory将 16 字节和 256 字节复制到共享内存之间的区别
【发布时间】：2015-07-06 09:52:48
【问题描述】：

我正在编写一个 CUDA 程序。代码将固定内存复制到共享内存，固定内存是指分配有cudaHostAlloc(., ., cudaHostAllocMapped) 的内存。复制 16 字节需要 600us，复制 256 字节需要 8ms。为什么差别这么大？

我的代码如下：

__global__
void kernel_func(char* dict, int dict_len)
{
    __shared__ char s_dict[256];

    /* dict_len = 16; */
    if(threadIdx.x == 0) {// copy once for each block
        memcpy((unsigned char*)s_dict, (unsigned char*)dict, dict_len);
    }
    __syncthreads(); 
}

kernel_func<<<32, 128>>>("256 bytes pinned memory", 256);

环境：GTX650 + CUDA 6.5 + Win7-32bit

【问题讨论】：

600 us - 16 bytes * 16 -> 9.6 ms - 256 bytes。 9.6 ms ~= 8 ms 没那么巨大。
600µs 远高于预期。内存副本是否执行 16 x 1 字节读取、1 字节写入。如果 dict 是 16 字节对齐的，那么您可以使用单个 LD.128 读取数据并使用 LDS.128 写入数据；否则，最好的选择是让前 16 个线程在一条指令中加载数据并将数据存储在一条指令中。对于 16 个字节，我的期望是

标签： c++ cuda shared-memory

【解决方案1】：

因为 256 字节是您之前复制的（16 字节）的 16 倍

现在 16 字节 占用了 600us

而16倍也就是9600us接近你观察到的8000us（1ms = 1000us）

【讨论】：

啊哈 :) 为什么内核每次传输 16 字节？有关于 16 字节的文档吗？
不，这只是您在原始问题中所说的数字的巧合。它只是：16 字节 = 600us 因此 1 字节 = 37.5us 然后因此 256 字节 = 9600us = 9.6ms ~ 8ms