CUDA Matrix разрывы умножения для больших матриц - PullRequest
7 голосов
/ 30 октября 2010

У меня следующий матричный код умножения, реализованный с использованием 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 или больше, я получаю следующую ошибку после черного экрана: screenshot

Я посмотрел онлайн и увидел, что у некоторых людей есть эта проблема, потому что сторожевой таймер убивал ядро ​​после того, как оно зависало более 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;
}

Ответы [ 3 ]

10 голосов
/ 30 октября 2010

Управление тайм-аутом WDDM

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

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

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

1 голос
/ 31 октября 2010

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

0 голосов
/ 01 ноября 2010

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

...