Warp layout in a 2D thread block?

[lhw1]

On a Fermi GPU, if each thread block has 16x16 threads, can anyone tell me how the 32 threads of a warp will be distributed? Will each warp cover two adjacent rows of 16 threads? That would seem logical, but I haven’t been able to find a definitive answer anywhere.

Further to this, obviously it will depend on the application somewhat, but in general on a Fermi GPU should blocks that are 32 threads wide give the best performance?

Thanks in advance.

[Skybuck]

16x16 = 256 threads per block, so occupancy is 100% per multi processor.

My guess is 16x16 is simply converted to 256x1, but this is indeed undocumented, perhaps because they might change how the hardware works in the future.

As long as 1536/256 = 6 blocks are present as a minimum (per multi-processor) then it should give max performance.

(at least 1536 is maximum ammount of threads per multi processor for my gt 520 gpu, what does yours say ? External Image)

[spadflyer12]

According to the programming guide, it goes by x_index first, then y_index, then z_index. For the purposes of warp grouping threads don’t have 3 dimensional indices, they just go by 1. This index is given by threadId = threadIdx.x+blockDim.x*(threadIdx.y+blockDim.y*threadIdx.z). Every 32 threads of this index is a new warp.

[Skybuck]

These formula’s are not in “CUDA C Programming Guide Version 4.0”, if you believe otherwise please state which section ! External Image :)

I have seen one little formula in the guide though, for just 2 dimensions.

Even your formula is still missing the grid.

None-the-less thanks for the formula… it seems the shortest one so far.

I still have to test it and make sure it’s valid, but it seems valid to me External Image

[kbam]

Main reason I can see for needing this is to ensure contiguous memory access, Section 5.3.2.1.2 of the programming guide has this

5.3.2.1.2 Two-Dimensional Arrays

A common global memory access pattern is when each thread of index (tx,ty) uses the following address to access one element of a 2D array of width width, located at address BaseAddress of type type* (where type meets the requirement described in Section 5.3.2.1.1):

BaseAddress + width * ty + tx

For these accesses to be fully coalesced, both the width of the thread block and the width of the array must be a multiple of the warp size (or only half the warp size for devices of compute capability 1.x).

In particular, this means that an array whose width is not a multiple of this size will be accessed much more efficiently if it is actually allocated with a width rounded up to the closest multiple of this size and its rows padded accordingly. The cudaMallocPitch() and cuMemAllocPitch() functions and associated memory copy functions described in the reference manual enable programmers to write non-hardware-dependent code to allocate arrays that conform to these constraints.

[kbam]

Found this which describes it precisely

2.2 Thread Hierarchy

For convenience, threadIdx is a 3-component vector, so that threads can be identified using a one-dimensional, two-dimensional, or three-dimensional thread index, forming a one-dimensional, two-dimensional, or three-dimensional thread block. This provides a natural way to invoke computation across the elements in a domain such as a vector, matrix, or volume.

The index of a thread and its thread ID relate to each other in a straightforward way:

For a one-dimensional block,

they are the same;

for a two-dimensional block of size (Dx, Dy),

the thread ID of a thread of index (x, y) is (x + y Dx);

for a three-dimensional block of size (Dx, Dy, Dz),

the thread ID of a thread of index (x, y, z) is (x + y Dx + z Dx Dy).

and further down

global void MatAdd(float A[N][N], float B[N][N], float C[N][N])

{

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

int j = blockIdx.y * blockDim.y + threadIdx.y;

if (i < N && j < N)

C[i][j] = A[i][j] + B[i][j];

}

[Skybuck]

This is not the same formula as spadflyer12.

His formula moves the dimensions outside of the parenthesis ( ).

So his appears to be more efficient.