CUDA - нет блоков, только темы для неопределенных размеров

Question

У меня есть несколько матриц с неизвестными размерами, варьирующимися от 10 до 20 000 в обоих направлениях.

Я разработал ядро ​​CUDA с (x; y) блоками и (x; y) потоками.

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

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

Мой вопрос: что если я полностью устраню блоки и просто создаю сетку из потоков x; y? Будут ли у блока SM проблемы без большого количества блоков?

Могу ли я удалить блоки и использовать большое количество потоков или необходимо разделить блоки?

Answer 1

Вы не можете просто создать «сетку потоков», так как вам нужно организовать потоки по блокам, и вы можете иметь максимум 512 потоков на блок.Тем не менее, вы можете эффективно сделать это, используя 1 поток на блок, что приведет к сетке X на Y блоков 1x1.Однако это приведет к довольно ужасной производительности из-за нескольких факторов:

1. Согласно Руководству по программированию CUDA, SM может обрабатывать максимум 8 блоков в любое время.Это ограничит вас до 8 потоков на SM, что недостаточно для заполнения даже одной основы.Если у вас, скажем, 48 ядер CUDA, вы сможете обрабатывать только 384 потока в любой момент времени.

2. Если на SM доступно только 8 потоков, их будет слишком мало.деформации, чтобы скрыть задержки памяти.Графический процессор будет проводить большую часть своего времени, ожидая завершения доступа к памяти, а не выполнения каких-либо вычислений.

3. Вы не сможете объединить операции чтения и записи в память, что приведет к плохому использованию пропускной способности памяти.

4. Вы не сможете эффективно использовать разделяемую память, поскольку это общий ресурс между потоками в блоке.

При необходимостиУбедитесь, что корректность потоков в блоке раздражает, ваша производительность будет значительно лучше, чем ваша идея "сетки потоков".

Answer 2

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

PS: у меня всегда есть block_size == 128, потому что это был хороший компромисс между занятостью многоядерности, использованием регистров, требованиями к общей памяти и коалесцентным доступом для всех моих ядер.

Код для расчета хорошего размера сетки (хост):

#define GRID_SIZE 65535

//calculate grid size (store result in grid/block)
void kernelUtilCalcGridSize(unsigned int num_threads, unsigned int block_size, dim3* grid, dim3* block) {


    //block
    block->x = block_size;
    block->y = 1;
    block->z = 1;


    //number of blocks
    unsigned int num_blocks = kernelUtilCeilDiv(num_threads, block_size);
    unsigned int total_threads = num_blocks * block_size;
    assert(total_threads >= num_threads);

    //calculate grid size
    unsigned int gy = kernelUtilCeilDiv(num_blocks, GRID_SIZE);
    unsigned int gx = kernelUtilCeilDiv(num_blocks, gy);
    unsigned int total_blocks = gx * gy;
    assert(total_blocks >= num_blocks);

    //grid
    grid->x = gx;
    grid->y = gy;
    grid->z = 1;
}

//ceil division (rounding up)
unsigned int kernelUtilCeilDiv(unsigned int numerator, unsigned int denominator) {
    return (numerator + denominator - 1) / denominator;
}

Код для расчета уникального идентификатора потока и проверки границ (устройства):

//some kernel
__global__ void kernelFoo(unsigned int num_threads, ...) {


    //calculate unique id
    const unsigned int thread_id = threadIdx.x;
    const unsigned int block_id = blockIdx.x + blockIdx.y * gridDim.x;
    const unsigned int unique_id = thread_id + block_id * blockDim.x;


    //check range
    if (unique_id >= num_threads) return;

    //do the actual work
    ...
}

Я не думаю, что нужно много усилий / регистров / строк кода, чтобы проверить правильность.