Почему существует CL_DEVICE_MAX_WORK_GROUP_SIZE?

Question

Я пытаюсь понять архитектуру устройств OpenCL, таких как графические процессоры, и не могу понять, почему существует явное ограничение на количество рабочих элементов в локальной рабочей группе, то есть константу CL_DEVICE_MAX_WORK_GROUP_SIZE.

Мне кажется, что об этом должен позаботиться компилятор, т. Е. Если (одномерное для простоты) ядро ​​выполняется с размером локальной рабочей группы 500, а его физический максимум равен 100, и ядро ​​выглядит, например, так:

__kernel void test(float* input) {
    i = get_global_id(0);
    someCode(i);
    barrier();
    moreCode(i);
    barrier();
    finalCode(i);
}

тогда его можно автоматически преобразовать в выполнение с размером рабочей группы 100 в этом ядре:

__kernel void test(float* input) {
    i = get_global_id(0);
    someCode(5*i);
    someCode(5*i+1);
    someCode(5*i+2);
    someCode(5*i+3);
    someCode(5*i+4);
    barrier();
    moreCode(5*i);
    moreCode(5*i+1);
    moreCode(5*i+2);
    moreCode(5*i+3);
    moreCode(5*i+4);
    barrier();
    finalCode(5*i);
    finalCode(5*i+1);
    finalCode(5*i+2);
    finalCode(5*i+3);
    finalCode(5*i+4);
}

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

Answer 1

Я думаю, что источник CL_DEVICE_MAX_WORK_GROUP_SIZE лежит в базовой аппаратной реализации.

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

Что касается компилятора, это можно сделать, чтобы сделать это возможным. Это может работать, но понимать, что это будет означать перекомпиляцию ядра заново. Что не всегда возможно. Я могу представить себе ситуации, когда разработчики выгружают скомпилированное ядро ​​для каждой платформы в двоичном формате и поставляют его со своим программным обеспечением только по причинам «не с открытым исходным кодом».

Answer 2

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