Comportamientos extraños de un kernel de CUDA con bucle infinito en diferentes GPU de NVIDIA.

#include

void loop(void) {
int smid = -1;
if (threadIdx.x == 0) {
asm volatile(“mov.u32 %0, %%smid;”: “=r”(smid));
printf(“smid: %d\n”, smid);
}
while (1);
}

int main() {
loop<<<1, 32>>>();
cudaDeviceSynchronize();
return 0;
}
“`
Este es mi código fuente, el kernel imprime smid cuando el índice del hilo es 0 y luego entra en un bucle infinito. El host invoca el kernel de cuda anterior y espera por él. Realicé algunos experimentos en 2 configuraciones diferentes, como sigue:

• 1. GPU(Geforce 940M) SO(Ubuntu 18.04) MPS(Habilitado) CUDA(v11.0)
• 2. GPU(Geforce RTX 3050Ti Mobile) SO(Ubuntu 20.04) MPS(Habilitado) CUDA(v11.4)

Experimento 1: Cuando ejecuto este código en configuración 1, el sistema de GUI parece congelarse porque no se pueden observar respuestas gráficas, pero al presionar ctrl+c, este fenómeno desaparece cuando el proceso de CUDA se mata.

Experimento 2: Cuando ejecuto este código en configuración 2, el sistema parece funcionar bien sin ningún fenómeno anormal y se puede mostrar la salida de smid como smid: 2\n.

Experimento 3: Al cambiar la configuración del bloque loop<<<1, 1024>>> y ejecutar este nuevo código dos veces en configuración 2, obtengo la misma salida de smid como smid: 2\nsmid: 2\n.(En el caso de Geforce RTX 3050Ti Mobile, la cantidad de SM es 20, el número máximo de hilos por multiprocesador es 1536 y el número máximo de hilos por bloque es 1024.)

Estoy confundido con estos resultados, y aquí están mis preguntas:

• 1. ¿Por qué el sistema no muestra la salida de smid en la configuración 1?
• 2. ¿Por qué el sistema GUI parece congelarse en la configuración 1?
• 3. A diferencia del experimento 1, ¿por qué el experimento 2 muestra smid normalmente?
• 4. En el tercer experimento, la configuración del bloque alcanza los 1024 hilos, lo que significa que dos bloques diferentes no pueden programarse en el mismo SM. Bajo el entorno de MPS, todos los contextos de CUDA se fusionarán en un solo contexto de CUDA y compartirán el recurso de la GPU sin timeslice, pero ¿por qué sigo obteniendo el mismo smid en el tercer experimento? (Además, al cambiar la configuración de la cuadrícula a 10 y ejecutarlo dos veces, el smid varía de 0 a 19 y cada smid aparece solo una vez)