CUDA Programming – 2D Thread Organization

by
Dr. Nileshchandra Pikle
Assistant Professor
&
“A certified CUDA instructor by NVIDIA”
Thread Organization Extended to 2D

● Threads can be organized in 2D or 3D

● dim3 is an integer vector type that can be used in
CUDA code.
● Its most common application is to pass the grid and
block dimensions in a kernel invocation.
● Eg.
dim3
dim3 grid(x,
grid(x, y,
y, z)
z)

Grid is a vector of 3 dimension and of type dim3

CUDA Programming by N K Pikle 2 / 31

Thread Organization Extended to 2D

● Threads can be organized in 2D or 3D

● dim3 is an integer vector type that can be used in
CUDA code.
● Its most common application is to pass the grid and
block dimensions in a kernel invocation.
● Eg.
dim3 grid( 512 ); // 512 x 1 x 1
dim3 block( 1024, 1024 ); // 1024 x 1024 x 1
fooKernel<<< grid, block >>>();

CUDA Programming by N K Pikle 3 / 31

 Thread Organization

● 1D grid and 1D block

int
int gid
gid =
= blockIdx.x
blockIdx.x *blockDim.x
*blockDim.x +
+ threadIdx.x;
threadIdx.x;

gid 0 1 2 4 5 6 7 8 9 10 11 12

ThreadIdx.x 0 1 2 3 0 1 2 3 0 1 2 3

BlockDim.x = 4 BlockDim.x = 4 BlockDim.x = 4

BlockIdx.x = 0 BlockIdx.x = 1 BlockIdx.x = 2

CUDA Programming by N K Pikle 4 / 31

 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

blockDim.x = 2 blockDim.x = 2 blockDim.x = 2

threadIdx.x

int
int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y
threadIdx.y

T(0,0) T(1,0) .y
threadIdx.y
threadIdx.y ** blockDim
blockDim.x
.x +
+ threadIdx.x
threadIdx.x;;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 5 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 BlockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

blockDim.x = 2 blockDim.x = 2 blockDim.x = 2

threadIdx.x

int
int gid
gid =
= blockIdx.x
blockIdx.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y
threadIdx.y

T(0,0) T(1,0) .y
threadIdx.y
threadIdx.y ** blockDim
blockDim.x
.x +
+ threadIdx.x
threadIdx.x;;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 6 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

blockDim.x = 2 blockDim.x = 2 blockDim.x = 2

ThreadIdx.x
ThreadIdx.y

T(0,0) T(1,0) int

int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim.x
blockDim.x *blockDim.y
*blockDim.y
threadIdx.y
threadIdx.y ** blockDim
blockDim.x
.x +
+ threadIdx.x
threadIdx.x;;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 7 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

BlockDim.x = 2 blockDim.x = 2 blockDim.x = 2

threadIdx.x
ThreadIdx.y

T(0,0) T(1,0) int

int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y.y
threadIdx.y
threadIdx.y ** blockDim.x
blockDim.x ++ threadIdx.x
threadIdx.x;;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 8 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

blockDim.x = 2 blockDim.x = 2 blockDim.x = 2

ThreadIdx.x

int
int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y
threadIdx.y

T(0,0) T(1,0) .y
threadIdx
threadIdx.y.y ** blockDim
blockDim.x
.x +
+ threadIdx.x;
threadIdx.x;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 9 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

BlockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

This computes the

Cooresponding block
blockDim.x = 2 blockDim.x = 2 BlockDim
[Link]
= 2thread

ThreadIdx.x

int
int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y
threadIdx.y

T(0,0) T(1,0) .y
threadIdx
threadIdx.y.y ** blockDim
blockDim.x
.x +
+ threadIdx
threadIdx.x;
.x;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 10 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

This computes the

Cooresponding y
blockDim.x = 2 blockDim.x = 2 Coordinate
BlockDim.x = 2 in a
Thread block
ThreadIdx.x

int
int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y
threadIdx.y

T(0,0) T(1,0) .y
threadIdx
threadIdx.y.y ** blockDim
blockDim.x
.x +
+ threadIdx
threadIdx.x;
.x;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 11 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

T(0,0) T(1,0) T(0,0) T(1,0) T(0,0) T(1,0)

blockDim.y

Grid
Grid
T(0,1) T(1,1) T(0,1) T(1,1) T(0,1) T(1,1)

This computes the

Cooresponding x
blockDim.x = 2 blockDim.x = 2 Coordinate
BlockDim.x = 2 in a
Thread block
ThreadIdx.x

int
int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y
threadIdx.y

T(0,0) T(1,0) .y
threadIdx
threadIdx.y.y ** blockDim
blockDim.x
.x +
+ threadIdx
threadIdx.x;
.x;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 12 / 31
 Thread Organization
● 1D grid and 2D block
.y = 2 blockIdx.x = 0 blockIdx.x = 1 blockIdx.x = 2

0 1 4 5 8 9
blockDim.y

Grid
Grid
2 3 6 7 10 11

blockDim.x = 2 blockDim.x = 2 blockDim.x = 2

ThreadIdx.x
ThreadIdx.y

T(0,0) T(1,0) int

int gid
gid =
= blockIdx
blockIdx.x
.x ** blockDim
blockDim.x
.x *blockDim
*blockDim.y.y
threadIdx
threadIdx.y.y ** blockDim
blockDim.x
.x +
+ threadIdx
threadIdx.x;
.x;

T(0,1) T(1,1)
CUDA Programming by N K Pikle 13 / 31
Thread Organization 2D block & 2D grid
blockIdx.x

Block(0,0) Block(1,0)

threadIdx .y
dIdx.y T(0,0) T(1,0) T(0,0) T(1,0)
.y
x.y

T(0,1) T(1,1) T(0,1) T(1,1)

bloc kIdx

threadIdx.x

T(0,0) T(1,0) T(0,0) T(1,0)

T(0,1) T(1,1) T(0,1) T(1,1)

Block(0,1) Block(1,1)

int
int row
row == threadIdx.y
threadIdx.y +
+ blockIdx.y
blockIdx.y ** blockDim.y
blockDim.y
int
int col
col =
= threadIdx.x
threadIdx.x +
+ blockIdx.x
blockIdx.x ** blockDim.x
blockDim.x

CUDA Programming by N K Pikle 14 / 31

Thread Organization 2D block & 2D grid
blockIdx.x

Block(0,0) Block(1,0)

.y
dIdx.y
threadIdx (0,0) (0,1) (0,2) (0,3)
.y
x.y

(1,0) (1,1) (1,2) (1,3)

bloc kIdx

threadIdx.x

(2,0) (2,1) (2,2) (2,3)

(3,0) (3,1) (3,2) (3,3)

Block(0,1) Block(1,1)

int
int row
row == threadIdx.y
threadIdx.y +
+ blockIdx.y
blockIdx.y ** blockDim.y
blockDim.y
int
int col
col =
= threadIdx.x
threadIdx.x +
+ blockIdx.x
blockIdx.x ** blockDim.x
blockDim.x

CUDA Programming by N K Pikle 15 / 31

Thread Organization 2D block & 2D grid
blockIdx.x

Block(0,0) Block(1,0)

.y
dIdx.y
threadIdx 0 1 2 3
.y
x.y

4 5 6 7
bloc kIdx

threadIdx.x

8 9 10 11

12 13 14 15

Block(0,1) Block(1,1)

int
int gid
gid =
= row
row ** blockDim.x
blockDim.x ** gridDim.x
gridDim.x +
+ col
col

CUDA Programming by N K Pikle 16 / 31

Program 2D convolution Parallel
1 1 1

Convolution mask 1 1 1

1 1 1

1 1 1 1 1

1 1 1 1 1

Input Image 1 1 1 1 1

1 1 1 1 1

1 1 1 1 1

CUDA Programming by N K Pikle 17 / 31

Program 2D convolution Parallel

To
To take
take care
care of
of boundary
boundary conditions
conditions either
either
apply
apply zeroPadding
zeroPadding or
or avoid
avoid the
the out
out of
of
range
range computations
computations

0 0 0 0 0 0 0
0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 18 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 4 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 19 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 6 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 20 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 6 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 21 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 6 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 22 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 1 4 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 23 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 1 1 0

0 6 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 24 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 1 1 0

0 1 9 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 25 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 1 1 0

0 1 1 9 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 26 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 1 1 0

0 1 1 1 9 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 27 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

0 0 0 0 0 0 0
0 1 1 1 1 1 0

0 1 1 1 1 6 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 1 1 1 1 1 0

0 0 0 0 0 0 0
CUDA Programming by N K Pikle 28 / 31
 Program 2D convolution Parallel
1 1 1
Convolution +a +a

mask 1 1 1 O[i ][ j ] = ∑ ∑ mask [ s][t ]∗A [i+ s][ j+t ]

s=−a t =−a

1 1 1

4 6 6 6 4

6 9 9 9 6

Final Output 6 9 9 9 6
Image
6 9 9 9 6

4 6 6 6 4

CUDA Programming by N K Pikle 29 / 31

Program 2D convolution Parallel

Refer
Refer program
program 2D_Convolution.cu
2D_Convolution.cu

CUDA Programming by N K Pikle 30 / 31

Program 2D convolution Parallel

● Summary
1. CUDA programming model – Heterogeneous
Computing
2. CUDA Thread organization: Threads grouped
together to form blocks and blocks to grid.
-Threads can be organized as 1D, 2D or 3D.
3. Case studies: To understand basic CUDA
programming model, thread and data mapping in
1D and 2D

CUDA Programming by N K Pikle 31 / 31