Вот обработанный пример, демонстрирующие 2 из 3 концепций, обсуждаемых в комментариях:
Первая оптимизация учитывать, что 512 нитей не достаточно, чтобы держать любой GPU занят. Мы хотели бы настроить 10000 потоков или больше. Графический процессор является скрытой машиной, и когда у вас слишком мало потоков, чтобы помочь латентности скрывать GPU, ваше ядро становится привязанным к задержке, что является проблемой, связанной с памятью. Самый простой способ сделать это - каждый поток обрабатывает один пиксель на изображении (с учетом общего количества потоков 512 * 1024), а не один столбец (допускает только 512 потоков). Однако, поскольку это, похоже, «ломает» наш алгоритм «обнаружения первой поверхности», мы должны сделать еще одну модификацию (2).
Как только все пиксели обрабатываются параллельно, то простая адаптация пункта 1 выше означает, что мы больше не знаем, какая поверхность была «первой», то есть какая «яркая» поверхность (за столбец) была ближе всего к строке 0 . Эта характеристика алгоритма изменяет задачу из простой трансформации в сокращение (на самом деле сокращение на один столбец изображения). Мы разрешим обработку каждого столбца параллельно, используя 1 поток, назначенный каждому пиксель, но мы выберем результирующий пиксель, который удовлетворяет критерию яркости, который равен ближайшим к нулевой строке. Относительно простой метод для этого состоит в том, чтобы просто использовать atomicMin в массиве с одним столбцом из минимальной строки (в каждом столбце), где обнаружена подходящая яркая пиксельная окрестность.
Следующий код демонстрирует выше 2 изменения (и не включает в себя использование общей памяти) и демонстрирует (для меня, на Tesla K40) о диапазоне 1x-20x от ускорения против первоначального ядра ФПА в , Диапазон ускорений связан с тем, что алгоритмы работают в зависимости от структуры изображения. Оба алгоритма имеют стратегии выхода из предыдущего. Исходное ядро может выполнять значительно более или менее работу из-за структуры раннего выхода на контурах, в зависимости от того, где (если есть) яркие пиксельные окрестности обнаруживаются в каждом столбце.Поэтому, если все столбцы имеют яркие окрестности вблизи строки 0, я вижу «улучшение» около 1x (то есть мое ядро работает примерно с той же скоростью, что и оригинал). Если все столбцы имеют яркие окрестности (только) рядом с другим «концом» изображения, я вижу улучшение около 20x. Это может сильно различаться в зависимости от GPU, поскольку кеплерные графические процессоры улучшают глобальную пропускную способность атома, которые я использую. EDIT: Из-за этой переменной работы я добавил грубую стратегию «раннего выхода» в качестве тривиальной модификации моего кода. Это приводит к кратчайшему времени выполнения для приближения паритета между обоими ядрами (то есть около 1x).
Остающиеся оптимизации могут включать в себя:
Использование разделяемой памяти. Это должно быть тривиальной адаптацией того же подхода на основе разделяемой памяти на основе плитки, который используется, например, для умножения матрицы. Если вы используете квадратную изоляцию, тогда вам нужно будет отрегулировать размеры блока ядра/сетки, чтобы сделать их «квадратными».
Улучшенная техника сокращения. Для некоторых структур изображений атомный метод может быть несколько медленным. Это можно было бы улучшить, переключившись на правильное параллельное сокращение на столбец. Вы можете выполнить «худший» тест на моем ядре, установив тестовое изображение в качестве порогового значения везде. Это должно привести к тому, что исходное ядро будет работать быстрее, а мое ядро будет работать медленнее, но в этом случае я не наблюдал значительного замедления моего ядра. Время выполнения моего ядра довольно постоянное. Опять же, это может быть связано с GPU.
Пример:
#include <stdlib.h>
#include <stdio.h>
#define SKIP_ROWS 10
#define THRESH 1.0f
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL
unsigned long long dtime_usec(unsigned long long start){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}
__global__ void firstSurfaceDetection (float *threshImg, float *firstSurfaceImg, int height, int width) {
int col = threadIdx.x + (blockDim.x*blockIdx.x);
int rows2skip = SKIP_ROWS;
float thresh = THRESH;
//thread Index: (0 -> 511)
if (col < width) {
if(col == 0) { // first col - 0
for (int row = 0 + rows2skip; row < height; row++) { // skip first 30 rows
int cnt = 0;
float neibs[6]; // not shared mem as it reduces speed
// get six neighbours - three in same col, and three to the right
neibs[0] = threshImg[((row)*width) +(col)]; if(neibs[0] == thresh) { cnt++; } // current position
neibs[1] = threshImg[((row)*width) +(col+1)]; if(neibs[1] == thresh) { cnt++; } // right
neibs[2] = threshImg[((row+1)*width) +(col)]; if(neibs[2] == thresh) { cnt++; } // bottom
neibs[3] = threshImg[((row+1)*width) +(col+1)]; if(neibs[3] == thresh) { cnt++; } // bottom right
neibs[4] = threshImg[((row+2)*width) +(col)]; if(neibs[4] == thresh) { cnt++; } // curr offset by 2 - bottom
neibs[5] = threshImg[((row+2)*width) +(col+1)]; if(neibs[5] == thresh) { cnt++; } // curr offset by 2 - bottom right
if(cnt == 6) { // if all neighbours are bright, we are at the edge boundary
firstSurfaceImg[(row)*width + col] = 1.0f;
row = height;
}
}
}
else if (col == (width-1)) { // last col
for (int row = 0 + rows2skip; row < height; row++) {
int cnt = 0;
float neibs[6]; // not shared mem as it reduces speed
// get six neighbours - three in same col, and three to the left
neibs[0] = threshImg[((row)*width) +(col)]; if(neibs[0] == thresh) { cnt++; } // current position
neibs[1] = threshImg[((row)*width) +(col-1)]; if(neibs[1] == thresh) { cnt++; } // left
neibs[2] = threshImg[((row+1)*width) +(col)]; if(neibs[2] == thresh) { cnt++; } // bottom
neibs[3] = threshImg[((row+1)*width) +(col-1)]; if(neibs[3] == thresh) { cnt++; } // bottom left
neibs[4] = threshImg[((row+2)*width) +(col)]; if(neibs[4] == thresh) { cnt++; } // curr offset by 2 - bottom
neibs[5] = threshImg[((row+2)*width) +(col-1)]; if(neibs[5] == thresh) { cnt++; } // curr offset by 2 - bottom left
if(cnt == 6) { // if all neighbours are bright, we are at the edge boundary
firstSurfaceImg[(row)*width + col] = 1.0f;
row = height;
}
}
}
// remaining threads are: (1 -> 510)
else { // any col other than first or last column
for (int row = 0 + rows2skip; row < height; row++) {
int cnt = 0;
float neibs[9]; // not shared mem as it reduces speed
// for threads < width/4, get the neighbors
// get nine neighbours - three in curr col, three each to left and right
neibs[0] = threshImg[((row)*width) +(col-1)]; if(neibs[0] == thresh) { cnt++; }
neibs[1] = threshImg[((row)*width) +(col)]; if(neibs[1] == thresh) { cnt++; }
neibs[2] = threshImg[((row)*width) +(col+1)]; if(neibs[2] == thresh) { cnt++; }
neibs[3] = threshImg[((row+1)*width) +(col-1)]; if(neibs[3] == thresh) { cnt++; }
neibs[4] = threshImg[((row+1)*width) +(col)]; if(neibs[4] == thresh) { cnt++; }
neibs[5] = threshImg[((row+1)*width) +(col+1)]; if(neibs[5] == thresh) { cnt++; }
neibs[6] = threshImg[((row+2)*width) +(col-1)]; if(neibs[6] == thresh) { cnt++; }
neibs[7] = threshImg[((row+2)*width) +(col)]; if(neibs[7] == thresh) { cnt++; }
neibs[8] = threshImg[((row+2)*width) +(col+1)]; if(neibs[8] == thresh) { cnt++; }
if(cnt == 9) { // if all neighbours are bright, we are at the edge boundary
firstSurfaceImg[(row)*width + col] = 1.0f;
row = height;
}
}
}
}
__syncthreads();
}
__global__ void firstSurfaceDetection_opt (const float * __restrict__ threshImg, int *firstSurfaceImgRow, int height, int width) {
int col = threadIdx.x + (blockDim.x*blockIdx.x);
int row = threadIdx.y + (blockDim.y*blockIdx.y);
int rows2skip = SKIP_ROWS;
float thresh = THRESH;
if ((row >= rows2skip) && (row < height-2) && (col < width) && (row < firstSurfaceImgRow[col])) {
int cnt = 0;
int inc = 0;
if (col == 0) inc = +1;
if (col == (width-1)) inc = -1;
if (inc){
cnt = 3;
if (threshImg[((row)*width) +(col)] == thresh) cnt++;
if (threshImg[((row)*width) +(col+inc)] == thresh) cnt++;
if (threshImg[((row+1)*width) +(col)] == thresh) cnt++;
if (threshImg[((row+1)*width) +(col+inc)] == thresh) cnt++;
if (threshImg[((row+2)*width) +(col)] == thresh) cnt++;
if (threshImg[((row+2)*width) +(col+inc)] == thresh) cnt++;
}
else {
// get nine neighbours - three in curr col, three each to left and right
if (threshImg[((row)*width) +(col-1)] == thresh) cnt++;
if (threshImg[((row)*width) +(col)] == thresh) cnt++;
if (threshImg[((row)*width) +(col+1)] == thresh) cnt++;
if (threshImg[((row+1)*width) +(col-1)] == thresh) cnt++;
if (threshImg[((row+1)*width) +(col)] == thresh) cnt++;
if (threshImg[((row+1)*width) +(col+1)] == thresh) cnt++;
if (threshImg[((row+2)*width) +(col-1)] == thresh) cnt++;
if (threshImg[((row+2)*width) +(col)] == thresh) cnt++;
if (threshImg[((row+2)*width) +(col+1)] == thresh) cnt++;
}
if(cnt == 9) { // if all neighbours are bright, we are at the edge boundary
atomicMin(firstSurfaceImgRow + col, row);
}
}
}
int main(int argc, char *argv[]){
float *threshImg, *h_threshImg, *firstSurfaceImg, *h_firstSurfaceImg;
int *firstSurfaceImgRow, *h_firstSurfaceImgRow;
int actualImHeight = 1024;
int actualImWidth = 512;
int row_set = 512;
if (argc > 1){
int my_val = atoi(argv[1]);
if ((my_val > SKIP_ROWS) && (my_val < actualImHeight - 3)) row_set = my_val;
}
h_firstSurfaceImg = (float *)malloc(actualImHeight*actualImWidth*sizeof(float));
h_threshImg = (float *)malloc(actualImHeight*actualImWidth*sizeof(float));
h_firstSurfaceImgRow = (int *)malloc(actualImWidth*sizeof(int));
cudaMalloc(&threshImg, actualImHeight*actualImWidth*sizeof(float));
cudaMalloc(&firstSurfaceImg, actualImHeight*actualImWidth*sizeof(float));
cudaMalloc(&firstSurfaceImgRow, actualImWidth*sizeof(int));
cudaMemset(firstSurfaceImgRow, 1, actualImWidth*sizeof(int));
cudaMemset(firstSurfaceImg, 0, actualImHeight*actualImWidth*sizeof(float));
for (int i = 0; i < actualImHeight*actualImWidth; i++) h_threshImg[i] = 0.0f;
// insert "bright row" here
for (int i = (row_set*actualImWidth); i < ((row_set+3)*actualImWidth); i++) h_threshImg[i] = THRESH;
cudaMemcpy(threshImg, h_threshImg, actualImHeight*actualImWidth*sizeof(float), cudaMemcpyHostToDevice);
dim3 grid(1,1024);
//warm-up run
firstSurfaceDetection_opt <<< grid, 512 >>> (threshImg, firstSurfaceImgRow, actualImHeight, actualImWidth);
cudaDeviceSynchronize();
cudaMemset(firstSurfaceImgRow, 1, actualImWidth*sizeof(int));
cudaDeviceSynchronize();
unsigned long long t2 = dtime_usec(0);
firstSurfaceDetection_opt <<< grid, 512 >>> (threshImg, firstSurfaceImgRow, actualImHeight, actualImWidth);
cudaDeviceSynchronize();
t2 = dtime_usec(t2);
cudaMemcpy(h_firstSurfaceImgRow, firstSurfaceImgRow, actualImWidth*sizeof(float), cudaMemcpyDeviceToHost);
unsigned long long t1 = dtime_usec(0);
firstSurfaceDetection <<< 1, 512 >>> (threshImg, firstSurfaceImg, actualImHeight, actualImWidth);
cudaDeviceSynchronize();
t1 = dtime_usec(t1);
cudaMemcpy(h_firstSurfaceImg, firstSurfaceImg, actualImWidth*actualImHeight*sizeof(float), cudaMemcpyDeviceToHost);
printf("t1 = %fs, t2 = %fs\n", t1/(float)USECPSEC, t2/(float)USECPSEC);
// validate results
for (int i = 0; i < actualImWidth; i++)
if (h_firstSurfaceImgRow[i] < actualImHeight)
if (h_firstSurfaceImg[(h_firstSurfaceImgRow[i]*actualImWidth)+i] != THRESH)
{printf("mismatch at %d, was %f, should be %d\n", i, h_firstSurfaceImg[(h_firstSurfaceImgRow[i]*actualImWidth)+i], THRESH); return 1;}
return 0;
}
$ nvcc -arch=sm_35 -o t667 t667.cu
$ ./t667
t1 = 0.000978s, t2 = 0.000050s
$
Примечание:
приведенные выше пример вставляет "яркие окрестности" весь путь через изображение в строке = 512, таким образом давая среднее в моем случае коэффициент ускорения почти 20x (K40c).
для краткости представления, я отказался от proper cuda error checking. Однако я рекомендую его.
Производительность обоих ядер сильно зависит от того, выполняется ли это первый запуск или нет. Это, вероятно, связано с кешированием и эффектами разминки. Поэтому, чтобы дать разумные результаты, я запускал свое ядро сначала как дополнительный бессрочный разминочный запуск.
Одна из причин, по которой я не преследовал оптимизацию с общей памятью, заключается в том, что эта проблема уже довольно мала и, по крайней мере, для большого кеплерного графического процессора, такого как K40, она почти полностью подходит для кеша L2 (особенно для моего ядро, поскольку он использует меньшую структуру выходных данных.) Учитывая, что разделяемая память может не дать большой импульс.
EDIT: Я изменил код (снова) так, чтобы линия (строка) в тестовом изображении, где яркая граница вставлена может быть передана в качестве параметра командной строки, и я тестировал код на 3-х различных устройств, используя 3 различные настройки для яркой строки:
execution time on: K40 C2075 Quadro NVS 310
bright row = 15: 31/33 23/45 29/314
bright row = 512: 978/50 929/112 1232/805
bright row = 1000: 2031/71 2033/149 2418/1285
all times are microseconds (original kernel/optimized kernel)
CUDA 6.5, CentOS 6.2
Я не могу помочь с вопросом, но убедитесь, что вы используете [последние драйверы CUDA] (https: // developer.nvidia.com/cuda-zone) – taco
512 потоков недостаточно, чтобы GPU был занят.И если вы заинтересованы в производительности, вы никогда не захотите запускать ядра, которые бы вот так: '<<<1,...> >>' или это: '<<<...,1> >>' Вы выставили параллелизм по ширине изображения, теперь пришло время выставить параллелизм по высоте изображения. Избавьтесь от ваших for-loops и увеличьте сетку до достаточного количества потоков (возможно, захотите перейти в 2D-сетку), чтобы каждый поток обрабатывал один пиксель вместо того, чтобы каждый поток обрабатывал один столбец. –
После того, как вы подсчитали количество потоков, вы можете использовать разделяемую память очень просто, загрузив блок данных изображения в общую память и выработайте каждый поток из области разделяемой памяти для нагрузок и тестов. Я не знаю, почему люди помещают '__syncthreads()' в конец ядра. Здесь нет никакой цели. –