Una computación se considera determinista si múltiples ejecuciones con los mismos datos de entrada producen el mismo resultado bit a bit. Aunque esto parezca una propiedad sencilla de garantizar, puede ser complicado en la práctica, especialmente en la programación paralela y la aritmética de punto flotante. Esto se debe a que la suma y la multiplicación de punto flotante no son estrictamente asociativas, lo que significa que (a + b) + c puede no ser igual a a + (b + c) debido al redondeo que ocurre al almacenar resultados intermedios con precisión finita.
Con la versión 3.1 de las Bibliotecas de Cálculo de Núcleo CUDA de NVIDIA (CCCL), CUB, una biblioteca CUDA de bajo nivel para algoritmos paralelos, añadió una nueva API de fase única que acepta un entorno de ejecución. Esto permite a los usuarios personalizar el comportamiento del algoritmo. Podemos usar este entorno para configurar la propiedad de determinismo del algoritmo reduce, lo cual solo es posible a través de esta nueva API de fase única, ya que la API de dos fases no acepta un entorno de ejecución.
Niveles de determinismo en CUB
El siguiente código muestra cómo especificar el nivel de determinismo en CUB (encuentra el ejemplo completo en línea utilizando compiler explorer).
auto input = thrust::device_vector{0.0f, 1.0f, 2.0f, 3.0f}; auto output = thrust::device_vector(1); auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); // puede ser not_guaranteed, run_to_run (predeterminado) o gpu_to_gpu auto error = cub::DeviceReduce::Sum(input.begin(), output.begin(), input.size(), env); if (error != cudaSuccess) { std::cerr << "cub::DeviceReduce::Sum falló con estado: " << error << std::endl; } assert(output[0] == 6.0f);
Comenzamos especificando los vectores de entrada y salida. Luego utilizamos cuda::execution::require() para construir un objeto cuda::std::execution::env, estableciendo el nivel de determinismo a not_guaranteed.
Determinismo no garantizado
En las reducciones de punto flotante, el resultado puede depender del orden en que se combinan los elementos. Si dos ejecuciones aplican el operador de reducción en diferentes órdenes, los valores finales pueden diferir ligeramente. En muchas aplicaciones, estas diferencias menores son aceptables. Al relajar el requisito de determinismo estricto, la implementación de reducción puede reorganizar las operaciones en cualquier orden, lo que puede mejorar el rendimiento en tiempo de ejecución.
En CUB, not_guaranteed relaja el nivel de determinismo, lo que permite operaciones atómicas. Estas operaciones, cuya ejecución desordenada entre hilos resulta en un orden diferente de operaciones entre ejecuciones, pueden calcular tanto los agregados parciales a nivel de bloque como el valor final de reducción. La reducción completa también puede realizarse en un solo lanzamiento de kernel, ya que las operaciones atómicas combinan los agregados parciales a nivel de bloque en el resultado.
La variante de reducción no determinista suele ser más rápida que la versión determinista de ejecución a ejecución, especialmente para arreglos de entrada más pequeños, donde realizar la reducción en un solo kernel reduce la latencia de múltiples lanzamientos de kernel, minimiza el movimiento adicional de datos y evita sincronización adicional.
