Как найти эпсилон, минимальную и максимальную константы для CUDA? - PullRequest
Новая
       11

Как найти эпсилон, минимальную и максимальную константы для CUDA?

13 голосов
/ 11 января 2012

Я ищу значения epsilon (наименьший шаг между двумя числами), min (наименьшая величина) и max (наибольшая величина) для устройств CUDA.

IE эквиваленты FLT_EPSILON (DBL_EPSILON), FLT_MIN (DBL_MIN) и FLT_MAX (DBL_MAX), определенные в <float.h> в компиляторах gcc.

Есть ли константы в некоторых файлах включения CUDA?Любое руководство, объясняющее их?Любой способ написать ядро ​​для их вычисления?

Заранее спасибо.

1 Ответ

15 голосов
/ 11 января 2012

Да, вы можете рассчитать это сами, если хотите. пара примеров для расчета эпсилона машины даны в C на странице википедии;Точно так же вы можете найти мин / макс, разделив / умножив на два, пока вы не опуститесь / переполнитесь.(вам следует выполнить поиск между последним действительным значением и следующим множителем двух, чтобы найти «истинное» минимальное / максимальное значение, но это даст вам хорошую отправную точку).

Если у вас есть устройство соднако можно рассчитать до 2.0 или более, тогда математика в основном IEEE 754 с некоторыми небольшими отклонениями (например, поддерживаются не все режимы округления), но этих отклонений недостаточно для воздействия на основные числовые константы, подобные этим;так что вы получите стандартный emach для сингла 5.96e-08 и double для 1.11e-16;FLT_MIN / MAX 1.175494351e-38 / 3.402823466e + 38 и DBL_MIN / MAX 2.2250738585072014e-308 / 1.7976931348623158e + 308.

При вычислительной способности 1,3 станка денормализованные числа не поддерживаются с одинарной точностьюпоэтому ваш FLT_MIN будет значительно больше, чем на процессоре.

Быстрый тест на компьютере с вычислительной способностью 2.0, с быстрыми и грязными вычислениями для 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;
}

Компиляция /выполнение показывает, что ответы соответствуют версии CPU (кроме значений min; 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
...