cuda利用共享内存高效矩阵转置

PartyAnimal_w

已于 2023-12-25 16:48:41 修改

阅读量1.5k

点赞数 20

CC 4.0 BY-SA版权

文章标签：矩阵 c++

于 2023-12-25 16:21:21 首次发布

本文链接：https://blog.csdn.net/partyanimalw/article/details/135193834

文章目录

前言

共享内存的存在可以帮助我们更快速的写入和读取数据，缩短进程时间，以矩阵转置为例，记录一些自己的学习笔记。

一、参考链接

老规矩列出本文的参考链接：

CUDA - CUDA C/C++中的高效矩阵转置_cuda实现矩阵转置-CSDN博客

CUDA学习之共享内存和常量内存--part2-CSDN博客

二、矩阵转置

1.一些概念

①cuda的概念

在main函数中定义的grid和block：

int main(int argc,char** argv)
{
  printf("strating...\n");
  initDevice(0);
  int nx=1<<12;
  int ny=1<<12;
  int dimx=BDIMX;
  int dimy=BDIMY;
  int nxy=nx*ny;
  int nBytes=nxy*sizeof(float);
  int transform_kernel=0;
  if(argc==2)
    transform_kernel=atoi(argv[1]);
  if(argc>=4)
  {
      transform_kernel=atoi(argv[1]);
      dimx=atoi(argv[2]);
      dimy=atoi(argv[3]);
  }

  //Malloc
  float* A_host=(float*)malloc(nBytes);
  float* B_host_cpu=(float*)malloc(nBytes);
  float* B_host=(float*)malloc(nBytes);
  initialData(A_host,nxy);

  //cudaMalloc
  float *A_dev=NULL;
  float *B_dev=NULL;
  CHECK(cudaMalloc((void**)&A_dev,nBytes));
  CHECK(cudaMalloc((void**)&B_dev,nBytes));

  CHECK(cudaMemcpy(A_dev,A_host,nBytes,cudaMemcpyHostToDevice));
  CHECK(cudaMemset(B_dev,0,nBytes));



  // cpu compute
  double iStart=cpuSecond();
  transformMatrix2D_CPU(A_host,B_host_cpu,nx,ny);
  double iElaps=cpuSecond()-iStart;
  printf("CPU Execution Time elapsed %f sec\n",iElaps);

  // 2d block and 2d grid
  dim3 block(dimx,dimy);
  dim3 grid((nx-1)/block.x+1,(ny-1)/block.y+1);

这里用nx和ny，定义了一个2^12x2^12的矩阵；

下面用dim3定义了一个8x8的block(BDIMX和BDIMY在代码开始宏定义的）；

下面的grid被定义为512x512；

也就是说block的维度为8x8个thread，grid的尺寸为512x512个block,

那么一共含有乘起来的thread（我实在是有点懒得画图）

那么每个block负责转置2^6x2^6的数据块

也就是每个thread负责1个元素