CUDA Matrix прерывает умножение для больших матриц

Question

У меня есть следующий код умножения матриц, реализованный с использованием CUDA 3.2 и VS 2008.Я работаю на предприятии Windows Server 2008 R2.У меня 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);

Это то, что я использую для возврата набора результатов с устройства после вызова функции ядра умножения матриц.Кажется, до этого момента все шло нормально.Я считаю, что правильно распределяю память и не могу понять, почему это происходит.Я подумал, может быть, у меня на карте недостаточно памяти для этого, но разве 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().Когда вы запускаете ядро, графический процессор отключается и выполняет работу асинхронно с процессором, поэтому только при синхронизации с графическим процессором вам придется ждать завершения работы. cudaMemcpy() предполагает неявную синхронизацию, поэтому именно здесь вы видите проблему.

Вы можете проверить это, позвонив cudaThreadSynchronize() после ядра и проблема появится на cudaThreadSynchronize() вместо cudaMemcpy().

После изменения тайм-аута TDR вы перезагрузили компьютер?К сожалению, чтобы изменить настройки TDR, необходимо перезагрузить Windows. Этот документ Microsoft имеет довольно хорошее описание всех доступных настроек.

Проблемы с ядром

В этом случае проблема на самом деле не в тайм-ауте WDDM.В ядре есть ошибки, которые вам необходимо устранить (например, вы должны иметь возможность увеличивать i более чем на одну итерацию) и проверяя matrixMul образец в SDK может быть полезен.Кстати, я надеюсь, что это обучающее упражнение, поскольку на самом деле вам было бы лучше (с точки зрения производительности) использовать CUBLAS для выполнения умножения матриц.

Самая серьезная проблема в коде заключается в том, что вы используете общую память, фактически ее не выделяя.В вашем ядре у вас есть:

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

Но при запуске ядра вы не указываете, сколько общей памяти выделить для каждого блока:

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

Синтаксис <<<>>> фактически принимает четыре аргумента, третий и четвертый из которых являются необязательными.Четвертый — это индекс потока, который используется для обеспечения перекрытия между вычислениями и передачей данных (а также для одновременного выполнения ядра). третий Аргумент указывает объем общей памяти на блок.В этом случае я предполагаю, что вы хотите сохранить TileWidth * TileWidth плавает в общей памяти, поэтому вы должны использовать:

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

Главная проблема

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

Я должен был предложить это раньше, но часто бывает полезно запустить cuda-memcheck для проверки подобных нарушений доступа к памяти.

Answer 2

Вам необходимо изменить настройки тайм-аута драйвера — это функция Windows, которая предотвращает зависание системы из-за неисправных драйверов.Проверить Страница Майкрософт описывающее, как это сделать.

Answer 3

Вам также следует проверить настройку флага «тайм-аут» на вашем устройстве графического процессора.Если у вас установлен CUDA SDK, я полагаю, что приложение «deviceQuery» сообщит об этом свойстве.