Как определить производительность CUDA gpu?

Question

Я пишу программу cuda для соответствия каждому входному изображению разрешения ~ 180X180, с примерно 10000 изображений шаблонов с разрешением ~ 128 * 128. Целью является достижение производительности в реальном времени, т. Е. Сопоставление шаблонов из 25 ~ 30 входных изображений (каждый со всеми 10000 шаблонами) за 1 секунду.

В настоящее время я использую следующий подход

1. Поджатая все шаблоны на GPU глобальной памяти, чтобы сохранить время выполнения операций ввода/вывода.
2. Создал единое ядро ​​для соответствия одному источнику изображения со всеми изображениями шаблонов и возвращает массив для всех положительных совпадений.
3. Выполнение всех операций во временной области (без использования БПФ). причина была, я пробовал реализацию Radix-4 fft, но для этого требуется много промежуточных глобальных чтений и записей, в результате чего требуется больше времени.

до сих пор для 1 входного изображения до 10000 шаблонов, он занимает около 2 секунд.

Мои вопросы:

1. ли есть способ определить, является ли эта задача achieveable в реальном времени или нет? Я имею в виду с помощью maxmimum FLOPS и ограничений пропускной способности ввода-вывода e.t.c.
2. Как вычислить, будет ли GPU полностью утилизирован максимум?
3. Возможные способы повышения производительности?

машина спецификации: [i7-4770, 8GB, GTX-680]

Explaination текущего кода ядра:

1. все изображения шаблона [размер примерно 128х128 в RGB] являются в - загружается в память GPU. Идея заключается в том, чтобы сохранить ввод-вывод во время работы.
2. Каждое входное изображение загружается в память текстуры, поэтому текстура является хорошим вариантом для 2D-адресации.
3. Каждый «Блок» имеет 1024 потолка.
4. Каждый поток вычисляет значение для каждого выходного пикселя, размер вывода - [31X31 = 961 пиксель].
5. Число запущенных блоков равно количеству совпадений шаблонов.

Kernel Код:

__global__ void cudaMatchTemplate(TemplateArray *templates, uchar *Match) 
{ 
    int global = blockIdx.x*blockDim.x + threadIdx.x; 

    __shared__ int idx[TEMPLATE_MATCH_DIM]; 
    __shared__ float out_shared[TEMPLATE_MATCH_DIM]; 

    //halving the template size.... 
    int rows = (templates[blockIdx.x].nHeight)/2; 
    int cols = (templates[blockIdx.x].nWidth)/2; 

    int fullCol = templates[blockIdx.x].nWidth; 

    int x = templates[blockIdx.x].nMatchLeft; 
    int y = templates[blockIdx.x].nMatchTop; 

    int offset_y = (threadIdx.x/TEMPLATE_MATCH_SIZE); 
    int offset_x = (threadIdx.x - offset_y*TEMPLATE_MATCH_SIZE); 

    // *************** Performing match in time domain *****************************// 
    int sum = 0; 
    float temp; 
    int idxXFactor = 3*(2*(offset_x) + x); 
    int idxYFactor = 2*(offset_y) + y ; 

    for (int i = 0; i < rows; i++) 
    { 
     int I=3*i*fullCol; 
     int sourceIdxY = idxYFactor + 2*i; 
     for (int j = 0; j < cols; j++) 
     { 
      int J=3*j; 
      int sourceIdxX = idxXFactor + 2*J;   
      int templateIdx = 2*I+2*J; 
      //**** R *****// 
      temp = float(tex2D(SourceImgColorTex,sourceIdxX,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx]); 
      sum = sum + temp*temp; 
      //**** G *****// 
      temp = float(tex2D(SourceImgColorTex,sourceIdxX+1,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +1]); 
      sum = sum + temp*temp; 
      //**** B *****// 
      temp = float(tex2D(SourceImgColorTex,sourceIdxX+2,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +2]); 
      sum = sum + temp*temp; 
     } 
    } 

    __syncthreads(); 

//placing all values in shared memory for comparison. 
    if(threadIdx.x < TEMPLATE_MATCH_DIM) 
    { 
     idx[threadIdx.x] = threadIdx.x; 
     out_shared[threadIdx.x] = sum; 
    } 
    __syncthreads(); 


// //computing the Min location.....// 

#pragma unroll 
    for(int s=512; s>0; s>>=1) 
    { 
     if ((threadIdx.x < s) &&((threadIdx.x + s)<TEMPLATE_MATCH_DIM)) 
     { 
      idx[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? idx[threadIdx.x] : idx[threadIdx.x + s]; 
      out_shared[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? out_shared[threadIdx.x] : out_shared[threadIdx.x + s];   
     } 

    } 

    __syncthreads(); 

    if(threadIdx.x <1) 
    { 
     int half_Margin = MARGIN_FOR_TEMPLATE_MATCH/2; 
     int matchY = idx[0]/TEMPLATE_MATCH_SIZE ; 
     int matchX = idx[0] - matchY * TEMPLATE_MATCH_SIZE; 

     int diff = absolute(half_Margin - matchX) + absolute(half_Margin - matchY); 
     if(diff < THRESHOLD) 
     { 
      Match[blockIdx.x] = 1; 
     } 
     else 
      Match[blockIdx.x] = 0; 

    } 
} 

Answer 1

Я постараюсь ответить на большинство ваших вопросов

ли есть способ определить, является ли эта задача achieveable в реальном времени или нет? Я имею в виду с помощью maxmimum FLOPS и ограничений пропускной способности ввода-вывода e.t.c.

Я понятия не имею, как определить, является ли ядро ​​реального времени достижимо, вы можете максимизировать CUDA ядро ​​с помощью CUDA Occupancy Calculator, Вы можете рассмотреть возможность использования текстуры, поверхность памяти, постоянную памяти, прикрепленный хост память и больше , Это зависит от реализации вашего алгоритма.

Как вычислить, будет ли GPU полностью утилизирован максимум?

Вы можете использовать калькулятор занятости CUDA и визуальный профайлер CUDA. Я бы очень рекомендовал использовать визуальный профайлер, который поможет вам понять CUDA.

Возможные способы повышения производительности?

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

Если это не сработает, попробуйте работать с задержкой, одновременно используйте несколько потоков, обращаясь к графическому процессору, поскольку CC 3.5 CUDA запустил HyperQ, это может помочь вам выполнить несколько вызовов параллельно.

Если это не сработает, рассмотрите возможность использования нескольких устройств GPU.

Пожалуйста держите нас в курсе