تختلف نتائج CUDA kernel في وضع الإصدار

Question

أقوم باختبار بعض الأكواد البرمجية في CUDA (أنا جديد على CUDA وهذا هو طلبي الأول).لقد حققت حتى الآن نتائج في CUDA وهي نفس النتائج التي تم الحصول عليها عن طريق تشغيل الكود بشكل تسلسلي على وحدة المعالجة المركزية.أنا أستخدم Visual Studio 2010 وتكوين البناء هو Debug.ولكن بمجرد أن أقوم بتغيير تكوين البناء إلى "الإصدار"، أبدأ في الحصول على نتائج خاطئة.لم أتمكن من استخدام منتديات 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 في إصدار 2008 من معيار الفاصلة العائمة IEEE-754.

بشكل افتراضي، بالنسبة لإصدارات الإصدار، يقوم مترجم CUDA بإنشاء عمليات مدمجة (FMAD، FMA) بقوة لتحقيق أفضل أداء.بمعنى آخر، الإعداد الافتراضي للمترجم هو -fmad=true والذي يسمح للمترجم بدمج مضاعفات الفاصلة العائمة وإضافتها.من خلال تحديد -fmad=false، يتم منع دمج الضربات والإضافات، مما يوفر عادةً تناسقًا أكبر مع نتائج وحدة المعالجة المركزية، نظرًا لأن معظم وحدات المعالجة المركزية لا توفر عملية FMA.من الواضح أن تعطيل استخدام العمليات المدمجة له ​​تأثير سلبي على الأداء، لذا فإن -fmad=false مفيد بشكل أساسي كفحص سلامة.

عند الاشتباه في وجود مشكلات تتعلق بالدقة، أوصي عمومًا بالمقارنة مع تطبيق مرجعي عالي الدقة (على سبيل المثال.واحدة تعتمد على الدقة الرباعية أو التقنيات المزدوجة المزدوجة) لتقييم الخطأ بدقة على كل من وحدة المعالجة المركزية ووحدة معالجة الرسومات، بدلاً من استخدام إصدار وحدة المعالجة المركزية كمرجع (حيث تتأثر نتائج وحدة المعالجة المركزية أيضًا بخطأ التقريب).