كسر الضرب مصفوفة CUDA للمصفوفات الكبيرة

Question

لدي رمز مضاعفة المصفوفة التالي ، الذي تم تنفيذه باستخدام CUDA 3.2 و VS 2008. أقوم بتشغيل Windows Server 2008 R2 Enterprise. أقوم بتشغيل NVIDIA GTX 480. يعمل الكود التالي بشكل جيد مع قيم "العرض" (عرض المصفوفة) حتى حوالي 2500 أو نحو ذلك.

int size = Width*Width*sizeof(float);
float* Md, *Nd, *Pd;
cudaError_t err = cudaSuccess;

//Allocate Device Memory for M, N and P
err = cudaMalloc((void**)&Md, size);
err = cudaMalloc((void**)&Nd, size);
err = cudaMalloc((void**)&Pd, size);

//Copy Matrix from Host Memory to Device Memory
err = cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
err = cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);

//Setup the execution configuration
dim3 dimBlock(TileWidth, TileWidth, 1);
dim3 dimGrid(ceil((float)(Width)/TileWidth), ceil((float)(Width)/TileWidth), 1);

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

//Free Device Memory
cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);

عندما أقوم بتعيين "العرض" على 3000 أو أكثر ، أحصل على الخطأ التالي بعد شاشة سوداء:

نظرت إلى الإنترنت ورأيت أن بعض الأشخاص لديهم هذه القضية لأن الوكالة الدولية للطاقة كانت تقتل النواة بعد معلقها لأكثر من 5 ثوان. حاولت تحرير "Tdrdelay" في السجل وهذا ما أدى إلى تأخير الوقت قبل ظهور الشاشة السوداء ونفس الخطأ. لذلك خلصت إلى أن هذه لم تكن مشكلتي.

لقد قمت بتصحيح الأخطاء في الكود الخاص بي ووجدت أن هذا السطر هو الجاني:

err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

هذا ما أستخدمه لإرجاع مجموعة النتيجة الخاصة بي من الجهاز بعد استدعاء وظيفة kernel مضاعفة المصفوفة. كل شيء حتى هذه النقطة يبدو أنه جيد. أعتقد أنني أقوم بتخصيص الذاكرة بشكل صحيح ولا يمكنني معرفة سبب حدوث ذلك. اعتقدت أنه ربما لم يكن لدي ما يكفي من الذاكرة على بطاقتي لهذا ، ولكن بعد ذلك لا ينبغي أن يعيد Cudamalloc خطأ؟ (أكدت ذلك لم يحدث أثناء تصحيح الأخطاء).

أي أفكار/مساعدة ستكون موضع تقدير كبير! ... شكرًا جزيلاً يا رفاق !!

رمز النواة:

//Matrix Multiplication Kernel - Multi-Block Implementation
__global__ void MatrixMultiplicationMultiBlock_Kernel (float* Md, float* Nd, float* Pd, int Width) 
{
int TileWidth = blockDim.x;

//Get row and column from block and thread ids
int Row = (TileWidth*blockIdx.y) + threadIdx.y;
int Column = (TileWidth*blockIdx.x) + threadIdx.x;

//Pvalue store the Pd element that is computed by the thread
float Pvalue = 0;

for (int i = 0; i < Width; ++i)
{
    float Mdelement = Md[Row * Width + i];
    float Ndelement = Nd[i * Width + Column];
    Pvalue += Mdelement * Ndelement;
}

//Write the matrix to device memory each thread writes one element
Pd[Row * Width + Column] = Pvalue;
}

لدي أيضًا هذه الوظيفة الأخرى التي تستخدم الذاكرة المشتركة ، كما أنها تعطي نفس الخطأ:

مكالمة:

            MatrixMultiplicationSharedMemory_Kernel<<<dimGrid, dimBlock, sizeof(float)*TileWidth*TileWidth*2>>>(Md, Nd, Pd, Width);

رمز النواة:

 //Matrix Multiplication Kernel - Shared Memory Implementation
 __global__ void MatrixMultiplicationSharedMemory_Kernel (float* Md, float* Nd, float* Pd, int Width) 
 {
int TileWidth = blockDim.x;

//Initialize shared memory
extern __shared__ float sharedArrays[];
float* Mds = (float*) &sharedArrays;
float* Nds = (float*) &Mds[TileWidth*TileWidth];

int tx = threadIdx.x;
int ty = threadIdx.y;

//Get row and column from block and thread ids
int Row = (TileWidth*blockIdx.y) + ty;
int Column = (TileWidth*blockIdx.x) + tx;
float Pvalue = 0;

//For each tile, load the element into shared memory
for( int i = 0; i < ceil((float)Width/TileWidth); ++i)
{
    Mds[ty*TileWidth+tx] = Md[Row*Width + (i*TileWidth + tx)];
    Nds[ty*TileWidth+tx] = Nd[(ty + (i * TileWidth))*Width + Column]; 

    __syncthreads();

    for( int j = 0; j < TileWidth; ++j)
    {
        Pvalue += Mds[ty*TileWidth+j] * Nds[j*TileWidth+tx];
    }

    __syncthreads();
}

//Write the matrix to device memory each thread writes one element
Pd[Row * Width + Column] = Pvalue;
}

Answer 1

السيطرة على مهلة WDDM

المشكلة هي في الواقع النواة وليس cudaMemcpy(). عندما تقوم بتشغيل kernel ، تنفجر GPU وتفعل العمل بشكل غير متزامن مع وحدة المعالجة المركزية ، لذلك فقط عندما تتم مزامنة مع وحدة معالجة الرسومات التي يجب أن تنتظرها حتى ينتهي العمل. cudaMemcpy() ينطوي على تزامن ضمني ، وبالتالي هذا هو المكان الذي ترى فيه المشكلة.

يمكنك التحقق المزدوج عن طريق الاتصال cudaThreadSynchronize() بعد kernel وستظهر المشكلة على cudaThreadSynchronize() بدلا من ال cudaMemcpy().

بعد تغيير مهلة TDR ، هل قمت بإعادة تشغيل جهازك؟ لسوء الحظ ، يجب إعادة تشغيل Windows لتغيير إعدادات TDR. وثيقة Microsoft هذه لديه وصف جيد إلى حد ما للإعدادات الكاملة المتاحة.

مشاكل نواة

في هذه الحالة ، فإن المشكلة ليست في الواقع مهلة WDDM. هناك أخطاء في النواة التي ستحتاج إلى حلها (على سبيل المثال ، يجب أن تكون قادرًا على الزيادة i بأكثر من واحد في كل تكرار) والتحقق من matrixMul قد تكون العينة في SDK مفيدة. بالمناسبة ، آمل أن يكون هذا تمرينًا تعليميًا لأنه في الواقع سيكون أفضل حالًا (للأداء) باستخدام Cublas لأداء مضاعفة المصفوفة.

المشكلة الأكثر أهمية في الكود هي أنك تستخدم الذاكرة المشتركة دون تخصيص أي بالفعل. في kernel الخاص بك:

//Initialize shared memory
extern __shared__ float sharedArrays[];

ولكن عند تشغيل kernel ، لا تحدد مقدار الذاكرة المشتركة لتخصيصها لكل كتلة:

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

بناء الجملة << >>> يأخذ في الواقع أربع وسيطات حيث يكون الثالث والرابع اختياريين. الرابع هو فهرس الدفق الذي يتم استخدامه للتداخل بين الحساب ونقل البيانات (ولتنفيذ النواة المتزامنة) ولكن الثالث تحدد الوسيطة مقدار الذاكرة المشتركة لكل كتلة. في هذه الحالة أفترض أنك تريد تخزينها TileWidth * TileWidth يطفو في الذاكرة المشتركة ، لذلك يمكنك استخدام:

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock, dimBlock.x * dimBlock.x * sizeof(float)>>>(Md, Nd, Pd, Width);

المشكلة الأساسية

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

كان ينبغي أن اقترح هذا في وقت سابق ، ولكن غالبًا ما يكون من المفيد الركض cuda-memcheck للتحقق من حدوث انتهاكات وصول ميم مثل هذا.

Answer 2

يجب عليك تغيير إعدادات مهلة برنامج التشغيل ، هل ميزة Windows لمنع برامج التشغيل المعيبة لجعل النظام غير مستجيب. افحص ال صفحة Microsoft وصف كيفية القيام بذلك.

Answer 3

يجب عليك أيضًا التحقق من إعداد علامة "المهلة" على جهاز GPU الخاص بك. إذا تم تثبيت CUDA SDK ، أعتقد أن تطبيق "DeviceQuery" سيقوم بالإبلاغ عن هذه الخاصية.