Результаты ядра CUDA различаются в режиме выпуска

Question

Я тестирую некоторый код в CUDA (я новичок в CUDA и это мое первое приложение).На данный момент я добился результатов в CUDA, которые были получены при последовательном запуске кода на ЦП.Я использую Visual Studio 2010, а конфигурация сборки — отладочная.Но как только я меняю конфигурацию сборки на «Выпуск», я начинаю получать неправильные результаты.Я не смог пользоваться форумами Nvidia, так как они в данный момент недоступны.Может ли кто-нибудь с опытом CUDA указать на проблему.Код выглядит следующим образом

__global__ void MyKernel(int *Nptr,int *deltaptr, double *gravityptr, double *separationptr, double *fconptr, double *xForce, double *yForce, double *zForce,
double *xPos, double *yPos, double *zPos )
{
int N = *Nptr;
int delta= *deltaptr;
double gravity= *gravityptr;
double separation = *separationptr;
double fcon = *fconptr;

double len=0.0;
double r12X =0.0;
double r12Y =0.0;
double r12Z =0.0;
double PE=0.0;


int nx = blockDim.x * blockIdx.x + threadIdx.x;//use this place of nx
//int ny = blockDim.x * blockIdx.x + threadIdx.y;//use this place of ny
int ny = blockDim.y * blockIdx.y + threadIdx.y;
//printf("nx:%d ny:%d\n", nx,ny);

if(!(nx< N && ny <N))
    return;
//printf("nx:%d ny:%d\n", nx,ny);


xForce[nx*N+ny] = 0.0;
yForce[nx*N+ny] = -gravity;
zForce[nx*N+ny] = 0.0;

int lowerValuedx = maxOnDevice(nx-delta,0);
int upperValuedx=minOnDevice(nx+delta+1,N);
for(int dx=lowerValuedx; dx<upperValuedx;dx++)
{
    int lowerValuedy=maxOnDevice(ny-delta,0);
    int upperValuedy=minOnDevice(ny+delta+1,N);
    for(int dy=lowerValuedy; dy<upperValuedy;dy++)
    {
        len=sqrt((double)((nx-dx)*(nx-dx)+(ny-dy)*(ny-dy)) ) *separation;
        bool condition = ny!=dy;
        bool condition1 = nx!=dx;

        //if (nx!=dx || ny!=dy)
        if (condition || condition1)
        {
            r12X = xPos[dx*N+dy] - xPos[nx*N+ny];
            r12Y = yPos[dx*N+dy] - yPos[nx*N+ny];
            r12Z = zPos[dx*N+dy] - zPos[nx*N+ny];
            xForce[nx*N+ny] = xForce[nx*N+ny] +fcon*normxOnDevice(r12X,r12Y,r12Z)*(magOnDevice(r12X,r12Y,r12Z)-len);
            yForce[nx*N+ny]= yForce[nx*N+ny] +fcon*normyOnDevice(r12X,r12Y,r12Z)*(magOnDevice(r12X,r12Y,r12Z)-len);
            zForce[nx*N+ny]= zForce[nx*N+ny] +fcon*normzOnDevice(r12X,r12Y,r12Z)*(magOnDevice(r12X,r12Y,r12Z)-len);


        }
    }
}   

}

Спасибо

Answer 1

Нет ничего необычного в том, что числовые различия существуют между результатами ЦП и ГП, а также между отладочными и выпускными сборками на ГП.Это не означает, что какой-либо набор результатов неверен, но один может быть более точным, чем другой.См. следующий технический документ от NVIDIA, в котором обсуждаются различные механизмы, которые могут привести к числовым расхождениям:

http://developer.download.nvidia.com/assets/cuda/files/NVIDIA-CUDA-Floating-Point.pdf

Вы можете проверить, устраняет ли флаг nvcc -fmad=false различия, которые вы видите, что указывает на то, что они вызваны слиянием FMA/FMAD и, следовательно, вероятно, безвредны.

Графические процессоры обеспечивают операции FMAD и FMA (слитное умножение-сложение), которые объединяют умножение с плавающей запятой с зависимым сложением с плавающей запятой в одну операцию.Это повышает производительность, поскольку комбинированная операция обычно занимает столько же времени, сколько и каждая из ее составных частей.Однако поведение округления любой комбинированной операции отличается от использования двух операций округления по отдельности:

FMAD одинарной точности (вычислительная способность < 2,0) усекает результат умножения, а затем округляет результат окончательного сложения в соответствии со стандартом IEEE-754, округляя до ближайшего или четного.Напротив, FMA (одинарная точность при вычислительных возможностях >= 2,0 и двойная точность) вычисляет неокругленное произведение двойной ширины, добавляет к нему третий операнд и округляет окончательную сумму в соответствии с IEEE-754, округляя до ближайшего или -даже.Из-за этого единственного округления средняя точность, обеспечиваемая FMA, превосходит использование двух операций округления по отдельности.Операции FMA определены в версии стандарта IEEE-754 с плавающей запятой 2008 года.

По умолчанию для сборок выпуска компилятор CUDA активно генерирует объединенные операции (FMAD, FMA) для достижения наилучшей производительности.Другими словами, значением по умолчанию для компилятора является -fmad=true, что позволяет компилятору объединять умножения и сложения чисел с плавающей запятой.При указании -fmad=false слияние операций умножения и сложения запрещается, что обычно обеспечивает большую согласованность с результатами ЦП, поскольку большинство ЦП не поддерживают операцию FMA.Очевидно, что отключение использования объединенных операций отрицательно влияет на производительность, поэтому -fmad=false в первую очередь полезен для проверки работоспособности.

При подозрении на проблемы с точностью я обычно рекомендую сравнение с эталонной реализацией более высокой точности (например,один основан на методах четырехкратной точности или дабл-дабл), чтобы точно оценить ошибку как на ЦП, так и на графическом процессоре, а не использовать версию ЦП в качестве эталона (поскольку на результаты ЦП также влияет ошибка округления).