CUDA复制测试

  这里主要是测试了内存数据读写操作的几种方式，记录了一些测试结果，对于二维数组(10244)(1024*4)。

• (1)二维线程格，每个线程对应一个元素。
• (2)转换为int2类型，线程宽度减半。
• (3)线程宽度和高度减半，单个线程操作邻近的4个元素。
• (4)线程宽度和高度减四分之一，单个线程操作邻近的16个元素。
• (5)线程宽度和高度减半，单个线程操作完一个元素后，跨区域循环。
• (6)为单个线程处理16个元素，采用表面引用。
• (7)为单个线程处理64*64个元素，采用表面引用。
• (8)跨区域循环，单线程处理16个元素，采用设备内存。
• (9)跨区域循环，单线程处理16个元素，采用表面引用。
• (10)跨区域循环，单线程处理64*64元素，采用表面引用。
• (11)跨区域循环，单线程处理64*64元素，采用设备内存。
• (12)跨区域循环，单线程处理2*2元素，采用表面引用。
• (13)单线程处理单个元素，采用表面引用。

  在上面测试结果中，第一个和第二个的读取速度数据中，int2类型的速度是int类型的速度的大约两倍，这是来自于内存事务的合并。为了达到读写数据时的最佳性能，CUDA内核必须执行内存事务合并处理。任何不满足合并所需的全套标准的内存事务称为“未合并的”，未合并内存事务的性能惩罚由2~8倍不等，在最近的硬件上，合并内存事务对性能的影响显著减少。
  事务在每一个线程束的基础上被合并，为了由束执行内存读写事务的合并，一些简化的标准必须要满足。

• 字至少为32位，读写字节或16位字的事务总是非合并的。
• 束上的线程访问的地址是连续并递增的（即，依线程ID偏移）。
• 束的基地址（束中第一个线程访问的地址）需要对齐。

template<class T, const int n>
__global__ void GlobalWrites(T *out, T value, size_t N){
    size_t i;
    for(i = n*blockIdx.x*blockDim.x+threadIdx.x;
        i < N-n*blockDim.x*gridDim.x;
        i += n*blockDim.x*gridDim.x){
        for(int j = 0; j < n; j++){
            size_t index = i+j*blockDim.x;
            out[index] = value;
        }
    }
    // to avoid the (index<N) conditional in the inner loop,
    // we left off some work at the end
    for(int j = 0; j < n; j++){
        size_t index = i+j*blockDim.x;
        if(index < N) out[index] = value;
    }
}

template<class T, const int n, bool bOffset>
double ReportRow(size_t, N, size_t threadStart, size_t threadStop, size_t cBlocks){
    int maxThreads = 0;
    double maxBW = 0.0;
    printf("%d\t", n);
    for(int cThreads = threadStart; cThreads <= threadStop; cThreads *= 2){
        double bw;
        bw = BandwidthWrites<T, n, bOffset>(N, cBlocks, cThreads);
        if(bw > maxBW){
            maxBW = bw;
            maxThreads = cThreads;
        }
        printf("%.2f\t", bw);
    }
    printf("%.2f\t%d\n", maxBW, maxThreads);
    return maxBW;
}