Come trovare epsilon, min e max costanti per CUDA?

Question

Sto cercando i valori di epsilon (il passo più piccolo tra due numeri), min (la più piccola grandezza) e massimo (la massima magnitudine) per i dispositivi CUDA.

I.E gli equivalenti a FLT_EPSILON (DBL_EPSILON), FLT_MIN (DBL_MIN) e FLT_MAX (DBL_MAX) definiti in <float.h> nei compilatori gcc.

Ci sono costanti in alcuni file CUDA inclusi? Qualche manuale che li spiega? Qualche modo per scrivere un kernel per calcolarli?

Grazie in anticipo.

Answer 1

Sì, si potrebbe sicuramente calcolare da soli se si desidera. A coupleexamples per come calcolare la macchina epsilon sono indicati in C nella pagina di wikipedia; allo stesso modo puoi trovare min/max dividendo/moltiplicando per due fino a quando non sei sotto/trabocco. (dovresti quindi cercare tra l'ultimo valore valido e il fattore successivo di due per trovare il valore "vero" min/max, ma questo ti dà un buon punto di partenza).

Se si dispone di un dispositivo con capacità di calcolo pari o superiore a 2,0, tuttavia, la matematica è per lo più IEEE 754 con alcune piccole deviazioni (ad esempio, non tutte le modalità di arrotondamento supportate) ma tali deviazioni non sono sufficienti per influire in modo fondamentale costanti numeriche come queste; quindi otterrai l'emach standard per singolo di 5.96e-08 e doppio di 1.11e-16; FLT_MIN/MAX di 1.175494351e-38/3.402823466e + 38 e DBL_MIN/MAX di 2.2250738585072014e-308/1.7976931348623158e + 308.

Su macchine con capacità di calcolo 1.3, i numeri denormalizzati non erano supportati in una precisione singola, quindi il tuo FLT_MIN sarebbe notevolmente maggiore rispetto alla CPU.

Un test rapido su una macchina capacità di elaborazione 2.0, con i calcoli rapido e sporco per 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; 
} 

Compilazione/corsa mostra che le risposte siano coerenti con la versione della CPU (ad eccezione dei valori min; FLT_MIN fornisce il valore normale minimo anziché denormato nella 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