Почему простая функция CUDA требует так много локальной памяти?

Question

Я написал простую функцию на CUDA. Это изменить размер изображения в два раза. Для изображения с разрешением 1920*1080 для выполнения этой функции требуется ~20 мс. Я пробовал другой способ оптимизировать эту функцию. И я обнаружил, что, возможно, именно локальная память является ключевой причиной.

Я пробовал три разных метода для получения изображения.

• Модуль Gpu в OpenCV
• Привязка текстур к GpuMat в OpenCV
• Прямой выбор GpuMat из глобальной памяти

Ни один из них не может принести мне немного улучшения.

Тогда я использовал nvvp, чтобы выяснить причину. И накладные расходы локальной памяти составляют ~95% во всех трех вышеупомянутых условиях.

Поэтому я перехожу к своему коду, чтобы узнать, как nvcc использует память. Затем я обнаружил, что простая функция, как это:

__global__ void performDoubleImage(float* outData, size_t step, const int cols, const int rows)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    if (x >= cols)
        return;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if (y >= rows)
        return;
    ((float*)((size_t)outData+y*step))[x] = tex2D(texRef, x/2, y/2);
}

нужен кадр стека 80 байт (они находятся в локальной памяти).

И еще одна функция, как это:

__global__ void performFinalDoubleImage(const PtrStepSz<float> in, PtrStepSz<float> out)
{
    out(out.rows-1, out.cols-1) = out(in.rows-1, in.cols-1);
}

также необходим кадр стека 88 байтов.

Вопрос в том, почему моя функция использует так много локальной памяти и записывается в этой простой задаче? И почему функция в OpenCV может выполнять ту же функцию, не используя локальную память (это тест nvvp, загрузка локальной памяти равна нулю)?

Мой код скомпилирован в режиме отладки. А у меня карта GT650(192 SP/SM, 2 SM).

Answer 1

Две опубликованные вами функции слишком просты, чтобы использовать такой большой стек, фактически они вообще не должны использовать стек. Наиболее вероятная причина такого большого разлива заключается в том, что вы компилируете с отключенными оптимизациями (например, в режиме отладки).

Для справки, Роберт Кровелла скомпилировал ваше первое ядро ​​в релизе и в режиме отладки:

Debug:

Информация ptxas: Свойства функции для _Z18performDoubleImagePfmii Кадр стека 256 байтов, 0 хранилищ байтов, 0 разливов байтов Информация ptxas: Использовано 23 регистра, совокупный размер стека 296 байтов, 56 байтов cmem[0], 1 текстура

Релиз:

Информация ptxas: Свойства функции для _Z18performDoubleImagePfmii 0-байтовый фрейм стека, 0 байтовых хранилищ разлива, 0 байтовых разливов загружает ptxas информация: Используются 9 регистров, 56 байт cmem[0], 1 текстура

Обратите внимание на разницу в использовании стека и регистра. Как отмечено в комментариях, при измерении производительности программы вы всегда должны компилировать для максимального уровня оптимизации, иначе измерения будут бессмысленными.