Как определить производительность графического процессора CUDA?

StackOverflow https://stackoverflow.com//questions/21062607

Вопрос

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

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

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

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

Мои вопросы таковы:

  1. Есть ли способ определить, достижима ли эта задача в реальном времени или нет?Я имею в виду с помощью максимальных сбоев и ограничений пропускной способности ввода-вывода и т.д.
  2. Как вычислить, полностью ли используется графический процессор на максимуме?
  3. Возможные способы улучшить производительность?

Технические характеристики машины:[i7-4770, 8 ГБ, GTX-680]

Объяснение текущего кода ядра:

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

Код ядра:

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

    }
}
Это было полезно?

Решение

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

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

Я понятия не имею, как определить, достижимо ли ядро в реальном времени, вы можете максимизировать свое ядро CUDA, используя Калькулятор заполняемости CUDA, Вы можете рассмотреть возможность использования текстуры, поверхностной памяти, постоянной памяти, закрепленной памяти хоста и многого другого, это зависит от реализации вашего алгоритма.

Как вычислить, полностью ли используется графический процессор на максимуме?

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

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

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

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

Если это не сработает, рассмотрите возможность использования нескольких графических процессоров.

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

Лицензировано под: CC-BY-SA с атрибуция
Не связан с StackOverflow
scroll top