Convolution Kernel Practice

What distinguishes Strategy 1?

S1: block covers OUTPUT tile, threads also load the halo (multiple loads/thread)

What distinguishes Strategy 2?

S2: block covers INPUT tile, 1 load/thread, edge threads idle during compute.

What distinguishes Strategy 3?

3: block covers OUTPUT tile, loads only interior; halo read straight from global (cache).

What is strategy 1s block coverage/ sizing rules ?

S1 = output tile + halo loops and shared = output + (MW−1).

What is strategy 2s block coverage/ sizing rules?

block covers the input tile, and output = shared − (MW−1).

What is strategy 3s block coverage/sizing rules?

block covers the output tile, only the interior goes in shared memory, and the halo comes from global memory.

How much output does a block cover along one axis?

threads along axis × outputs per thread, for example BW * TTW. If a thread makes 1 output, it's just the output tile width.

For strategy 1/2, using the info given, how do you find dim3 gridDim?

gridDim is how many blocks it takes to cover the whole output per dimension. Since block size = output tile, and the shared memory = full input tile,

ceil(total OUTPUT size/ output per block)

Put the whole denominator in parentheses and use 1.0 * or (n + d - 1) / d.

Output size from input size with an MW mask?

input - MW + 1

How big is the input tile for a given output tile?

output tile + MW - 1, per axis.

How do you flatten and unflatten (row, col)?

Flatten: row * row_width + col. Unflatten: / row_width gives the row, % row_width gives the column.

Check: the first thread gives 0, the last gives count - 1.

Where does threadIdx appear in index math?

Only at the thread level. Block start is blockIdx * coverage. Thread start is threadIdx * per-thread tile. Element offset is element_x/y. A tile-relative coordinate plus the block start gives the absolute coordinate.

Shape of a strided loop?

for (i = my_id; i < total; i += num_threads)

Which coordinates index input and which index shared_tile?

input uses absolute coordinates. shared_tile uses tile-relative coordinates.

Which bounds checks do you need in "valid" convolution?

Only upper bounds (< size). Nothing starts left of or above the block, so there's no >= 0 check. Contrast with your 3D lab, where the halo reached left of the block.

S3-1. What defines Strategy 3?

The block covers the output tile, one thread per output. Only the "core" goes into shared memory, one load per thread. Halo cells are read straight from global memory.

Three-way test for each neighbor.

For every neighbor position, check in this order:

1. Outside the image (ghost cell): skip it, and here also reduce the divisor.

2. Inside the block's tile: read from tile[...].

3. Otherwise (a halo cell in a neighboring block's tile): read from in[...] in global memory.

The order matters. The image-bounds check comes first, and the tile check only means something for in-bounds positions.

Tile-relative vs global coordinates.

Global coordinates index in. Tile-relative coordinates index tile. To convert, subtract the block's start: z - bz * TILE_WIDTH. This is the same "absolute vs relative" pattern as P8 from the earlier cards.

Why no early return for out-of-image threads?

Every thread in the block must reach __syncthreads(). Guard the load and the write instead of returning.