【1 Cuda编程】

188****7632

一、使用块和线程建立矩阵索引
多线程的优点就是每个线程处理不同的数据计算，那么怎么分配好每个线程处理不同的数据，而不至于
多个不同的线程处理同一个数据，或者避免不同的线程没有组织的乱访问内存。下图说明了块和线程索
引、矩阵坐标以及线性全局内存索引之间的对应关系：

这样我们就得到了每个线程的唯一标号，并且在运行时kernel是可以访问这个标号的。为了让不同线程
处理不同的数据，CUDA常用的做法是让不同的线程对应不同的数据，也就是用线程的全局标号对应不
同组的数据。
设备内存或者主机内存都是线性存在的，比如一个二维矩阵 (8×6)(8×6)，存储在内存中是这样的：

#include <cuda_runtime.h>
#include <stdio.h>
#include "dependence.h"

__global__ void printThreadIndex(float *A,const int nx,const int ny)
{
  int ix=threadIdx.x+blockIdx.x*blockDim.x;
  int iy=threadIdx.y+blockIdx.y*blockDim.y;
  unsigned int idx=iy*nx+ix;
  printf("thread_id(%d,%d) block_id(%d,%d) coordinate(%d,%d)"
          "global index %2d ival %2d\n",threadIdx.x,threadIdx.y,
          blockIdx.x,blockIdx.y,ix,iy,idx,A[idx]);
}
int main(int argc,char** argv)
{
  initDevice(0);
  int nx=8,ny=6;
  int nxy=nx*ny;
  int nBytes=nxy*sizeof(float);

  //Malloc
  float* A_host=(float*)malloc(nBytes);
  initialData(A_host,nxy);
  printMatrix(A_host,nx,ny);

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

  cudaMemcpy(A_dev,A_host,nBytes,cudaMemcpyHostToDevice);

  dim3 block(4,2);
  dim3 grid((nx-1)/block.x+1,(ny-1)/block.y+1);

  printThreadIndex<<<grid,block>>>(A_dev,nx,ny);

  CHECK(cudaDeviceSynchronize());
  cudaFree(A_dev);
  free(A_host);

  cudaDeviceReset();
  return 0;
}

这段代码输出了一组我们随机生成的矩阵，并且核函数打印自己的线程标号，输出结果：

未完待续
参考：
https://face2ai.com/program-blog/