如何查找CUDA的epsilon，最小和最大常量？

Question

我正在寻找CUDA设备的epsilon值（两个数字之间的最小步长），min（最小值）和max（最大值）。

I.E等同于gcc编译器中<float.h>中定义的FLT_EPSILON（DBL_EPSILON），FLT_MIN（DBL_MIN）和FLT_MAX（DBL_MAX）。

在某些CUDA包含文件中是否有常量？ 任何手动解释他们？任何方式来编写内核来计算它们？

在此先感谢。

Answer 1

是的，你可以自己计算这些，如果你想。关于如何计算机器epsilon的coupleexamples在维基百科页面上的C中给出;类似地，您可以通过除以/ 2来找到最小值/最大值，直到您在/溢出时为止。 （然后，您应该搜索最后一个有效值和下一个两个因素，以找到“真正的”最小/最大值，但这会给您一个很好的起点）。如果你有一个计算能力为2.0或更高的设备，那么数学主要是IEEE 754，有一些小的偏差（例如，不是所有的舍入模式都支持），但是这些偏差不足以影响基本的这些数值常数;所以你会得到5.96e-08单个的标准emach和1.11e-16的双倍; 1.175494351e-38/3.402823466e + 38的FLT_MIN/MAX和2.2250738585072014e-308/1.7976931348623158e + 308的DBL_MIN/MAX。

在计算能力1.3机器上，单精度不支持非规格化数字，所以您的FLT_MIN将比CPU上的大得多。

一个计算能力2.0机器上的快速测试，以快速和肮脏的计算最小/最大：

#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; 
} 

编译/运行表明，答案是与CPU版本（除了最小值相一致;是FLT_MIN给出最小正常值而不是在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