CUDA开发：了解设备属性
　作者: 流浪念枫雪　编辑:
&&&&【IT168&技术】之前我们为大家介绍了《》这篇文章，今天介绍一下CUDA设备的相关属性，只有熟悉了硬件是相关属性，是怎么工作的，就能写出更适合硬件工作的代码。cudaDeviceProp这个结构体记录了设备的相关属性。 1 struct cudaDeviceProp 2 { 3
name[256];
/**& 设备的ASCII标识 */ 4
size_t totalGlobalM
/**& 可用的全局量，单位字节 */ 5
size_t sharedMemPerB
/**& 每个block可用的共享量，单位字节 */ 6
/**& 每个block里可用32位寄存器数量 */ 7
/**& 在线程warp块大小*/ 8
size_t memP
/**& 允许的内存复制最大修正，单位字节*/ 9
maxThreadsPerB
/**& 每个block最大进程数量 */10
maxThreadsDim[3];
/**& 一block里每个维度最大线程量 */11
maxGridSize[3];
/**& 一格里每个维度最大数量 */12
/**& 时钟频率，单位千赫khz */13
size_t totalConstM
/**& 设备上可用的常量内存，单位字节 */14
/**& 计算功能主版本号*/15
/**& 计算功能次版本号*/16
size_t textureA
/**& 对齐要求的纹理 */17
/**& 判断设备是否可以同时拷贝内存和执行内核。已过时。改用asyncEngineCount */18
multiProcessorC
/**& 设备上的处理器数量 */19
kernelExecTimeoutE
/**& 内核函数是否运行受时间限制*/20
/**& 设备是不是独立的 */21
canMapHostM
/**& 设备能否映射主机cudaHostAlloc/cudaHostGetDevicePointer */22
/**& 计算模式，有默认，独占，禁止，独占进程(See ::cudaComputeMode) */23
maxTexture1D;
/**& 1D纹理最大值 */24
maxTexture2D[2];
/**& 2D纹理最大维数*/25
maxTexture3D[3];
/**& 3D纹理最大维数 */26
maxTexture1DLayered[2];
/**& 最大的1D分层纹理尺寸 */27
maxTexture2DLayered[3];
/**& 最大的2D分层纹理尺寸
size_t surfaceA
/**& 表面的对齐要求*/29
concurrentK
/**& 设备是否能同时执行多个内核*/30
/**& 设备是否支持ECC */31
/**& 设备的PCI总线ID */32
pciDeviceID;
/**& PCI设备的设备ID*/33
pciDomainID;
/**&PCI设备的域ID*/34
/**& 1如果设备是使用了TCC驱动的Tesla设备，否则就是0 */35
asyncEngineC
/**& 异步Engine数量 */36
/**& 设备是否共享统一的地址空间与主机*/37
memoryClockR
/**&峰值内存时钟频率，单位khz*/38
memoryBusW
/**& 全局内存总线宽度，单位bit*/39
/**& L2 cache大小，单位字节 */40
maxThreadsPerMultiP/**& 每个多处理器的最大的常驻线程 */41 };　　通过cudaGetDeviceProperties()得到设备属性，cudaGetDeviceCount()来获取设备的个数，通过cudaChooseDevice()选择符合条件的设备，通过cudaGetDevice()可以得到当前的设备，通过cudaSetDevice()设置选择设备，SLI技术支持多个GPU。&&&&&&& 更多内容请点击：&&&&&&& CUDA专区：&&&&&&& CUDA论坛：&
IT168企业级线程的组织形式对程序的性能影响是至关重要的，本篇博文主要以下面一种情况来介绍线程组织形式：
2D grid 2D block
矩阵在memory中是row-major线性存储的：
在kernel里，线程的唯一索引非常有用，为了确定一个线程的索引，我们以2D为例：
线程和block索引
矩阵中元素坐标
线性global memory 的偏移
首先可以将thread和block索引映射到矩阵坐标：
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y
之后可以利用上述变量计算线性地址：
idx = iy * nx + ix
上图展示了block和thread索引，矩阵坐标以及线性地址之间的关系，谨记，相邻的thread拥有连续的threadIdx.x，也就是索引为(0,0)(1,0)(2,0)(3,0)...的thread连续，而不是(0,0)(0,1)(0,2)(0,3)...连续，跟我们线代里玩矩阵的时候不一样。
现在可以验证出下面的关系：
thread_id（2,1）block_id（1,0） coordinate（6,1） global index 14 ival 14
下图显示了三者之间的关系：
int main(int argc, char **argv) {
　　printf("%s Starting...\n", argv[0]);
　　// set up device
　　int dev = 0;
　　cudaDeviceProp deviceP
　　CHECK(cudaGetDeviceProperties(&deviceProp, dev));
　　printf("Using Device %d: %s\n", dev, deviceProp.name);
　　CHECK(cudaSetDevice(dev));　　// set up date size of matrix
　　int nx = 1&&14;
　　int ny = 1&&14;
　　int nxy = nx*
　　int nBytes = nxy * sizeof(float);
　　printf("Matrix size: nx %d ny %d\n",nx, ny);
　　// malloc host memory
　　float *h_A, *h_B, *hostRef, *gpuR
　　h_A = (float *)malloc(nBytes);
　　h_B = (float *)malloc(nBytes);
　　hostRef = (float *)malloc(nBytes);
　　gpuRef = (float *)malloc(nBytes);
　　　　// initialize data at host side
　　double iStart = cpuSecond();
　　initialData (h_A, nxy);
　　initialData (h_B, nxy);
　　double iElaps = cpuSecond() - iS
　　memset(hostRef, 0, nBytes);
　　memset(gpuRef, 0, nBytes);
　　// add matrix at host side for result checks
　　iStart = cpuSecond();
　　sumMatrixOnHost (h_A, h_B, hostRef, nx,ny);
　　iElaps = cpuSecond() - iS
　　// malloc device global memory
　　float *d_MatA, *d_MatB, *d_MatC;
　　cudaMalloc((void **)&d_MatA, nBytes);
　　cudaMalloc((void **)&d_MatB, nBytes);
　　cudaMalloc((void **)&d_MatC, nBytes);
　　　　// transfer data from host to device
　　cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);
　　cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice);
　　// invoke kernel at host side
　　int dimx = 32;
　　int dimy = 32;
　　dim3 block(dimx, dimy);
　　dim3 grid((nx+block.x-1)/block.x, (ny+block.y-1)/block.y);
　　iStart = cpuSecond();
　　sumMatrixOnGPU2D &&& grid, block &&&(d_MatA, d_MatB, d_MatC, nx, ny);
　　cudaDeviceSynchronize();
　　iElaps = cpuSecond() - iS
　　printf("sumMatrixOnGPU2D &&&(%d,%d), (%d,%d)&&& elapsed %f sec\n", grid.x,
　　grid.y, block.x, block.y, iElaps);
　　// copy kernel result back to host side
　　cudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost);
　　// check device results
　　checkResult(hostRef, gpuRef, nxy);
　　　　// free device global memory
　　cudaFree(d_MatA);
　　cudaFree(d_MatB);
　　cudaFree(d_MatC);
　　// free host memory
　　free(h_A);
　　free(h_B);
　　free(hostRef);
　　free(gpuRef);
　　// reset device
　　cudaDeviceReset();
　　return (0);
编译运行：
$ nvcc -arch=sm_20 sumMatrixOnGPU-2D-grid-2D-block.cu -o matrix2D
$ ./matrix2D
./a.out Starting...
Using Device 0: Tesla M2070
Matrix size: nx 16384 ny 16384
sumMatrixOnGPU2D &&&(512,512), (32,32)&&& elapsed 0.060323 sec
Arrays match.
接下来，我们更改block配置为32x16，重新编译，输出为：
sumMatrixOnGPU2D &&&(512,1024), (32,16)&&& elapsed 0.038041 sec
可以看到，性能提升了一倍，直观的来看，我们会认为第二个配置比第一个多了一倍的block所以性能提升一倍，实际上也确实是因为block增加了。但是，如果你继续增加block的数量，则性能又会降低：
sumMatrixOnGPU2D &&& (), (16,16) &&& elapsed 0.045535 sec
下图展示了不同配置的性能;
关于性能的分析将在之后的博文中总结，现在只是了解下，本文在于掌握线程组织的方法。
代码下载：
阅读(...) 评论()CUDA计算时间问题
[问题点数：20分]
CUDA计算时间问题
[问题点数：20分]
不显示删除回复
显示所有回复
显示星级回复
显示得分回复
只显示楼主
匿名用户不能发表回复！|