Skip to content

Instantly share code, notes, and snippets.

Embed
What would you like to do?
cuda のid計算がよくわからないので、まとめてみる。
#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
のとき(3D 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;
}
}
@onionmk2

This comment has been minimized.

Copy link
Owner Author

onionmk2 commented Jan 5, 2018

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

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
You can’t perform that action at this time.