تحسين أداء التقليص في CUDA باستخدام ذاكرة مشتركة وatomicAdd

في محاولتك لتحويل الشيفرة البرمجية السلسلية إلى نسخة تستخدم تقنية CUDA، تظهر أن هناك بعض الأخطاء التي يجب التعامل معها. يبدو أن السبب الرئيسي وراء عدم عمل الشيفرة هو استخدام atomicAdd بشكل غير صحيح.

لتصحيح هذا الخطأ، يجب تغيير طريقة استخدام atomicAdd. في الواقع، atomicAdd تستخدم لجمع قيمة معينة بشكل ذري، وليس فقط لتحديث قيمة واحدة. في حالتك، يجب عليك استخدام متغير مشترك لتخزين النتيجة النهائية وتحديثه ذريًا.

فيما يلي تعديل لكودك باستخدام متغير ذري لتخزين النتيجة:

cppCopy code
__global__ void kernel_foo(float* result, const float* in1, const float* in2, float in3, unsigned int size) {
    extern __shared__ float shared_tmp[]; // Shared memory for temporary results
    unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < size) {
        float tmp = 0.f;
        if (in2[i] > 0) {
            tmp = in1[i] / in3 - (in2[i] / in3) * (in2[i] / in3);
        }

        // Store the local result in shared memory
        shared_tmp[threadIdx.x] = tmp;

        // Synchronize threads within the block
        __syncthreads();

        // Perform reduction within the block
        for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
            if (threadIdx.x < stride) {
                shared_tmp[threadIdx.x] += shared_tmp[threadIdx.x + stride];
            }
            __syncthreads();
        }

        // The first thread in the block updates the global result atomically
        if (threadIdx.x == 0) {
            atomicAdd(result, shared_tmp[0]);
        }
    }
}

void launch_kernel_foo(float* result, const float* in1, const float* in2, float in3, unsigned int size) {
    // Initialize result to 0
    cudaMemset(result, 0, sizeof(float));

    // Launch the kernel with shared memory
    kernel_foo<<<(size + 255) / 256, 256, 256 * sizeof(float)>>>(result, in1, in2, in3, size);

    // Synchronize to ensure the kernel has finished
    cudaDeviceSynchronize();
}

تحتاج إلى تخصيص الذاكرة المشتركة باستخدام extern __shared__ لتخزين النتائج المؤقتة لكل موضوع. بعد ذلك، تقوم بإجراء عملية تقليص (reduction) داخل المجموعة (block) باستخدام الذاكرة المشتركة، ثم تقوم بتحديث النتيجة النهائية بشكل ذري.

تعتمد الشيفرة البرمجية التي قدمتها على مفهوم التقليص (reduction) داخل المجموعة (block) في CUDA. في هذا السياق، يتم تخزين نتائج مؤقتة محلية في الذاكرة المشتركة، ثم يتم تنفيذ عملية تقليص على هذه النتائج داخل المجموعة. هذا يسمح بتقليل العبء على الذاكرة العامة وتحسين أداء الكود.

في الشيفرة، يتم تخزين النتائج المؤقتة لكل موضوع في مصفوفة shared_tmp في الذاكرة المشتركة. بعد ذلك، يتم تنفيذ عملية تقليص باستخدام حلقة تكرارية تقوم بجمع النتائج في نصف الحجم في كل تكرار حتى تتم العملية على نطاق المجموعة بأكملها.

التزامن بين المواضيع داخل المجموعة يتم عبر __syncthreads() لضمان أن كل المواضيع قد قامت بتخزين النتائج المؤقتة قبل بدء عملية التقليص.

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

يرجى مراعاة أن هذا المقال يركز على التحسينات في التنفيذ باستخدام تقنيات CUDA، ويفترض أنك قد قمت بالتأكد من الشروط والقيود المتعلقة بالتعامل مع الذاكرة والتزامن في برامج CUDA.