gpt4 book ai didi

performance - CUDA 理论带宽与有效带宽

转载 作者:行者123 更新时间:2023-12-04 06:31:11 25 4
gpt4 key购买 nike

我有一个 CUDA 内核,它将两个矩阵相乘,其中 Width 和 Height 是我使用的块大小的倍数。

我使用的 Nvidia Quadro Fx 3800 的理论带宽为 50 Gb/s,我得到了一些奇怪的结果(有效带宽大于理论带宽)

我将在这里发布一些结果:

With Blocksize 2

[10][10] * [10][10] -> BW=0,02 Gb/s [1000][1000]*[1000][1000] -> BW=69,4 Gb/s

With Blocksize 64

[1000][1000] * [1000][1000] -> BW=486,4 Gb/s [10000][10000] * [10000][10000] -> BW= 45072,12 Gb/s



我从 Nvidia 最佳实践指南中获取了有效带宽公式(我已经简化了它,但它是等效的(除非有一个愚蠢的错误))。
我认为内核很好,因为它与我读过的一些 Nvidia Lectures 非常相似(如果不相等),也因为它工作正常(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);
}

提前致谢

最佳答案

我认为内核只是默默地失败了。

  • 你检查后是否有任何错误
    内核调用 ?
  • 代码有效吗?
  • 你有什么结果
    时间?
  • 关于performance - CUDA 理论带宽与有效带宽,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/5395901/

    25 4 0
    Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
    广告合作:1813099741@qq.com 6ren.com