CUDA编程入门介绍了CUDA编程的基本编程方法，首先介绍了CUDA模型的架构，然后从CUDA的内存模型和线程模型分别进

CUDA简要介绍

CUDA是建立在Nvidia GPU上的一套计算平台，并扩展出多语言的支持

CUDA C是标准ANSI C语言的扩展，扩展除一些语法和关键字来编写设备端代码，而且CUDA库本身提供了大量API来操作设备完成计算。

CUDA驱动API
CUDA运行时API

一个CUDA应用通常可以分解为两部分，

CPU主机端代码，
GPU设备端代码 CUDA nvcc编译器会自动分离代码里的不同部分，主机代码用C编写，使用本地的C语言编译器编译，设备端代码，也就是核函数，用CUDA C编写，通过nvcc编译，链接阶段，在内核程序调用或者明显的GPU设备操作时，添加运行时库。 核函数是我们后面主要接触的代码，也就是GPU上执行的程序段

CUDA “Hello World”

Hello World是所有编程语言的入门程序，对于CUDA学习我们也来完成一个Hello World。之前GPU是不能printf的，我当时就很懵，GPU是个做显示的设备，为啥不能输出，后来就可以直接在CUDA核里面打印信息了，我们写下面程序

/*
*hello_world.cu
*/
#include<stdio.h>
__global__ void hello_world(void)
{
  printf("GPU: Hello world!\n");
}
int main(int argc,char **argv)
{
  printf("CPU: Hello world!\n");
  hello_world<<<1,10>>>();
  cudaDeviceReset();//if no this line ,it can not output hello world from gpu
  return 0;
}

简单介绍其中几个关键字 __global__：用来告诉编译器这个是个可以在设备上执行的和函数

hello_world<<<1,10>>>();这句C语言中的“<<<>>>”是对设备进行配置的参数，也是CUDA扩展出来的部分。

cudaDeviceReset();这句话是对主机和设备进行同步的语句，原本GPU和CPU执行程序时异步的，核函数调用后不等到核函数执行完成，程序立刻回到主机线程继续执行。

一般的CUDA程序分成下面这些步骤：

分配GPU内存
拷贝内存到设备
调用CUDA内核函数来执行计算
把计算完成的数据拷贝回主机端
内存销毁

但是要编写高性能的CUDA程序可没这么简单，CUDA中有两个模型时决定性能的：

内存层次结构
线程层次结构

CUDA编程模型

编程模型可以理解为，我们编程要用到的语法，内存结构，线程结构等这些我们写程序时我们自己控制的部分。 GPU中大致可以分为：

核函数
内存管理
线程管理
流

CUDA编程结构

一个异构环境，通常有多个CPU多个GPU，他们都通过PCIE总线相互通信，也是通过PCIE总线分割开的。所以我们要区分一下两种设备的内存：

主机：CPU及其内存
设备：GPU及其内存

注意这两个内存从硬件到软件都是隔离的（CUDA6.0 以后支持统一寻址），我们目前先不研究统一寻址，我们现在还是用内存来回拷贝的方法来编写调试程序，以巩固大家对两个内存隔离这个事实的理解。

一个完整的CUDA应用可能的执行顺序如下图：

内存管理

内存管理在传统串行程序是非常常见的，CUDA程序也同样，只是CUDA提供的API可以分配管理设备上的内存，也可以用CUDA管理主机上的内存，主机上的传统标准库也能完成主机内存管理。下面表格有一些主机API和CUDA C的API对比：

标准C函数	CUDA C 函数	说明
malloc	cudaMalloc	内存分配
memcpy	cudaMemcpy	内存复制
memset	cudaMemset	内存设置
free	cudaFree	释放内存
我们先研究最关键得一步，这一步要走总线的

cudaError_t cudaMemcpy(void * dst,const void * src,size_t count,
  cudaMemcpyKind kind)

这个函数是内存拷贝过程，可以完成以下几种过程(cudaMemcpyKind kind)

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice 这四个过程的方向可以清楚的从字面上看出来，如果函数执行成功，则会返回cudaSuccess否则返回cudaErrorMemoryAllocation

调用下面这个函数可以把错误码转换为错误详细信息：

char* cudaGetErrorString(cudaError_t error)

内存是分层次的，下图可以简单地描述，但是不够准确，后续会写文章介绍每一个具体的环节：

共享内存（shared Memory）和全局内存（global Memory）后面我们会特别详细深入的研究，这里我们来个向量加法的例子：代码库：github.com/Tony-Tan/CU…

/*
* https://github.com/Tony-Tan/CUDA_Freshman
* 3_sum_arrays
*/
#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"


void sumArrays(float * a,float * b,float * res,const int size)
{
  for(int i=0;i<size;i+=4)
  {
    res[i]=a[i]+b[i];
    res[i+1]=a[i+1]+b[i+1];
    res[i+2]=a[i+2]+b[i+2];
    res[i+3]=a[i+3]+b[i+3];
  }
}
__global__ void sumArraysGPU(float*a,float*b,float*res)
{
  int i=threadIdx.x;
  res[i]=a[i]+b[i];
}
int main(int argc,char **argv)
{
  int dev = 0;
  cudaSetDevice(dev);

  int nElem=32;
  printf("Vector size:%d\n",nElem);
  int nByte=sizeof(float)*nElem;
  float *a_h=(float*)malloc(nByte);
  float *b_h=(float*)malloc(nByte);
  float *res_h=(float*)malloc(nByte);
  float *res_from_gpu_h=(float*)malloc(nByte);
  memset(res_h,0,nByte);
  memset(res_from_gpu_h,0,nByte);

  float *a_d,*b_d,*res_d;
  CHECK(cudaMalloc((float**)&a_d,nByte));
  CHECK(cudaMalloc((float**)&b_d,nByte));
  CHECK(cudaMalloc((float**)&res_d,nByte));

  initialData(a_h,nElem);
  initialData(b_h,nElem);

  CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));

  dim3 block(nElem);
  dim3 grid(nElem/block.x);
  sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d);
  printf("Execution configuration<<<%d,%d>>>\n",block.x,grid.x);

  CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));
  sumArrays(a_h,b_h,res_h,nElem);

  checkResult(res_h,res_from_gpu_h,nElem);
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);

  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);

  return 0;
}

线程管理

内核函数执行时，一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程，这种分层组织结构使得我们的并行过程更加灵活：一个线程块block中可以完成下述协作：

同步
共享内存

不同块内的线程不能相互影响，他们是物理隔离的！

需要给每一个线程分配一个编号，每个线程都执行同样的一段串行代码，怎么让这段相同的代码对应不同的数据？得让这些线程批次区分开，才能对应到相应的线程。主要依靠下面的两个内置结构体确定线程标号：

blockIdx（线程块在线程网格内的位置索引）
threadIdx（线程在线程块内的位置索引）

这两个内置结构体基于uint3定义，包含三个无符号整数的结构，通过三个字段来指定：

blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z

上面这两个是坐标，我们还需要同样对应的两个结构体来保存其范围，也就是threadIdx和blockIdx中三个字段的范围：

blockDim
gridDim

它们是dim3类型（基于uint3定义的数据结构）的变量，也包含三个字段x，y，z。

blockDim.x
blockDim.y
blockDim.z

网格和块的维度一般是二维和三维的，也就是说一个网格通常被分成二维的块，而每个块常被分成三维的线程。 dim3是手工定义的，主机端可见。uint3是设备端在执行的时候可见的，不可以在核函数运行时修改，初始化完成后uint3值就不变了。他们是有区别的！这一点必须要注意。下面有一段代码，表明块的索引和维度的关系：

/*
*1_check_dimension
*/
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void checkIndex(void)
{
  printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)\
  gridDim(%d,%d,%d)\n",threadIdx.x,threadIdx.y,threadIdx.z,
  blockIdx.x,blockIdx.y,blockIdx.z,blockDim.x,blockDim.y,blockDim.z,
  gridDim.x,gridDim.y,gridDim.z);
}
int main(int argc,char **argv)
{
  int nElem=6;
  dim3 block(3);
  dim3 grid((nElem+block.x-1)/block.x);
  printf("grid.x %d grid.y %d grid.z %d\n",grid.x,grid.y,grid.z);
  printf("block.x %d block.y %d block.z %d\n",block.x,block.y,block.z);
  checkIndex<<<grid,block>>>();
  cudaDeviceReset();
  return 0;
}

可以运行得到不同线程分解方式

接下来这段代码是检查网格和块的大小的：

/*
*2_grid_block
*/
#include <cuda_runtime.h>
#include <stdio.h>
int main(int argc,char ** argv)
{
  int nElem=1024;
  dim3 block(1024);
  dim3 grid((nElem-1)/block.x+1);
  printf("grid.x %d block.x %d\n",grid.x,block.x);

  block.x=512;
  grid.x=(nElem-1)/block.x+1;
  printf("grid.x %d block.x %d\n",grid.x,block.x);

  block.x=256;
  grid.x=(nElem-1)/block.x+1;
  printf("grid.x %d block.x %d\n",grid.x,block.x);

  block.x=128;
  grid.x=(nElem-1)/block.x+1;
  printf("grid.x %d block.x %d\n",grid.x,block.x);

  cudaDeviceReset();
  return 0;
}

CUDA核函数介绍

核函数就是在CUDA模型上诸多线程中运行的那段串行代码，这段代码在设备上运行，用NVCC编译，产生的机器码是GPU的机器码，所以我们写CUDA程序就是写核函数，第一步我们要确保核函数能正确的运行产生正确的结果，第二优化CUDA程序的部分，无论是优化算法，还是调整内存结构，线程结构都是要调整核函数内的代码，来完成这些优化的。
我们一直把我们的CPU当做一个控制者，运行核函数，要从CPU发起。

启动核函数

通过下面的ANSI C扩展出的CUDA C指令：

kernel_name<<<grid,block>>>(argument list);

其标准C的原型就是C语言函数调用

function_name(argument list);

这三个尖括号'<<<grid, block>>>'内是对设备代码执行的线程结构的配置（或者简称为对内核进行配置），也是前面提到的线程结构中的网格，块。我们通过CUDA C内置的数据类型dim3类型的变量来配置grid和block。

通过指定grid和block的维度，我们可以配置：

内核中线程的数目
内核中使用的线程布局

我们可以使用dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化：

kernel_name<<<4,8>>>(argument list);

上面这条指令的线程布局是：

核函数是同时复制到多个线程执行的，为了让多线程按照我们的需求对应到不同的数据，就要给线程一个唯一的标识。由于设备内存是线性的，可以用threadIdx.x和blockIdx.x来组合获得对应线程的唯一标识，threadIdx和blockIdx可以组合出很多不一样的效果。

改变核函数的配置，产生运行结果一样，但效率不同的代码： 1.一个块：

kernel_name<<<1,32>>>(argument list);

2.32个块：

kernel_name<<<32,1>>>(argument list);

上述代码如果在核函数中没有特殊结构，执行结果应该一致，但是有些效率会一直比较低。

当主机启动了核函数，控制权马上回到主机，而不是主机等待设备完成核函数的运行（前面提到过设备代码是异步执行的）。

想要主机等待设备端执行完毕可以用下面这个指令：

cudaError_t cudaDeviceSynchronize(void);

这是一个显示的方法，对应的也有隐式方法。隐式方法就是不明确说明主机要等待设备端，二十设备端不执行完，主机没办法执行，比如内存拷贝函数：

cudaError_t cudaMemcpy(void* dst,const void * src,
  size_t count,cudaMemcpyKind kind);

当核函数启动后的下一条指令就是从设备复制数据回主机端，那么主机端必须要等待设备端计算完成。

所有CUDA核函数的启动都是异步的，这点与C语言是完全不同的

编写核函数

核函数是我们自己声明定义的，声明核函数有一个比较模板化的方法：

__global__ void kernel_name(argument list);

声明和定义是不同的，这点CUDA C与C语言是一致的 在C语言函数前没有的限定符global，CUDA C中还有一些其他我们在C中没有的限定符，如下：

限定符	执行	调用	备注
global	设备端执行	可以从主机调用也可以从计算能力3以上的设备调用	必须有一个void的返回类型
device	设备端执行	设备端调用
host	主机端执行	主机调用	可以省略

这里有个特殊的情况是有些函数可以同时被定义为device和host，这种函数可以同时被设备和主机端的代码调用，主机端代码调用函数很平常，设备端调用函数与C语言一致，但是要申明成设备端代码，告诉nvcc编译成设备机器码，同时声明主机端设备端函数，那么就要告诉编译器，生成两份不同设备的机器码。

Kernel核函数编写有以下限制

只能访问设备内存
必须有void返回类型
不支持可变数量的参数
不支持静态变量
显示异步行为

并行程序中经常的一种现象，把串行代码并行化时对串行代码块for的操作，也就是把for并行化。例如：

串行：

void sumArraysOnHost(float *A, float *B, float *C, const int N) {
    for (int i = 0; i < N; i++)
        C[i] = A[i] + B[i];
}

并行：

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
  int i = threadIdx.x;
  C[i] = A[i] + B[i];
}

验证核函数

验证核函数就是验证其正确性，下面这段代码上文出现过，但是同样包含验证核函数的方法：代码库：github.com/Tony-Tan/CU…

/*
* https://github.com/Tony-Tan/CUDA_Freshman
* 3_sum_arrays
*/
#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"


void sumArrays(float * a,float * b,float * res,const int size)
{
  for(int i=0;i<size;i+=4)
  {
    res[i]=a[i]+b[i];
    res[i+1]=a[i+1]+b[i+1];
    res[i+2]=a[i+2]+b[i+2];
    res[i+3]=a[i+3]+b[i+3];
  }
}
__global__ void sumArraysGPU(float*a,float*b,float*res)
{
  int i=threadIdx.x;
  res[i]=a[i]+b[i];
}
int main(int argc,char **argv)
{
  int dev = 0;
  cudaSetDevice(dev);

  int nElem=32;
  printf("Vector size:%d\n",nElem);
  int nByte=sizeof(float)*nElem;
  float *a_h=(float*)malloc(nByte);
  float *b_h=(float*)malloc(nByte);
  float *res_h=(float*)malloc(nByte);
  float *res_from_gpu_h=(float*)malloc(nByte);
  memset(res_h,0,nByte);
  memset(res_from_gpu_h,0,nByte);

  float *a_d,*b_d,*res_d;
  CHECK(cudaMalloc((float**)&a_d,nByte));
  CHECK(cudaMalloc((float**)&b_d,nByte));
  CHECK(cudaMalloc((float**)&res_d,nByte));

  initialData(a_h,nElem);
  initialData(b_h,nElem);

  CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));

  dim3 block(nElem);
  dim3 grid(nElem/block.x);
  sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d);
  printf("Execution configuration<<<%d,%d>>>\n",block.x,grid.x);

  CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));
  sumArrays(a_h,b_h,res_h,nElem);

  checkResult(res_h,res_from_gpu_h,nElem);
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);

  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);

  return 0;
}

在开发阶段，每一步都进行验证是绝对高效的，比把所有功能都写好，然后进行测试这种过程效率高很多，同样写CUDA也是这样的每个代码小块都进行测试，看起来慢，实际会提高很多效率。
CUDA小技巧，当我们进行调试的时候可以把核函数配置成单线程的：

kernel_name<<<1,1>>>(argument list)

错误处理

所有编程都需要对错误进行处理，早期的编码错误，编译器会帮我们搞定，内存错误也能观察出来，但是有些逻辑错误很难发现，甚至到了上线运行时才会被发现，而且有些厉害的bug复现会很难，不总出现，但是很致命，而且CUDA基本都是异步执行的，当错误出现的时候，不一定是哪一条指令触发的，这一点非常头疼；这时候我们就需要对错误进行防御性处理了，例如我们代码库头文件里面的这个宏：

#define CHECK(call)\
{\
  const cudaError_t error=call;\
  if(error!=cudaSuccess)\
  {\
      printf("ERROR: %s:%d,",__FILE__,__LINE__);\
      printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
      exit(1);\
  }\
}

就是获得每个函数执行后的返回结果，然后对不成功的信息加以处理，CUDA C 的API每个调用都会返回一个错误代码，这个代码我们就可以好好利用了，当然在release版本中可以去除这部分，但是开发的时候一定要有的。

编译执行

编译指令非常简单，只需输入下面的代码即可：

nvcc sum_arrays.cu -o sum_arrays

在我的Windows电脑下，运行结果如下：

组织并行线程

前面介绍每一个线程时怎么确定唯一的索引，然后建立并行计算，并且不同的线程组织形式是怎样影响性能的：

二维网格二维线程块
一维网格一维线程块
二维网格一维线程块

使用块和线程建立矩阵索引

对于线程模型，前文已经有了大概的介绍，但是下图可以非常形象的反映线程模型，不过注意硬件实际的执行和存储不是按照图中的模型来的：

这里的(ix, iy)是整个线程模型中任意一个线程的索引，或者叫做全局地址，局部地址就是(threadIdx.x, threadIdx.y)了，整个局部地址目前还没什么用处，他只能索引线程块内的线程，不同线程块中有相同的局部索引值，比如同一个小区，A栋有16楼，B栋也有16楼，A栋和B栋就是blockIdx，而16就是threadIdx。

图中横坐标是：

ix = threadIdx.x + blockIdx.x \times blockDim.x

纵坐标是：

iy = threadIdx.y + blockIdx.y \times blockDim.y

这样我们得到了每个线程的唯一标号，并且在运行时kernel是可以访问这个标号的。前面讲过CUDA每一个线程执行相同的代码，CUDA常用的做法是让不同的线程对应不同的数据，也就是用线程的全局标号对应不同组的数据。

设备内存或者主机内存都是线性存在的，比如一个二维矩阵(8×6)，存储在内存中是这样的：

我们要做管理的就是：

线程和块索引(来计算线程的全局索引)
矩阵中给定点的坐标(ix, iy)
(ix, iy)对应的线性内存的位置

线性位置的计算方法是：

idx = ix + iy \times nx

通过上面的式子计算出线程的全局坐标，用线程的全局坐标对应矩阵的坐标，也就是说，线程的坐标(ix, iy)对应矩阵中(ix, iy)的元素，这样形成了一一对应，不同的线程处理矩阵中不同的数据。接下来的代码来输出每个线程的标号信息：

#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.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;
}

这段代码输出了一组我们随机生成的矩阵，并且核函数打印自己的线程标号，注意，核函数能调用printf这个特性是CUDA后来加的，最早的版本里面不能printf，输出结果：由于截图不完全，上面有一段打印信息没贴全，但是我们可以知道每一个线程已经对应到了不同的数据，接着我们就要用这个方法来进行计算了，最简单的当然就是二维矩阵加法啦。

二维矩阵加法

定义核函数：

__global__ void sumMatrix(float * MatA,float * MatB,float * MatC,int nx,int ny)
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx=ix+iy*ny;
    if (ix<nx && iy<ny)
    {
      MatC[idx]=MatA[idx]+MatB[idx];
    }
}

下面我们调整不同的线程组织形式，测试一下不同的效率并保证得到正确的结果，但是什么时候得到最好的效率是后面要考虑的，我们要做的就是用各种不同的相乘组织形式得到正确结果.

二维网格和二维块

代码：

// 2d block and 2d grid
dim3 block_0(dimx,dimy);
dim3 grid_0((nx-1)/block_0.x+1,(ny-1)/block_0.y+1);
iStart=cpuSecond();
sumMatrix<<<grid_0,block_0>>>(A_dev,B_dev,C_dev,nx,ny);
CHECK(cudaDeviceSynchronize());
iElaps=cpuSecond()-iStart;
printf("GPU Execution configuration<<<(%d,%d),(%d,%d)>>> Time elapsed %f sec\n",
      grid_0.x,grid_0.y,block_0.x,block_0.y,iElaps);
CHECK(cudaMemcpy(C_from_gpu,C_dev,nBytes,cudaMemcpyDeviceToHost));
checkResult(C_host,C_from_gpu,nxy);

运行结果：

红色框内是运行结果，用cpu写一个矩阵计算，然后比对结果，发现我们的运算结果是正确的，用时0.002152秒。

一维网格和一维块

代码：

// 1d block and 1d grid
dimx=32;
dim3 block_1(dimx);
dim3 grid_1((nxy-1)/block_1.x+1);
iStart=cpuSecond();
sumMatrix<<<grid_1,block_1>>>(A_dev,B_dev,C_dev,nx*ny ,1);
CHECK(cudaDeviceSynchronize());
iElaps=cpuSecond()-iStart;
printf("GPU Execution configuration<<<(%d,%d),(%d,%d)>>> Time elapsed %f sec\n",
      grid_1.x,grid_1.y,block_1.x,block_1.y,iElaps);
CHECK(cudaMemcpy(C_from_gpu,C_dev,nBytes,cudaMemcpyDeviceToHost));
checkResult(C_host,C_from_gpu,nxy);

运行结果：

二维网格和一维块

代码：

// 2d block and 1d grid
dimx=32;
dim3 block_2(dimx);
dim3 grid_2((nx-1)/block_2.x+1,ny);
iStart=cpuSecond();
sumMatrix<<<grid_2,block_2>>>(A_dev,B_dev,C_dev,nx,ny);
CHECK(cudaDeviceSynchronize());
iElaps=cpuSecond()-iStart;
printf("GPU Execution configuration<<<(%d,%d),(%d,%d)>>> Time elapsed %f sec\n",
      grid_2.x,grid_2.y,block_2.x,block_2.y,iElaps);
CHECK(cudaMemcpy(C_from_gpu,C_dev,nBytes,cudaMemcpyDeviceToHost));
checkResult(C_host,C_from_gpu,nxy);

运行结果：

总结

用不同的线程组织形式会得到正确结果，但是效率有所区别：

线程配置	执行时间
CPU单线程	0.060022
(128,128),(32,32)	0.002152
(524288,1),(32,1)	0.002965
(128,4096),(32,1)	0.002965

观察结果没有多大差距，但是明显比CPU快了很多，而且最主要的是我们本文用不同的线程组织模式都得到了正确结果，并且：

改变执行配置（线程组织）能得到不同的性能
传统的核函数可能不能得到最好的效果
一个给定的核函数，通过调整网格和线程块大小可以得到更好的效果