how to determine current warp / half-warp?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

Say i have a kernel with a 2 dimensional block grid:

THREAD_w = 8
THREAD_h = 16

dim3 block(THREAD_w, THREAD_h, 1);

This gives me 2x16 = 128 threads per block.

I’d like to figure out from threadIdx which warp or half-warp i’m in,
in order to optimize my code by selecting random branching to alternate but follow the same paths inside each warp.

The manual is unclear as to how this is indexed for 2d grids like my example above.

Can someone clarify ?

Thank you,
Sando

2

please see page 9 in programming guide 2.3

__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];

}

you can determine warp_id as

__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;

	

	int threadID = j * N + i;

	int warpID = threadID / 32;

	 

  ....

}
3

LSChien, you shouldn’t hard code the warp size when determining a value, as that could (and probably will) introduce errors into your code on future devices. Just use the built-in variable “warpSize”:

int warpId = threadId / warpSize;