¿Por qué hay CL_DEVICE_MAX_WORK_GROUP_SIZE?

Question

Estoy intentando comprender la arquitectura de los dispositivos OpenCL, como las GPU, y no veo por qué hay un límite explícito en la cantidad de elementos de trabajo en un grupo de trabajo local, es decir, la constante CL_DEVICE_MAX_WORK_GROUP_SIZE.

Me parece que esto debe ser resuelto por el compilador, es decir, si se ejecuta un kernel (unidimensional por simplicidad) con un tamaño de grupo de trabajo local 500 mientras que su máximo físico es 100, y el kernel busca un ejemplo de esta manera:

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

entonces podría ser convertido automáticamente a una ejecución con el tamaño del grupo de trabajo 100 en este núcleo:

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

sin embargo, parece que esto no se hace por defecto. Por qué no? ¿Hay alguna manera de automatizar este proceso (que no sea escribir un precompilador para mí)? ¿O hay un problema intrínseco que puede hacer que mi método falle en ciertos ejemplos (y puede darme uno)?

Answer 1

Creo que el origen de CL_DEVICE_MAX_WORK_GROUP_SIZE se encuentra en la implementación del hardware subyacente.

Se están ejecutando varios subprocesos simultáneamente en unidades de cómputo y cada uno de ellos necesita mantener el estado (para llamada, jmp, etc.). La mayoría de las implementaciones usan una pila para esto y si nos fijamos en la familia Evergreen de AMD, se trata de un límite de hardware para la cantidad de entradas de la pila que están disponibles (cada entrada de la pila tiene subinserciones). Lo que en esencia limita el número de hilos que cada unidad de computación puede manejar simultáneamente.

En cuanto al compilador, puede hacer esto para que sea posible. Podría funcionar, pero entender que significaría volver a compilar el kernel. Lo cual no siempre es posible. Puedo imaginar situaciones en las que los desarrolladores vuelcan el kernel compilado para cada plataforma en un formato binario y lo envían con su software solo por razones "no de código abierto".

Answer 2

Estas constantes son consultadas desde el dispositivo por el compilador para determinar un tamaño de grupo de trabajo adecuado en tiempo de compilación (donde la compilación por supuesto se refiere a la compilación del kernel). Puede que te esté equivocando, pero parece que estás pensando en establecer esos valores por ti mismo, lo que no sería el caso.

La responsabilidad está dentro de su código para consultar las capacidades del sistema para estar preparado para cualquier hardware en el que se ejecutará.