cuda - ¿Cuál es la forma más eficiente de calcular el id warp id/lane en una cuadrícula 1-D?

ptx

Question

En CUDA, cada hilo conoce su índice de bloque en la grilla y el índice de hilo dentro del bloque. Pero dos valores importantes no parecen estar explícitamente disponibles para él:

• Su índice como un carril dentro de su warp (su "identificación de carril")
• El índice de la urdimbre del cual es un carril dentro del bloque (su "id warp")

Asumiendo que la grilla es unidimensional (aka lineal, es decir, blockDim.y y blockDim.z son 1), obviamente uno puede obtenerlos de la siguiente manera:

enum : unsigned { warp_size = 32 };
auto lane_id = threadIdx.x % warp_size;
auto warp_id = threadIdx.x / warp_size;

y si no confías en el compilador para optimizar eso, puedes reescribirlo como:

enum : unsigned { warp_size = 32, log_warp_size = 5 };
auto lane_id = threadIdx.x & (warp_size - 1);
auto warp_id = threadIdx.x >> log_warp_size;

¿Es eso lo más eficiente que puedes hacer? Todavía parece una gran pérdida para cada hilo tener que calcular esto.

_{(inspirado por esta pregunta )}

Answer 1

Advertencia: @tera sugiere que esta respuesta puede ser incorrecta en cuanto a rendimiento.

La forma más eficiente de calcularlos debería ser no computarlos. Ya están disponibles si miras debajo del capó.

Verá, las GPU nVIDIA tienen registros especiales que su código (compilado) puede leer para acceder a varios tipos de información útil. Uno de estos registros contiene threadIdx.x ; otro contiene blockDim.x ; otro: el conteo del tic del reloj; y así. C ++ como lenguaje no tiene estos expuestos, obviamente; y, de hecho, tampoco lo hace CUDA. Sin embargo, la representación intermedia en la que se compila el código CUDA, denominado PTX , expone estos registros especiales (desde PTX 1.3, es decir, con versiones CUDA> = 2.1).

Afortunadamente para nosotros, dos de estos registros especiales son %warpid y %laneid , y contienen lo que sugiere su nombre.

Ahora, CUDA admite la creación de un código PTX dentro del código CUDA con la palabra clave asm , del mismo modo que se puede usar para que el código del lado del servidor emita instrucciones de ensamblaje de la CPU directamente. Con este mecanismo uno puede usar los registros especiales sin costo, es decir, sin calcular su valor. Así es cómo:

__forceinline__ __device__ unsigned lane_id()
{
    unsigned ret; 
    asm volatile ("mov.u32 %0, %laneid;" : "=r"(ret));
    return ret;
}

__forceinline__ __device__ unsigned warp_id()
{
    unsigned ret; 
    asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
    return ret;
}

Estas funciones siempre estarán en línea, así que cuando asignes

auto x = lane_id();

simplemente está asignando el valor del registro especial a la variable x .

Nota: El uso de esta función no debería terminar haciendo nada en PTX, sino que cuando usas x más adelante, el compilador debe usar el registro %lane_id directamente. En la práctica, eso no parece suceder, y obtenemos una asignación de registro, por ejemplo:

mov.u32 %r1, %laneid;

que todavía es bastante bueno

_{Si quiere un código un poco más robusto para todos los registros especiales que aún no están disponibles para usted, intente esto (está en el repositorio de mi biblioteca de descompresión GPU pero en realidad es un código de utilidad completamente genérico).}