哪些网站平台可以做推广,怎样在百度免费做网站,溧水做网站,建设一个网站平台需要哪些技术员本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/#xff0c;仅用于学习。
前言
线程的组织形式对程序的性能影响是至关重要的#xff0c;本篇博文主要以下面一种情况来介绍线程组织形式#xff1a;
2D grid 2D block
线程索引
矩阵在memory中是row-major线性…本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/仅用于学习。
前言
线程的组织形式对程序的性能影响是至关重要的本篇博文主要以下面一种情况来介绍线程组织形式
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_id2,1block_id1,0 coordinate6,1 global index 14 ival 14
下图显示了三者之间的关系 代码 int main(int argc, char **argv) {printf(%s Starting...\n, argv[0]);// set up deviceint dev 0;cudaDeviceProp deviceProp;CHECK(cudaGetDeviceProperties(deviceProp, dev));printf(Using Device %d: %s\n, dev, deviceProp.name);CHECK(cudaSetDevice(dev)); // set up date size of matrixint nx 114;int ny 114;int nxy nx*ny;int nBytes nxy * sizeof(float);printf(Matrix size: nx %d ny %d\n,nx, ny);// malloc host memoryfloat *h_A, *h_B, *hostRef, *gpuRef;h_A (float *)malloc(nBytes);h_B (float *)malloc(nBytes);hostRef (float *)malloc(nBytes);gpuRef (float *)malloc(nBytes);// initialize data at host sidedouble iStart cpuSecond();initialData (h_A, nxy);initialData (h_B, nxy);double iElaps cpuSecond() - iStart;memset(hostRef, 0, nBytes);memset(gpuRef, 0, nBytes);// add matrix at host side for result checksiStart cpuSecond();sumMatrixOnHost (h_A, h_B, hostRef, nx,ny);iElaps cpuSecond() - iStart;// malloc device global memoryfloat *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 devicecudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice);// invoke kernel at host sideint dimx 32;int dimy 32;dim3 block(dimx, dimy);dim3 grid((nxblock.x-1)/block.x, (nyblock.y-1)/block.y);iStart cpuSecond();sumMatrixOnGPU2D grid, block (d_MatA, d_MatB, d_MatC, nx, ny);cudaDeviceSynchronize();iElaps cpuSecond() - iStart;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 sidecudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost);// check device resultscheckResult(hostRef, gpuRef, nxy);// free device global memorycudaFree(d_MatA);cudaFree(d_MatB);cudaFree(d_MatC);// free host memoryfree(h_A);free(h_B);free(hostRef);free(gpuRef);// reset devicecudaDeviceReset();return (0);
}
编译运行
$ nvcc -archsm_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 (1024,1024), (16,16) elapsed 0.045535 sec
下图展示了不同配置的性能; 关于性能的分析将在之后的博文中总结现在只是了解下本文在于掌握线程组织的方法。
