¿cómo puede una función __global__ DEVOLVER un valor o BREAK como C/C++ hace

Question

Recientemente he estado haciendo trabajos de comparación de cuerdas en CUDA, y me pregunto cómo puede una función __global__ devolver un valor cuando encuentra la cadena exacta que yo Estoy buscando.

Quiero decir, necesito la función __global__ que contiene una gran cantidad de subprocesos para encontrar una cierta cadena entre un gran grupo de cadenas al mismo tiempo, y espero que una vez que se capture la cadena exacta, la función __global__ pueda detener todo los hilos y volver a la función principal, y me dice "¡lo hizo!"

Estoy usando CUDA C. ¿Cómo puedo lograr esto?

Answer 1

No hay manera en CUDA (o en las GPU NVIDIA) para que un hilo interrumpa la ejecución de todos los hilos en ejecución. No se puede tener la salida inmediata del kernel tan pronto como se encuentre un resultado, simplemente no es posible hoy.

Pero puede hacer que todos los hilos salgan tan pronto como sea posible después de que un hilo encuentre un resultado. Aquí hay un modelo de cómo lo haría.

__global___ void kernel(volatile bool *found, ...) 
{ 
    while (!(*found) && workLeftToDo()) { 

     bool iFoundIt = do_some_work(...); // see notes below 

     if (iFoundIt) *found = true; 
    } 
} 

Algunas notas sobre esto.

1. Tenga en cuenta el uso de volatile. Esto es importante.
2. ¡Asegúrese de inicializar found — que debe ser un puntero de dispositivo — a false antes de iniciar el kernel!
3. Los hilos no se cerrarán instantáneamente cuando otro hilo actualice found. Saldrán solo la próxima vez que regresen a la parte superior del ciclo while.
4. Cómo implementar do_some_work asuntos. Si es demasiado trabajo (o muy variable), entonces el retraso para salir luego de que se encuentre un resultado será largo (o variable).Si es demasiado poco trabajo, entonces sus hilos pasarán la mayor parte de su tiempo chequeando found en lugar de hacer un trabajo útil.
5. do_some_work también es responsable de asignar tareas (es decir, calcular/incrementar índices), y cómo lo hace es específico del problema.
6. Si la cantidad de bloques que ejecuta es mucho mayor que la ocupación máxima del núcleo en la presente GPU, y no se encuentra una coincidencia en la primera "oleada" de bloques de hilos, entonces este kernel (y el siguiente) puede un punto muerto Si se encuentra una coincidencia en la primera oleada, los bloques posteriores solo se ejecutarán después de found == true, lo que significa que se iniciarán y luego saldrán inmediatamente. La solución es lanzar solo tantos bloques como puedan residir simultáneamente (también conocido como "lanzamiento máximo") y actualizar su asignación de tareas en consecuencia.
7. Si el número de tareas es relativamente pequeño, puede reemplazar el while con un if y ejecutar los hilos suficientes para cubrir el número de tareas. Entonces no hay posibilidad de un punto muerto (pero se aplica la primera parte del punto anterior).
8. workLeftToDo() es específico de un problema, pero devolvería falso cuando no queda trabajo por hacer, para que no se interrumpa en el caso de que no se encuentre coincidencia.

Ahora, lo anterior puede dar lugar a un excesivo vaciado de particiones (todos los subprocesos golpean la misma memoria), especialmente en arquitecturas antiguas sin caché L1. Así que es posible que desee escribir una versión un poco más complicada, utilizando un estado compartido por bloque.

__global___ void kernel(volatile bool *found, ...) 
{ 
    volatile __shared__ bool someoneFoundIt; 

    // initialize shared status 
    if (threadIdx.x == 0) someoneFoundIt = *found; 
    __syncthreads(); 

    while(!someoneFoundIt && workLeftToDo()) { 

     bool iFoundIt = do_some_work(...); 

     // if I found it, tell everyone they can exit 
     if (iFoundIt) { someoneFoundIt = true; *found = true; } 

     // if someone in another block found it, tell 
     // everyone in my block they can exit 
     if (threadIdx.x == 0 && *found) someoneFoundIt = true; 

     __syncthreads(); 
    } 
} 

De esta manera, un hilo por encuestas en bloque la variable global, y sólo hilos que encontrar una coincidencia alguna vez escribir en él, por lo que el tráfico de memoria global se reduce al mínimo.

A un lado: las funciones __global__ son nulas porque es difícil definir cómo devolver valores de 1000 s de hilos en un solo subproceso de CPU. Es trivial para el usuario idear una matriz de retorno en el dispositivo o memoria de copia cero que se adapte a su propósito, pero difícil de hacer un mecanismo genérico.

Descargo de responsabilidad: Código escrito en el navegador, no probado, no verificado.

Answer 2

La función global realmente no contiene una gran cantidad de subprocesos como piensas. Es simplemente un kernel, función que se ejecuta en el dispositivo, que se llama al pasar los parámetros que especifican el modelo de hilo. El modelo que CUDA emplea es un modelo de cuadrícula 2D y luego un modelo de hilo 3D dentro de cada bloque en la cuadrícula.

Con el tipo de problema que tiene, no es realmente necesario usar nada además de una cuadrícula 1D con 1D de subprocesos en cada bloque porque el grupo de cadenas realmente no tiene sentido dividirse en 2D como otros problemas (por ejemplo multiplicación de matrices)

Voy a ver un ejemplo simple de decir 100 cadenas en el grupo de cadenas y desea que todas ellas se comprueben de forma paralela en lugar de secuencialmente.

//main 
//Should cudamalloc and cudacopy to device up before this code 
dim3 dimGrid(10, 1); // 1D grid with 10 blocks 
dim3 dimBlocks(10, 1); //1D Blocks with 10 threads 
fun<<<dimGrid, dimBlocks>>>(, Height) 
//cudaMemCpy answerIdx back to integer on host 

//kernel (Not positive on these types as my CUDA is very rusty 
__global__ void fun(char *strings[], char *stringToMatch, int *answerIdx) 
{ 
    int idx = blockIdx.x * 10 + threadIdx.x; 

    //Obviously use whatever function you've been using for string comparison 
    //I'm just using == for example's sake 
    if(strings[idx] == stringToMatch) 
    { 
     *answerIdx = idx 
    } 
} 

Esto obviamente no es el más eficiente y es más probable no la forma exacta para pasar parametros y trabajar con la memoria w/CUDA, pero espero que se pone el punto a través de la división de la carga de trabajo y que el 'global 'Las funciones se ejecutan en muchos núcleos diferentes, por lo que no puede decirles a todos que se detengan. Puede que haya una forma en la que no estoy familiarizado, pero la velocidad que obtendrá al dividir la carga de trabajo en el dispositivo (por supuesto, de manera sensata) ya le dará enormes mejoras de rendimiento. Para tener una idea del modelo de hilo, recomiendo leer los documentos en el sitio de Nvidia para CUDA. Te ayudarán tremendamente y te enseñarán la mejor manera de configurar la cuadrícula y los bloques para un rendimiento óptimo.

Answer 3

Si se siente aventurero, un enfoque alternativo para detener la ejecución del núcleo sería sólo hay que ejecutar

// (write result to memory here) 
__threadfence(); 
asm("trap;"); 

si no se encuentra una respuesta.

Esto no requiere memoria de sondeo, pero es inferior a la solución que Mark Harris sugirió porque hace que el kernel salga con una condición de error. Esto puede enmascarar errores reales (así que asegúrese de escribir los resultados de una manera que claramente permita distinguir entre una ejecución exitosa y un error), y puede causar otros contratiempos o disminuir el rendimiento general ya que el conductor lo trata como una excepción.

Si busca una solución segura y simple, siga la sugerencia de Mark Harris.