Требования к памяти CUDA

Question

Я недавно написал очень простое ядро:

__device__ uchar elem(const Matrix m, int row, int col) {
    if(row == -1) {
        row = 0;
    } else if(row > m.rows-1) {
        row = m.rows-1;
    }

    if(col == -1) {
        col = 0;
    } else if(col > m.cols-1) {
        col = m.cols-1;
    }
    return *((uchar*)(m.data + row*m.step + col));
}

/**
* Each thread will calculate the value of one pixel of the image 'res'
*/
__global__ void resizeKernel(const Matrix img, Matrix res) {
    int row = threadIdx.y + blockIdx.y * blockDim.y;
    int col = threadIdx.x + blockIdx.x * blockDim.x;

    if(row < res.rows && col < res.cols) {
        uchar* e = res.data + row * res.step + col;

        *e = (elem(img, 2*row, 2*col) >> 2) +
             ((elem(img, 2*row, 2*col-1) + elem(img, 2*row, 2*col+1) 
             + elem(img, 2*row-1, 2*col) + elem(img, 2*row+1, 2*col)) >> 3) +
             ((elem(img, 2*row-1, 2*col-1) + elem(img, 2*row+1, 2*col+1)
             + elem(img, 2*row+1, 2*col-1) + elem(img, 2*row-1, 2*col+1)) >> 4);
    }
}

По сути, он вычисляет значение пикселя изображения уменьшенного размера, используя значения большего изображения. Внутри «если» в resizeKernel.

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

Моя теория заключалась в том, что это может иметь какое-то отношение к доступной памяти для хранения промежуточных результатов выражения. Итак, уменьшив количество потоков на блок, он начал работать отлично, без необходимости сокращать количество операций.

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

Спасибо!

Answer 1

Это, скорее всего, регистры, и вы можете узнать потребление регистров для каждого потока, добавив опцию -Xptxas="-v" к вызову nvcc, который компилирует ядро.Ассемблер возвращает количество регистров на поток, статическую разделяемую память, локальную память и постоянную память, используемую скомпилированным кодом.

NVIDIA создает электронную таблицу калькулятора занятости ( доступна здесь ), в которую можно подключить выход ассемблера, чтобы увидеть допустимый диапазон размеров блоков и их влияние на занятость графического процессора.Глава 3 руководства по программированию CUDA также содержит подробное обсуждение концепции занятости и того, как взаимодействуют размер блока и требования к ресурсам ядра.