rowheighPoint.y
colwidthPoint.x
目录
1 GPU硬件架构及运行机制
文章写的很全面
2 GPU计算基础知识
CUDA编程需要CPU和GPU协同工作，程序中既包含host程序，又包含device程序，他们分别在CPU和GPU上运行，二者间可以进行通信。
host - CPU及其内存 device - GPU及其内存
2.1 kernel 核函数
- kernel是在device上线程中并行执行的函数
- 核函数用__global__符号申明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量
- 在CUDA中，每一个线程都要执行核函数，每个县城分配一个唯一的线程号（thread ID），这个ID值可以通过核函数的内置变量threadldx获得
//Kernel 定义
__global__ void vec_add(double *x, double *y, double *z, int n)
{
int i = get_tid();// 全局变量i；user-defined macro/function
if(i<n) z[i] = x[i] + y[i];// 向量相加
}
int main()
{
int N = 1000000;// 1M
int bs = 256;// 每个块有256线程
int gs = (N + bs -1)/bs;// 网格
vec_add<<<gs,bs>>>(x, y, z, N);// kernel, call GPU，指定使用的线程数和结构
}
}
2.2 程序层次结构
- 第一层：一个kernel所启动的所有线程成为一个网格（grid）
- 同一个网格上的线程，共享相同的全局内存
- dim3类型
- 第二层：网格可以分为很多线程块（block），一个线程块包含很多线程
- dim3类型
- 第三层：32线程一组为线程束（warp ）
在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。
dim3类型 包含三个无符号整数(x, y, z)成员的结构体变量，在定义时，缺省值初始化为1；1-dim、2-dim、3-dim结构均可
global：在device上执行，从host中调用，返回类型必须是void，不支持可变参数，不能成为类成员函数；异步，host会在kernel执行时进行下一步
device：在device上执行，单次可以从device中调用；不可以与__global__同时用
host：在host上执行，仅可以从host上调用，一般省略；不可以与__global__同时用；可以与__device__同时用，此时函数会在devicw和host都编译
2.3 【常用】CUDA内置变量
一个线程需要两个内置的坐标变量（blockldx， threadldx）来唯一标识，它们都是dim3类型变量。
一个线程块上的线程，放在同一个流式多处理器（SM）上
单个SM资源有限，线程块中的线程数有限制
blockldx 指明线程所在grid中的位置：blockldx.x, blockldx.y, blockldx.z
threadldx 指明线程所在block中的位置：threadldx.x ,threadldx.y, threadldx.z
gridDim 网格块各个维度的大小
blockDim 线程块各个维度的大小
对于一个2-dim的block(Dx,Dy)，线程(x,y)的ID值为(x+y∗Dx)，如果是3-dim的block(Dx,Dy,Dz)，线程(x,y,z)的ID值为(x+y∗Dx+z∗Dx∗Dy)。
/* get thread id: 1D block and 2D grid */
# define get_tid()(blockDim.x*(blockldx.x + blockldx.y*gridDim.x) + threadldx.x)
/* get block id: 2D grid */
# define get_bid()(blockldx.x + blockldx.y*gridDim.x)
__global__void vec_add(double *x, double *y, double *z, int n)
{
int i = get_tid();// user-defined function
if(i<n)z[i] = x[i] + y[i];
}
3 并行编程
3.1 线程并行
#include "cuda_runtime.h" //CUDA运行时API
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };
// Add vectors in parallel.
cudaError_t cudaStatus;
int num = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);
for(int i = 0;i<num;i++)
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "addWithCuda failed!");
return 1;
}
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadExit failed!");
return 1;
}
return 0;
}
// 重点理解这个函数
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0; //GPU设备端数据指针
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus; //状态指示
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0); //选择运行平台
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// 分配GPU设备端内存
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// 拷贝数据到GPU
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
// 运行核函数
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize(); //同步线程
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost); //拷贝结果回主机
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaFree(dev_c); //释放GPU设备端内存
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
<<<>>>表示运行时配置符号，里面1表示只分配一个线程组（又称线程块、Block），size表示每个线程组有size个线程（Thread）。本程序中size根据前面传递参数个数应该为5，所以运行的时候，核函数在5个GPU线程单元上分别运行了一次，总共运行了5次。
这5个线程是如何知道自己“身份”的？是靠threadIdx这个内置变量，它是个dim3类型变量，接受<<<>>>中第二个参数，它包含x,y,z 3维坐标，而我们传入的参数只有一维，所以只有x值是有效的。通过核函数中int i = threadIdx.x;这一句，每个线程可以获得自身的id号，从而找到自己的任务去执行。
3.2 块并行
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = blockIdx.x;
c[i] = a[i] + b[i];
}
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };
// Add vectors in parallel.
cudaError_t cudaStatus;
int num = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);
for(int i = 0;i<num;i++)
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "addWithCuda failed!");
return 1;
}
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadExit failed!");
return 1;
}
return 0;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
// Launch a kernel on the GPU with one thread for each element.
addKernel<<<size,1 >>>(dev_c, dev_a, dev_b);
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
<<<>>>里第一个参数改成了size，第二个改成了1，表示我们分配size个线程块，每个线程块仅包含1个线程，总共还是有5个线程。这5个线程相互独立，执行核函数得到相应的结果，与上一节不同的是，每个线程获取id的方式变为int i = blockIdx.x；这是线程块ID。
3.3 线程并行与块并行比较
线程并行是细粒度并行，调度效率高；块并行是粗粒度并行，每次调度都要重新分配资源，有时资源只有一份，那么所有线程块都只能排成一队，串行执行。
我们的任务有时可以采用分治法，将一个大问题分解为几个小规模问题，将这些小规模问题分别用一个线程块实现，线程块内可以采用细粒度的线程并行，而块之间为粗粒度并行，这样可以充分利用硬件资源，降低线程并行的计算复杂度。适当分解，降低规模，在一些矩阵乘法、向量内积计算应用中可以得到充分的展示。
3.4 流并行
CUDA中的流用cudaStream_t类型实现，用到的API有以下几个：cudaStreamCreate(cudaStream_t * s)用于创建流，cudaStreamDestroy(cudaStream_t s)用于销毁流，cudaStreamSynchronize()用于单个流同步，cudaDeviceSynchronize()用于整个设备上的所有流同步，cudaStreamQuery()用于查询一个流的任务是否已经完成。具体的含义可以查询API手册。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = blockIdx.x;
c[i] = a[i] + b[i];
}
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };
// Add vectors in parallel.
cudaError_t cudaStatus;
int num = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);
for(int i = 0;i<num;i++)
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "addWithCuda failed!");
return 1;
}
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadExit failed!");
return 1;
}
return 0;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStream_t stream[5];
for(int i = 0;i<5;i++)
{
cudaStreamCreate(&stream[i]); //创建流
}
// Launch a kernel on the GPU with one thread for each element.
for(int i = 0;i<5;i++)
{
addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i); //执行流
}
cudaDeviceSynchronize();
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
for(int i = 0;i<5;i++)
{
cudaStreamDestroy(stream[i]); //销毁流
}
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
<<<>>>中的参数多了两个，其中前两个和块并行、线程并行中的意义相同，仍然是线程块数（这里为1）、每个线程块中线程数（这里也是1）。第三个为0表示每个block用到的共享内存大小，这个我们后面再讲；第四个为流对象，表示当前核函数在哪个流上运行。我们创建了5个流，每个流上都装载了一个核函数，同时传递参数有些不同，也就是每个核函数作用的对象也不同。这样就实现了任务级别的并行，当我们有几个互不相关的任务时，可以写多个核函数，资源允许的情况下，我们将这些核函数装载到不同流上，然后执行，这样可以实现更粗粒度的并行。
3.5 性能剖析
在CUDA中，我们有专门测量设备运行时间的API，下面一一介绍。
cudaEvent_t start,stop;//事件对象
cudaEventCreate(&start);//创建事件
cudaEventCreate(&stop);//创建事件
cudaEventRecord(start,stream);//记录开始
myKernel<<<dimg,dimb,size_smem,stream>>>(parameter list);//执行核函数
cudaEventRecord(stop,stream);//记录结束事件
cudaEventSynchronize(stop);//事件同步，等待结束事件之前的设备操作均已完成
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start,stop);//计算两个事件之间时长（单位为ms）
cudaEventDestroy(start);//销毁事件
cudaEventDestroy(stop);//销毁事件
核函数执行时间将被保存在变量elapsedTime中。通过这个值我们可以评估算法的性能。下面给一个例子，来看怎么使用计时功能。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
__global__ void addKernel_blk(int *c, const int *a, const int *b)
{
int i = blockIdx.x;
c[i] = a[i]+ b[i];
}
__global__ void addKernel_thd(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i]+ b[i];
}
int main()
{
const int arraySize = 1024;
int a[arraySize] = {0};
int b[arraySize] = {0};
for(int i = 0;i<arraySize;i++)
{
a[i] = i;
b[i] = arraySize-i;
}
int c[arraySize] = {0};
// Add vectors in parallel.
cudaError_t cudaStatus;
int num = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);
for(int i = 0;i<num;i++)
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "addWithCuda failed!");
return 1;
}
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadExit failed!");
return 1;
}
for(int i = 0;i<arraySize;i++)
{
if(c[i] != (a[i]+b[i]))
{
printf("Error in %d\n",i);
}
}
return 0;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
for(int i = 0;i<1000;i++)
{
// addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);
addKernel_thd<<<1,size>>>(dev_c, dev_a, dev_b);
}
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float tm;
cudaEventElapsedTime(&tm,start,stop);
cudaEventDestroy(start);//销毁事件
cudaEventDestroy(stop);//销毁事件
printf("GPU Elapsed time:%.6f ms.\n",tm);
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize();
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess)
{
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
addKernel_blk是采用块并行实现的向量相加操作，而addKernel_thd是采用线程并行实现的向量相加操作。分别运行，得到的结果如下图所示：
线程并行：
块并行：
可见性能竟然相差近16倍！因此选择并行处理方法时，如果问题规模不是很大，那么采用线程并行是比较合适的，而大问题分多个线程块处理时，每个块内线程数不要太少，像本文中的只有1个线程，这是对硬件资源的极大浪费。一个理想的方案是，分N个线程块，每个线程块包含512个线程，将问题分解处理，效率往往比单一的线程并行处理或单一块并行处理高很多。这也是CUDA编程的精髓。
上面这种分析程序性能的方式比较粗糙，只知道大概运行时间长度，对于设备程序各部分代码执行时间没有一个深入的认识，这样我们就有个问题，如果对代码进行优化，那么优化哪一部分呢？是将线程数调节呢，还是改用共享内存？这个问题最好的解决方案就是利用Visual Profiler。下面内容摘自《CUDA_Profiler_Users_Guide》
“Visual Profiler是一个图形化的剖析工具，可以显示你的应用程序中CPU和GPU的活动情况，利用分析引擎帮助你寻找优化的机会。”
其实除了可视化的界面，NVIDIA提供了命令行方式的剖析命令：nvprof。对于初学者，使用图形化的方式比较容易上手，所以本节使用Visual Profiler。
Visual Profiler
打开Visual Profiler，可以从CUDA Toolkit安装菜单处找到。主界面如下：
我们点击File->New Session，弹出新建会话对话框，如下图所示：
具体使用方法
1万+
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
