??这里主要是测试了内存数据读写操作的几种方式,记录了一些测试结果,对于二维数组(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)单线程处理单个元素,采用表面引用。
序号 | 读 | 读 | 读 | 写 | 写 | 写 |
---|
(1) | 123.38 | 129.02 | 129.22 | 295.72 | 304.6 | 261.97 | (2) | 48.11 | 54.82 | 67.93 | 306.24 | 305.16 | 305.50 | (3) | 24.31 | 22.40 | 29.42 | 393.10 | 394.25 | 386.28 | (4) | 6.32 | 5.85 | 6.02 | 485.72 | 469.24 | 467.52 | (5) | 24.71 | 22.20 | 22.27 | 385.06 | 356.74 | 365.97 | (6) | 6.30 | 5.82 | 5.81 | 845.41 | 815.74 | 811.05 | (7) | 0.596 | 0.592 | 0.601 | 3046.6 | 3056.9 | 2988.5 | (8) | 5.86 | 5.90 | 5.80 | 356.65 | 355.79 | 357.08 | (9) | 5.84 | 5.88 | 6.06 | 207.64 | 207.91 | 217.58 | (10) | 0.592 | 0.652 | 0.543 | 240.72 | 240.63 | 240.23 | (11) | 0.580 | 0.590 | 0.591 | 433.59 | 430.95 | 431.52 | (12) | 22.57 | 21.65 | 22.33 | 205.68 | 206.82 | 207.69 | (13) | 91.36 | 89.37 | 89.07 | 208.97 | 209.02 | 208.97 |
??在上面测试结果中,第一个和第二个的读取速度数据中,int2类型的速度是int类型的速度的大约两倍,这是来自于内存事务的合并。为了达到读写数据时的最佳性能,CUDA内核必须执行内存事务合并处理。任何不满足合并所需的全套标准的内存事务称为“未合并的”,未合并内存事务的性能惩罚由2~8倍不等,在最近的硬件上,合并内存事务对性能的影响显著减少。 ??事务在每一个线程束的基础上被合并,为了由束执行内存读写事务的合并,一些简化的标准必须要满足。
- 字至少为32位,读写字节或16位字的事务总是非合并的。
- 束上的线程访问的地址是连续并递增的(即,依线程ID偏移)。
- 束的基地址(束中第一个线程访问的地址)需要对齐。
字大小 | 对齐 |
---|
8位 | 未合并 | 16位 | 未合并 | 32位 | 64字节 | 64位 | 128字节 | 128位 | 256字节 |
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;
}
}
for(int j = 0; j < n; j++){
size_t index = i+j*blockDim.x;
if(index < N) out[index] = value;
}
}
|