Индексы трехмерных изображений

Question

У меня есть изображение размером 512 х 512 х 512. Мне нужно обработать все воксели индивидуально. Как я могу получить идентификатор потока, чтобы сделать это? Если я использую 1D идентификатор потока, количество блоков превысит 65536.

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

Примечание: - Моя карта не поддерживает 3D-сетки

Answer 1

Вы можете использовать трехмерные индикаторы в CUDA 4.0 и вычислительные возможности 2.0+. Пример кода:

int blocksInX = (nx+8-1)/8;
int blocksInY = (ny+8-1)/8;
int blocksInZ = (nz+8-1)/8;

dim3 Dg(blocksInX, blocksInY, blocksInZ);
dim3 Db(8, 8, 8);
foo_kernel<<Dg, Db>>(R, nx, ny, nz);

...

__global__ void foo_kernel( float* R, const int nx, const int ny, const int nz )
{
  unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
  unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
  unsigned int zIndex = blockDim.z * blockIdx.z + threadIdx.z;

  if ( (xIndex < nx) && (yIndex < ny) && (zIndex < nz) )
  {
    unsigned int index_out = xIndex + nx*yIndex + nx*ny*zIndex;
    ...
    R[index_out] = ...;
  }
}

Если ваше устройство не поддерживает вычислительные возможности 2.0, есть несколько хитростей:

int threadsInX = 16;
int threadsInY = 4;
int threadsInZ = 4;

int blocksInX = (nx+threadsInX-1)/threadsInX;
int blocksInY = (ny+threadsInY-1)/threadsInY;
int blocksInZ = (nz+threadsInZ-1)/threadsInZ;

dim3 Dg = dim3(blocksInX, blocksInY*blocksInZ);
dim3 Db = dim3(threadsInX, threadsInY, threadsInZ);

foo_kernel<<<Dg, Db>>>(R, nx, ny, nz, blocksInY, 1.0f/(float)blocksInY);

__global__ void foo_kernel(float *R, const int nx, const int ny, const int nz,
                           unsigned int blocksInY, float invBlocksInY)
{

    unsigned int blockIdxz = __float2uint_rd(blockIdx.y * invBlocksInY);
    unsigned int blockIdxy = blockIdx.y - __umul24(blockIdxz, blocksInY);
    unsigned int xIndex = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
    unsigned int yIndex = __umul24(blockIdxy, blockDim.y) + threadIdx.y;
    unsigned int zIndex = __umul24(blockIdxz, blockDim.z) + threadIdx.z;

    if ( (xIndex < nx) && (yIndex < xIndex) && (zIndex < nz) )
    {
        unsigned int index = xIndex + nx*yIndex + nx*ny*zIndex;
        ...
        R[index] = ...;
    }

}

Answer 2

Обратите внимание, что память вашего компьютера не в 3D.Это всего лишь вопрос визуализации, поэтому вы можете конвертировать свое 3D-изображение в один указатель.

Array[i][j][z] is same as Array2[ i*cols+j + rows*cols*z];

Теперь передайте Array2 в CUDA и работайте в одном измерении

Answer 3

Вы можете использовать сетки.Это дает вам гораздо больше индексов.

Answer 4

Я использовал что-то вроде этого:

В коде определите вашу сетку: dim3 altgrid, altthreads; altgrid.x = лк; altgrid.y = LY; altgrid.z = 1; altthreads.x = LZ; altthreads.y = 1; altthreads.z ​​= 1;

и в ядре

int idx = threadIdx.x;
int idy = blockIdx.x ;
int idz = blockIdx.y ;

Поскольку массив на устройстве имеет только 1D, вы извлекаете элемент [idx] [idy] [idz] из матрицы A как A [ind], где ind = idz + lz * (idy + ly * idx) ;

Надеюсь, это поможет

Answer 5

Если вам нужна сетка большего размера, CUDA поддерживает 2D-сетки на всем оборудовании, а самые последние версии набора инструментов CUDA также поддерживают 3D-сетки на текущем оборудовании Fermi.

Однако не обязательно иметь такие большие сетки.Если каждая операция вокселя независима, то почему бы просто не использовать одномерную сетку, а каждый поток обрабатывать более одного вокселя?Мало того, что такая схема не требует больших 2D или 3D сеток, она также может быть более эффективной, потому что постоянные затраты, связанные с планированием и инициализацией блока, могут амортизироваться при множественных вычислениях вокселей.