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

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

Согласно моему выводу deviceQuery, мой графический процессор имеет 16MP, 32 ядра / mp, максимальное количество блоков составляет 1024x1024x64, а у меня максимальное количество потоков / блок = 1024.

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

В существующем коде изображения хранятся в виде 2D-массива [строки] [столбцы]. Итак, это ядро ​​в C выглядит так, как вы ожидаете, с циклом по строкам и циклом по столбцам с вычислением посередине.

Итак, как мне настроить часть вычисления размеров этого кода в CUDA? Я посмотрел код сокращения в SDK, но это для одномерного массива. В нем нет никакого упоминания о том, как настроить количество блоков и потоков, когда у вас есть что-то 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];
}
}
}

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

Любые советы приветствуются! Спасибо!


c++ gpu cuda gpgpu
person Derek    schedule 25.01.2011    source источник

Ответы (3)


arrow_upward
3
arrow_downward

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

// 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 = Количество блоков
Blocksize = Количество потоков на блок

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

__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);   
}

Удачи!!!

person Community    schedule 29.04.2015

arrow_upward
1
arrow_downward

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

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

(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-массив и сохранить его в виде единой непрерывной части памяти, называется хранением данных в строчном порядке. См. Статью Википедии здесь. После того, как вы измените макет ваших данных на этот формат, вы можете использовать сокращение, которое было показано в SDK, и ваш код должен быть намного быстрее, так как вы сможете выполнять больше объединенных операций чтения в коде ядра графического процессора.

person Samsdram    schedule 26.01.2011
comment
Я согласен, что это самый простой (и, вероятно, наиболее эффективный) способ решения этой проблемы. Меня беспокоит только точность: если вы сокращаете сумму очень больших изображений с помощью высокоточных пикселей, то у вас могут закончиться биты, поэтому убедитесь, что вы используете достаточно большой тип данных. В качестве альтернативы вы можете изменить сокращение для вычисления скользящего среднего, а не суммы. - person harrism; 04.09.2012

arrow_upward
0
arrow_downward

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

Определяет и справочные функции:

#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;
}
}
person peakxu    schedule 26.01.2011
comment
Если я понимаю, что делает iDivUP, вы можете немного упростить логику благодаря целочисленному усечению: return (a + b-1) / b; - person Erich Mirabal; 15.02.2013