Windows 7, NVidia GeForce 425M.
Я написал простой код CUDA, который вычисляет суммы строк матрицы. Матрица имеет одномерное представление (указатель на float).
Последовательная версия кода ниже (она имеет 2
петли, как и ожидалось):
void serial_rowSum (float* m, float* output, int nrow, int ncol) {
float sum;
for (int i = 0 ; i < nrow ; i++) {
sum = 0;
for (int j = 0 ; j < ncol ; j++)
sum += m[i*ncol+j];
output[i] = sum;
}
}
Внутри кода CUDA я вызываю функцию ядра, подметая матрицу по строкам. Ниже фрагмент вызова ядра:
dim3 threadsPerBlock((unsigned int) nThreadsPerBlock); // has to be multiple of 32
dim3 blocksPerGrid((unsigned int) ceil(nrow/(float) nThreadsPerBlock));
kernel_rowSum<<<blocksPerGrid, threadsPerBlock>>>(d_m, d_output, nrow, ncol);
и функция ядра, которая выполняет параллельную сумму строк (все еще имеет цикл 1
):
__global__ void kernel_rowSum(float *m, float *s, int nrow, int ncol) {
int rowIdx = threadIdx.x + blockIdx.x * blockDim.x;
if (rowIdx < nrow) {
float sum=0;
for (int k = 0 ; k < ncol ; k++)
sum+=m[rowIdx*ncol+k];
s[rowIdx] = sum;
}
}
Пока все хорошо. Серийный и параллельный (CUDA) результаты равны.
Все дело в том, что версия CUDA занимает почти в два раза больше времени для последовательного вычисления, даже если я изменяю параметр nThreadsPerBlock
: я тестировал nThreadsPerBlock
от 32
до 1024
(максимальное количество потоки на блок разрешены для моей карты).
IMO, размер матрицы достаточно велик, чтобы оправдать распараллеливание: 90,000 x 1,000
.
Ниже я сообщаю время, прошедшее для серийной и параллельной версий, используя разные nThreadsPerBlock
. Время, указанное в msec
, в среднем по сравнению с 100
образцами:
Матрица: nrow = 90000 x ncol = 1000
Серийный: среднее время, прошедшее за выборку в msec (100
samples): 289.18
.
CUDA (32
ThreadsPerBlock): среднее время, прошедшее за выборку в msec (100
samples): 497.11
.
CUDA (1024
ThreadsPerBlock): среднее время, прошедшее за выборку в msec (100
samples): 699.66
.
На всякий случай, версия с 32
/1024
nThreadsPerBlock
является самой быстрой/медленной.
Я понимаю, что при копировании с Host на Device есть какой-то накладные расходы, но, может быть, медленность заключается в том, что я не реализую самый быстрый код.
Поскольку я далек от того, чтобы быть экспертом CUDA:
Я кодирую самую быструю версию для этой задачи? Как я могу улучшить свой код? Могу ли я избавиться от цикла в функции ядра?
Любые мысли оценили.
РЕДАКТИРОВАТЬ 1
Хотя я описываю стандарт rowSum
, меня интересует работа строк AND
/OR
, имеющих значения (0;1}
, такие как rowAND
/rowOR
. Тем не менее, это не позволяет мне использовать cuBLAS
умножить на трюк столбца тэга 1
COL
, как это было предложено некоторыми комментаторами.
РЕДАКТИРОВАТЬ 2
Как предлагают пользователи других пользователей и здесь одобрены:
ЗАБУДЬТЕ О ПЫТАХ НАПИСАТЬ СВОИ СОБСТВЕННЫЕ ФУНКЦИИ, вместо этого используйте библиотеку Thrust и придет волшебство.