cuda のid計算がよくわからないので、まとめてみる。

get_global_id.cu

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

namespace comment_thread_id
{
	/*
		blockIdx.x == n
		blockIdx.y == 1
		blockIdx.z == 1
		
		のときを考える。これを 1D grid と呼ぶことにする。
	 */
	//1D grid, 3D block
	__device__ int getGlobalIdx_1D_3D()
	{
		return     blockIdx.x * blockDim.x * blockDim.y * blockDim.z
			// x y zの3次元を1次元に直す。zは xとyの重みが、yはxの重みが、xは1の重みがあると考える。
			+ threadIdx.z * blockDim.y * blockDim.x 
			+ threadIdx.y * blockDim.x
			+ threadIdx.x;
	}

	//1D grid, 2D block
	__device__ int getGlobalIdx_1D_2D()
	{
		return	blockIdx.x * blockDim.x * blockDim.y * 1 /*blockDim.z*/
			// x yの2次元を1次元に直す。yはxの重みが、xは1の重みがあると考える。
			+ 0 /*threadIdx.z*/ * blockDim.y * blockDim.x
			+ threadIdx.y * blockDim.x
			+ threadIdx.x;
	}

	//1D grid, 1D block
	__device__ int getGlobalIdx_1D_1D()
	{
		return	blockIdx.x * blockDim.x * 1/* blockDim.y */ * 1 /*blockDim.z*/
			// x の1次元なので直す必要がないが、xは1の重みがあると考えられる。
			+ 0 /*threadIdx.z*/ * 1 /*blockDim.y */ * blockDim.x
			+ 0 /*threadIdx.y*/ * blockDim.x
			+ threadIdx.x;
	}


	/*
		ここまでのパターンを
	 
		blockIdx.x == n
		blockIdx.y == m
		blockIdx.z == n

		のとき（３D grid）に拡張する。
	*/

	//3D grid, 3D blocks
	__device__ int getGlobalIdx_3D_3D()
	{
		/*
		   gridDim = {1,1,1} のときは

		   __device__ int getGlobalIdx_1D_3D()
			{
				return     blockIdx.x * blockDim.x * blockDim.y * blockDim.z
				+ threadIdx.z * blockDim.y * blockDim.x 
				+ threadIdx.y * blockDim.x 
				+ threadIdx.x;
			}

			だった。
			上記は、 block_id == blockIdx.x である特別な場合であり以下のコードの略であると考える。
			{
				int block_id == blockIdx.x;
			 
				return block_id * (blockDim.x * blockDim.y * blockDim.z)
				+ threadIdx.z * blockDim.y * blockDim.x 
				+ threadIdx.y * blockDim.x 
				+ threadIdx.x;
			}			
		 */

		// x y zの3次元を1次元に直す。zは xとyの重みが、yはxの重みが、xは1の重みがあると考える。
		int block_id =    blockIdx.z * (gridDim.x * gridDim.y)
				+ blockIdx.y * (gridDim.x)
				+ blockIdx.x;
		int threadId = block_id * (blockDim.x * blockDim.y * blockDim.z)
			+ (threadIdx.z * (blockDim.x * blockDim.y))
			+ (threadIdx.y * blockDim.x)
			+ threadIdx.x;
		return threadId;
	}
}

任意1..3の次元数でidを求めるためのすべての関数は、最初のコミット
https://gist.github.com/onionmk2/854c333829f047a5e86cfab5a0ccae3a/7bcd4a696f1f08319ea0564975ed962ce61a484d
に記載がある。