выяснить, сколько блоков и потоков для ядра cuda и как их использовать

Question

Я пытался выяснить, как сделать то, что, как я думал, было бы простым ядром для вычисления средних значений в матрице 2d, но у меня возникают некоторые проблемы, связанные с моим мыслительным процессом.

Согласно моему выпуску deviceQuery, мой GPU имеет 16MP, 32cores/mp, max max - 1024x1024x64, а максимальный поток/block = 1024.

Итак, я работаю над обработкой некоторых больших изображений. Может, 5000px x 3500px или что-то в этом роде. Одно из моих ядер принимает среднее значение некоторых значений во всех пикселях изображения.

Существующий код имеет изображения, хранящиеся в виде 2D-массива [rows] [cols]. Таким образом, ядро ​​в C выглядит так, как вы ожидали бы, с циклом над строками и циклом над cols, с вычислением в середине.

Итак, как мне настроить часть вычисления размера этого кода в CUDA? Я посмотрел на код сокращения внутри SDK, но это для одного массива измерений. В нем нет упоминаний о том, как настроить количество блоков и потоков, если у вас есть soemthing 2D.

Я имею в виду, я бы на самом деле нужно, чтобы настроить его, как это так, и это, где я хотел бы кого-то, чтобы звонить в и помощь:

num_threads=1024; 
blocksX = num_cols/sqrt(num_threads); 
blocksY = num_rows/sqrt(num_threads); 
num_blocks = (num_rows*num_cols)/(blocksX*blocksY); 

dim3 dimBlock(blocksX, blocksY, 1); 
dim3 dimGrid(num_blocks, 1, 1); 

Кажется ли это имеет смысл для установки ?

И затем в ядре, чтобы работать на определенной строке или столбце, я должен был бы использовать

rowidx = (blockIdx.x * blockDim.x) + threadId.x colidx = (blockIdx. y * blockDim.y) + threadId.y

По крайней мере, я думаю, что это будет работать для получения строки и столбца.

Как я мог бы получить доступ к этой конкретной строке r и столбцу c в ядре? В руководстве по программированию CUDA я нашел следующий код:

// Host code int width = 64, height = 64; 
float* devPtr; size_t pitch; 
cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height); 
MyKernel<<<100, 512>>>(devPtr, pitch, width, height); 
// Device code __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height) 
{ 
for (int r = 0; r < height; ++r) 
{ 
float* row = (float*)((char*)devPtr + r * pitch); 
for (int c = 0; c < width; ++c) 
{ 
float element = row[c]; 
} 
} 
} 

который похож на то, как вы будете использовать таНос в C, чтобы объявить 2D массива, но он не имеет никакого упоминания о доступе к массиву в собственном ядре , Думаю, в моем коде я буду использовать этот вызов cudaMallocPitch, а затем выполнить memcpy, чтобы получить мои данные в 2D-массив на устройстве?

Любые советы оценены! Благодаря!

Answer 1

Ниже приведен короткий фрагмент с простым ядром из моего собственного кода. Указатели с плавающей точкой - это все указатели на устройства. Надеюсь, это полезно.

определяет и помощь функции: расчет

#define BLOCK_SIZE 16 

int iDivUp(int a, int b){ 
    return (a % b != 0) ? (a/b + 1) : (a/b); 
} 

Размер блока:

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
dim3 dimGridProj(iDivUp(width,BLOCK_SIZE), iDivUp(height,BLOCK_SIZE)); 

Хост вызов:

calc_residual<<<dimGridProj, dimBlock>>>(d_image1, d_proj1, d_raynorm1, d_resid1, width, height); 

Ядро:

__global__ void calc_residual(float *d_imagep, float *d_projp, float *d_raysump, float *d_residualp, int width, int height) 
{ 
    int iy = blockDim.y * blockIdx.y + threadIdx.y; 
if (iy >= height) { 
    return; 
} 
int ix = blockDim.x * blockIdx.x + threadIdx.x; 
if (ix >= width) { 
    return; 
} 
int idx = iy * width + ix; 
float raysumv = d_raysump[idx]; 
if (raysumv > 0.001) { 
    d_residualp[idx] = (d_projp[idx]-d_imagep[idx])/raysumv; 
} 
else{ 
    d_residualp[idx] = 0; 
} 
} 

Answer 2

Для таких приложений производительности вам необходимо хранить двумерную матричную информацию в виде единого массива в памяти. Поэтому, если у вас есть матрица M x N, вы можете сохранить ее в одном массиве длиной M * N.

Так что, если вы хотите хранить матрицу 2х2

(1 , 2) 
(3 , 4) 

Затем создать единый массив инициализации в строке I, и столбца J, используя следующие элементы.

int rows=2; 
int cols=2; 
float* matrix = malloc(sizeof(float)*rows*cols); 
matrix[i*cols+j]=yourValue; 
//element 0,0 
matrix[0*cols+0]=1.0; 
//element 0,1 
matrix[0*cols+1]=2.0; 
//element 1,0 
matrix[1*cols+0]=3.0; 
//element 1,1 
matrix[1*cols+1]=4.0; 

Этот способ принимает массив 2D и хранить его один непрерывный кусок памяти таким способом, называется хранения данных по строкам основного заказа. См. Статью Википедии here. Как только вы измените формат своих данных в этот формат, вы можете использовать сокращение, которое было показано в SDK, и ваш код должен быть намного быстрее, так как вы сможете делать более объединенные чтения в коде ядра графического процессора.

Answer 3

В последнее время я решил этот вопрос следующим образом.

// Grid and block size 
const dim3 blockSize(16,16,1); 
const dim3 gridSize(numRows, numCols, 1); 
// kernel call 
rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols 

gridsize = Количество блока
размер_блока = Threads на блоке

Вот соответствующее ядро ​​

__global__ void rgba_to_greyscale(const uchar4* const rgbaImage, 
         unsigned char* const greyImage, 
         int numRows, int numCols) 
{ 
    int idx = blockIdx.x + blockIdx.y * numRows; 
    uchar4 pixel  = rgbaImage[idx]; 
    float intensity = 0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z; 
    greyImage[idx] = static_cast<unsigned char>(intensity); 
} 

УДАЧИ !!!