CUDA : Global unique thread index in a 3D Grid

Question

As the question states, If I have a 3D Grid of Blocks what is the formula to get a global unique index for one thread ?

Lets keep the block itself as one dimensional.

Answer 1

// unique block index inside a 3D block grid
const unsigned long long int blockId = blockIdx.x //1D
        + blockIdx.y * gridDim.x //2D
        + gridDim.x * gridDim.y * blockIdx.z; //3D

// global unique thread index, block dimension uses only x-coordinate
const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x;

Answer 2

A bit late to the party but here is how I usually approach this in a pretty generic way in that it supports any number and size of block (even 2D):

// Compute the offset in each dimension
const size_t offsetX = blockDim.x * blockIdx.x + threadIdx.x;
const size_t offsetY = blockDim.y * blockIdx.y + threadIdx.y;
const size_t offsetZ = blockDim.z * blockIdx.z + threadIdx.z;

// Make sure that you are not actually outs
if (offsetX >= sizeX || offsetY >= sizeY || offsetZ >= sizeZ)
  return;

// Compute the linear index assuming that X,Y then Z memory ordering
const size_t idx = offsetZ * sizeX * sizeY + offsetY * sizeX + offsetX;

Mind you that I'm not a CUDA ninja.