¿Cómo encontrar las constantes epsilon, mínimas y máximas para CUDA?

Question

Estoy buscando los valores de epsilon (el paso más pequeño entre dos números), min (la magnitud más pequeña) y max (la mayor magnitud) para dispositivos CUDA.

I.E los equivalentes a FLT_EPSILON (DBL_EPSILON), FLT_MIN (DBL_MIN) y FLT_MAX (DBL_MAX) definidos en <float.h> en compiladores gcc.

¿Hay constantes en algún archivo de inclusión de CUDA? ¿Algún manual que los explique? ¿Alguna forma de escribir un kernel para calcularlos?

Gracias de antemano.

Answer 1

Sí, ciertamente podría calcularlas usted mismo si quisiera. Un coupleexamples para saber cómo calcular la máquina épsilon se da en C en la página de wikipedia; de forma similar, puede encontrar el mínimo/máximo dividiendo/multiplicando por dos hasta que suba/desborde. (Luego debe buscar entre el último valor válido y el siguiente factor de dos para encontrar el valor "verdadero" mínimo/máximo, pero esto le proporciona un buen punto de partida). No obstante, si tiene un dispositivo con una capacidad de cálculo de 2.0 o superior, las matemáticas son en su mayoría IEEE 754 con algunas desviaciones pequeñas (por ejemplo, no todos los modos de redondeo son compatibles) pero esas desviaciones no son suficientes para afectar a los fundamentales. constantes numéricas como estas; por lo que obtendrá el emach estándar para el único de 5.96e-08 y el doble de 1.11e-16; FLT_MIN/MAX de 1.175494351e-38/3.402823466e + 38, y DBL_MIN/MAX de 2.2250738585072014e-308/1.7976931348623158e + 308.

En la capacidad de cómputo 1.3 máquinas, los números desnormalizados no eran compatibles con una sola precisión, por lo que su FLT_MIN sería significativamente mayor que en la CPU.

Una prueba rápida en una máquina de capacidad de cálculo 2.0, con cálculos rápidos y sucios para min/max:

#include <stdio.h> 
#include <stdlib.h> 
#include <getopt.h> 
#include <cuda.h> 
#include <sys/time.h> 
#include <math.h> 
#include <assert.h> 
#include <float.h> 

#define CHK_CUDA(e) {if (e != cudaSuccess) {fprintf(stderr,"Error: %s\n", cudaGetErrorString(e)); exit(-1);}} 

/* from wikipedia page, for machine epsilon calculation */ 
/* assumes mantissa in final bits */ 
__device__ double machine_eps_dbl() { 
    typedef union { 
     long long i64; 
     double d64; 
    } dbl_64; 

    dbl_64 s; 

    s.d64 = 1.; 
    s.i64++; 
    return (s.d64 - 1.); 
} 

__device__ float machine_eps_flt() { 
    typedef union { 
     int i32; 
     float f32; 
    } flt_32; 

    flt_32 s; 

    s.f32 = 1.; 
    s.i32++; 
    return (s.f32 - 1.); 
} 

#define EPS 0 
#define MIN 1 
#define MAX 2 

__global__ void calc_consts(float *fvals, double *dvals) { 

    int i = threadIdx.x + blockIdx.x*blockDim.x; 
    if (i==0) { 
     fvals[EPS] = machine_eps_flt(); 
     dvals[EPS]= machine_eps_dbl(); 

     float xf, oldxf; 
     double xd, oldxd; 

     xf = 2.; oldxf = 1.; 
     xd = 2.; oldxd = 1.; 

     /* double until overflow */ 
     /* Note that real fmax is somewhere between xf and oldxf */ 
     while (!isinf(xf)) { 
      oldxf *= 2.; 
      xf *= 2.; 
     } 

     while (!isinf(xd)) { 
      oldxd *= 2.; 
      xd *= 2.; 
     } 

     dvals[MAX] = oldxd; 
     fvals[MAX] = oldxf; 

     /* half until overflow */ 
     /* Note that real fmin is somewhere between xf and oldxf */ 
     xf = 1.; oldxf = 2.; 
     xd = 1.; oldxd = 2.; 

     while (xf != 0.) { 
      oldxf /= 2.; 
      xf /= 2.; 
     } 

     while (xd != 0.) { 
      oldxd /= 2.; 
      xd /= 2.; 
     } 

     dvals[MIN] = oldxd; 
     fvals[MIN] = oldxf; 

    } 
    return; 
} 

int main(int argc, char **argv) { 
    float fvals[3]; 
    double dvals[3]; 
    float *fvals_d; 
    double *dvals_d; 

    CHK_CUDA(cudaMalloc(&fvals_d, 3*sizeof(float))); 
    CHK_CUDA(cudaMalloc(&dvals_d, 3*sizeof(double))); 

    calc_consts<<<1,32>>>(fvals_d, dvals_d); 

    CHK_CUDA(cudaMemcpy(fvals, fvals_d, 3*sizeof(float), cudaMemcpyDeviceToHost)); 
    CHK_CUDA(cudaMemcpy(dvals, dvals_d, 3*sizeof(double), cudaMemcpyDeviceToHost)); 

    CHK_CUDA(cudaFree(fvals_d)); 
    CHK_CUDA(cudaFree(dvals_d)); 

    printf("Single machine epsilon:\n"); 
    printf("CUDA = %g, CPU = %g\n", fvals[EPS], FLT_EPSILON); 
    printf("Single min value (CUDA - approx):\n"); 
    printf("CUDA = %g, CPU = %g\n", fvals[MIN], FLT_MIN); 
    printf("Single max value (CUDA - approx):\n"); 
    printf("CUDA = %g, CPU = %g\n", fvals[MAX], FLT_MAX); 

    printf("\nDouble machine epsilon:\n"); 
    printf("CUDA = %lg, CPU = %lg\n", dvals[EPS], DBL_EPSILON); 
    printf("Double min value (CUDA - approx):\n"); 
    printf("CUDA = %lg, CPU = %lg\n", dvals[MIN], DBL_MIN); 
    printf("Double max value (CUDA - approx):\n"); 
    printf("CUDA = %lg, CPU = %lg\n", dvals[MAX], DBL_MAX); 

    return 0; 
} 

Compilando/marcha muestra que las respuestas son consistentes con la versión CPU (excepto los valores min; ¿FLT_MIN da el valor normal mínimo en lugar de denormado en la CPU?)

$ nvcc -o foo foo.cu -arch=sm_20 
$ ./foo 
Single machine epsilon: 
CUDA = 1.19209e-07, CPU = 1.19209e-07 
Single min value (CUDA - approx): 
CUDA = 1.4013e-45, CPU = 1.17549e-38 
Single max value (CUDA - approx): 
CUDA = 1.70141e+38, CPU = 3.40282e+38 

Double machine epsilon: 
CUDA = 2.22045e-16, CPU = 2.22045e-16 
Double min value (CUDA - approx): 
CUDA = 4.94066e-324, CPU = 2.22507e-308 
Double max value (CUDA - approx): 
CUDA = 8.98847e+307, CPU = 1.79769e+308