Теоретическая пропускная способность CUDA против эффективной пропускной способности - PullRequest
Новая
       27

Теоретическая пропускная способность CUDA против эффективной пропускной способности

2 голосов
/ 22 марта 2011

У меня есть ядро ​​CUDA, которое умножает две матрицы, в которых ширина и высота являются кратными размера используемого мной блока.

Используемая мной видеокарта Nvidia Quadro Fx 3800 имеет теоретическую пропускную способность 50 Гбит / с, и яя получаю некоторые странные результаты (эффективная пропускная способность больше, чем теоретическая пропускная способность)

Я опубликую здесь некоторые результаты:

с размером блока 2

[10] [10] * [10] [10] -> BW = 0,02 Гб / с [1000] [1000] * [1000] [1000] -> BW = 69,4 Гб / с

с размером блока 64

[1000] [1000] * [1000] [1000] -> BW = 486,4 Гбит / с [10000] [10000] * [10000][10000] -> BW = 45072,12 Гбит / с

Я взял формулу эффективной полосы пропускания из Руководства по оптимальным методам Nvidia (я упростил ее, но ее эквивалент (если только нет глупой ошибки)).Я думаю, что ядро ​​прекрасно, так как оно очень похоже (если не равно) на некоторые лекции по Nvidia, которые я читаю, а также потому, что оно работает правильно (afaik). 

#define blocksize 64
#define HM (10000) 
#define WM (10000) 
#define WN (10000)
#define HN WM 
#define WP WN   
#define HP HM  
#define PTH WM
#define PTW HM

__global__ void nonsquare(float*M, float*N, float*P, int uWM,int uWN)
   {
__shared__ float MS[blocksize][blocksize];
__shared__ float NS[blocksize][blocksize];

int tx=threadIdx.x, ty=threadIdx.y, bx=blockIdx.x, by=blockIdx.y;
int rowM=ty+by*blocksize;
int colN=tx+bx*blocksize;
int Pvalue=0;

for(int m=0; m< uWM/blocksize;m++){
    MS[ty][tx]=M[rowM*uWM+(m*blocksize+tx)];
    NS[ty][tx]=M[colN + uWN*(m*blocksize+ty)];
    __syncthreads();
    for(int k=0;k<blocksize;k++)
        Pvalue+=MS[ty][k]*NS[k][tx];
    P[rowM*WP+colN]=Pvalue;
}

}
int main(){


cudaEvent_t evstart, evstop;
cudaEventCreate(&evstart);
cudaEventCreate(&evstop);

float*M=(float*)malloc(sizeof(float)*HM*WM);
float*N=(float*)malloc(sizeof(float)*HN*WN);

for(int i=0;i<WM*HM;i++)
    M[i]=(float)i;
for(int i=0;i<WN*HN;i++)
    N[i]=(float)i;




float*P=(float*)malloc(sizeof(float)*HP*WP);



float *Md,*Nd,*Pd;
cudaMalloc((void**)&Md,HM*WM*sizeof(float));

cudaMalloc((void**)&Nd,HN*WN*sizeof(float));

cudaMalloc((void**)&Pd,HP*WP*sizeof(float));



cudaMemcpy(Md,M,HM*WM*sizeof(float),cudaMemcpyHostToDevice);

cudaMemcpy(Nd,N,HN*WN*sizeof(float),cudaMemcpyHostToDevice);



dim3 dimBlock(blocksize,blocksize);//(tile_width , tile_width);
dim3 dimGrid(WN/dimBlock.x,HM/dimBlock.y);//(width/tile_width , width/tile_witdh);

cudaEventRecord(evstart,0);

nonsquare<<<dimGrid,dimBlock>>>(Md,Nd,Pd,WM,WN);

cudaEventRecord(evstop,0);
cudaEventSynchronize(evstop);
float time;
cudaEventElapsedTime(&time,evstart,evstop);

cudaMemcpy(P,Pd,WP*HP*sizeof(float),cudaMemcpyDeviceToHost);

    cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);


    printf("\ntime spent:%f",time);
float Bandwidth=(HM*WM*4+WN*HN*4+HP*WP*4)/(time*1000000); /
printf("\nEffective Bandwidth:%f Gb/s\n",Bandwidth);
    }

Заранее спасибо

Ответы [ 3 ]

2 голосов
/ 28 марта 2011

Я думаю, что ядро ​​просто тихо выходит из строя.

  1. Проверяли ли вы какие-либо ошибки после вызов ядра?

  2. Работает ли код?

  3. Какие результаты вы имеете на время?

1 голос
/ 23 марта 2011

Обратите внимание, что при использовании разделяемой памяти, памяти текстур и т. Д. Иногда возможно превышение теоретической пропускной способности.Это часто означает, что вы подключаетесь к некоторым выделенным аппаратно-поддерживаемым функциям (таким как встроенная билинейная интерполяция текстур и т. Д.), Возможно, непреднамеренно.

Помимо причин, упомянутых Робертом Харви, существует также потенциальный заводской разгонкарты от производителей (хотя чаще для GeForce, чем для Quadros).

В целом, я бы сказал, что у вас все хорошо, если вы приближаетесь к теоретической пропускной способности или превышаете ее (либо в памяти, либо в вычислениях).

0 голосов
/ 22 марта 2011

Я могу придумать несколько объяснений:

  1. Изменения в базовом коде, которые отрицательно влияют на измерения
  2. Неправильные предположения о производительности
  3. Неопознанные микрооптимизации.
  4. Нереальные тесты.

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

...