电子科技大学：《GPU并行编程 GPU Parallel Programming》课程教学资源（课件讲稿）Lecture 02 CUDA PARALLELISM MODEL

GPU Teachina Kit LECTURE 2-CUDA PARALLELISM MODEL Multidimensional Kernel Configuration

LECTURE 2 – CUDA PARALLELISM MODEL Multidimensional Kernel Configuration GPU Teaching Kit

Multidimensional Kernel Configuration Color-to-Greyscale Image Processing Example Blur Image Processing Example Universityf Electr Science and TachnoloChina

Multidimensional Kernel Configuration Color-to-Greyscale Image Processing Example Blur Image Processing Example

OBJECTIVE .To understand multidimensional Grids -Multi-dimensional block and thread indices Mapping block/thread indices to data indices 2 电子料发女学 Universityof Electr Science and TachnolopChina O

OBJECTIVE ▪ To understand multidimensional Grids ▪ Multi-dimensional block and thread indices ▪ Mapping block/thread indices to data indices 2

threadldx.x Grid ack{0,0) Block (1.0) B1ock(2,0) threadldx.y 0ck(0.1) Slock (1.1) B0ck(2.1) blockldx.x Block (1.1) blockldx.y blockDimy threadldxy Thread(0.01 Thre3d情，0 Thread(2.0)Tnreod (3.0) blockldx.y blockddx.x blockDim.x+threadidx.x Thread(21)Thread3.1) Th302hrd121 Thread(2.2)Thread (3.2) 电子料皮女学 University of Electraaie Science and Technolory of China O

A MULTI-DIMENSIONAL GRID EXAMPLE host device Grid 1 Block Block (0,0) (0,1) Kernel l Block Block (1,0) (1,1) Block(1,0) .0,0 1.0,1) (1.0,2) 人1,0,3) Thread Thread Thread Thread (0,0,0) (0,0,1) (0.0,2) (0.0,31 Threrd Thread Thread Thread 0,1,0 0,1,1) 01,2) (0,1,3) 电子科烛女学 niversitof Electr Science and TachnoloChina O

host device Kernel 1 Grid 1 Block (0, 0) Block (1, 1) Block (1, 0) Block (0, 1) Grid 2 Block (1,0) Thread Thread (0,0,0) (0,1,3) Thread (0,1,0) Thread (0,1,1) Thread (0,1,2) Thread (0,0,0) Thread (0,0,1) Thread (0,0,2) Thread (0,0,3) (1,0,0) (1,0,1) (1,0,2) (1,0,3) A MULTI-DIMENSIONAL GRID EXAMPLE 5

PROCESSING A PICTURE WITH A 2D GRID 16×16 blocks 62x76 picture 电子料线女学 University of Electreaie Science and Technolory of China O

PROCESSING A PICTURE WITH A 2D GRID

ROW-MAJOR LAYOUT IN C/C++ M Row*Vidth+Col=2*4+1≡9■ M Mi M2 M3 MM 5M6 M Ms Mo M10 Mi1 M12 Mi3 M Mi M Moo Mo1 Mo2 Mos Mio M11M12 Mis M20 M21 M22 M23 M3o M31 M32 M3s Moo Mor Mo2Mo3 Mio M12M13 M20M2,1M2.2M2.3 M30M3,M32M3,3 电子料做女学 Universityof ElectriScience and TachnolofChina O

M0,2 M1,1 M0,0 M0,1 M1,0 M0,3 M1,2 M1,3 M0,0 M0,1 M0,2 M0,3 M1,0 M1,1 M1,2 M1,3 M2,0 M2,1 M2,2 M2,3 M2,0 M2,1 M2,2 M2,3 M3,0 M3,1 M3,2 M3,3 M3,0 M3,1 M3,2 M3,3 M Row*Width+Col = 2*4+1 = 9 M0 M1 M2 M3 M4 M5 M6 M7 M8 M9 M10 M11 M12 M13 M14 M15 M ROW-MAJOR LAYOUT IN C/C++

SOURCE CODE OF A PICTUREKERNEL global void PictureKernel(float*d Pin,float*d Pout, int height,int width) { /Calculate the row of the d pin and d Pout element int Row blockIdx.y*blockDim.y threadIdx.y; /Calculate the column of the d pin and d Pout element int Col blockIdx.x*blockDim.x threadIdx.x; /each thread computes one element of d Pout if in range if ((Row height)&&(Col width)){ d Pout [Row*width+Col]2.0*d Pin[Row*width+Col]; 电子料发女学 University of Electreaie Science and Technolory of China O

SOURCE CODE OF A PICTUREKERNEL __global__ void PictureKernel(float* d_Pin, float* d_Pout, int height, int width) { // Calculate the row # of the d_Pin and d_Pout element int Row = blockIdx.y*blockDim.y + threadIdx.y; // Calculate the column # of the d_Pin and d_Pout element int Col = blockIdx.x*blockDim.x + threadIdx.x; // each thread computes one element of d_Pout if in range if ((Row < height) && (Col < width)) { d_Pout[Row*width+Col] = 2.0*d_Pin[Row*width+Col]; } } Scale every pixel value by 2.0

HOST CODE FOR LAUNCHING PICTUREKERNEL /assume that the picture is m x n, //m pixels in y dimension and n pixels in x dimension /input d_Pin has been allocated on and copied to device /output d Pout has been allocated on device dim3 DimGrid(n-1)/16+1,(m-1)/16+1,1); dim3 DimBlock(16,16,1); PictureKernel>>(d_Pin, d_Pout,m,n); 电子料效女学 University of Electricience and TachnolopChina O

HOST CODE FOR LAUNCHING PICTUREKERNEL // assume that the picture is m × n, // m pixels in y dimension and n pixels in x dimension // input d_Pin has been allocated on and copied to device // output d_Pout has been allocated on device … dim3 DimGrid((n-1)/16 + 1, (m-1)/16+1, 1); dim3 DimBlock(16, 16, 1); PictureKernel>>(d_Pin, d_Pout, m, n); …

COVERING A62×I6 PICTURE WITH16×16 BLOCKS 16×16 block Not all threads in a Block will follow the same control flow path. 电子料线女学 Universityof ElectriScience and TachnolopChina O

COVERING A 62×76 PICTURE WITH 16×16 BLOCKS Not all threads in a Block will follow the same control flow path