cuda和gpu的基本概念整理

1.基本介绍

SP（streaming Processor）是gpu最基本的硬件单元，俗称cuda核心（类似于cpu的一个核），多个sp组成一个SM（Streaming Multiprocessor），我的gtx 1060显卡是1920核心的，10个sm，每个sm有128个sp

grid、block、thread都是软件的概念，grid由很多个block组成，block又由多个thread组成，gpu执行时是把一个个block分配给SM执行，所以每个sm中的sp数量决定这个block的并行能力，在block里每32个thread组成一个wrap，执行相同的指令，只是数据不同，如果thread数理不是32的整数倍，将会有部分计算力浪费，因gpu是以wrap为单位来执行的，如果wrap里只有20个thread，那么会有8个sp在那里什么事都不干。如果SM里总的thread数量超过了sp的数量，那将会串行执行，因为sp的数理决定了最大并行能力。

一个SM 一次只会执行一个block 里的一个wrap，所以理论上真正在同时执行的thread数量是32sm数量，对gtx 1060来说是3210=320而不是1920，虽然在硬件上一个sm的多个wrap不是并行的，但是可能一个wrap存取global memory需要等待一段时间，那么就会切换到另一个wrap，这里的切换是没有损耗的。wrap是用来解决等待延迟的，我的一个block最多支持1024个线程，block里的线程容量越多，就能得到足够多的隐藏延迟操作

一个block里会通过share memory来通信，速度很快，share memory是存放在L2 Cache，一个线程里的栈是放在寄存器里的，如果寄存器放不下，会放到global memory里，这个稍微慢一点。

gpu独有的内存分配和拷贝cudaMalloc、cudaMemcpy，cpu和gpu可以共同访问的内存分配：cudaMallocManaged，见Programming Guide----》K. Unified Memory Programming这一节

第一个程序

#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
__global__ void hellofromGPU(void)
{
	printf("Hello world from GPU! by block %d thread %d\n",blockIdx.x,threadIdx.x);
}
int main()
{
	//hello from CPU
	printf("Hello world from CPU!\n");
	hellofromGPU <<<2, 5 >>> ();
	cudaDeviceReset();
	return 0;
}

2.测试结果

先说结论：不要特意去优化，照着官方给的demo直接去写就好了，性能差是规模不够造成的，当计算规模上去后各种优化的效果微乎其微，当然有些设置错误如单个block的线程超过1024这种明显的错误还是要改的

矩阵大小8192*8192
cpu E5-2683 v3双路 24核，编译-O3 -mavx -fopenmp 时间23s
gpu gtx 1050 2g：cuda 780ms cublas 800ms
gtx 1060 6g：cuda 780ms cublas 500ms
目前这块gtx 1060显卡质量有问题，结果可能不准确

测试代码：
https://github.com/pfysw/test/tree/master/mat_test

3.参考代码

首先是官方代码:
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.1

github上找了一些代码，没怎么测试，需要的自行测试:

https://github.com/NVIDIA/cutlass
https://github.com/Huanghongru/SGEMM-Implementation-and-Optimization
https://github.com/mz24cn/gemm_optimization
https://github.com/lzhengchun/matrix-cuda
https://github.com/kberkay/Cuda-Matrix-Multiplication

4.参考资料

其他的懒得写了，参考下面这些文章吧，因为我也是从下面的文章摘抄出来的。

官方手册
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.1\doc\html
CUDA中grid、block、thread、warp与SM、SP的关系
https://blog.csdn.net/wvh2007/article/details/49891363
cuda-Block和Grid设定
https://blog.csdn.net/yu132563/article/details/52548913
CUDA 核函数运行参数
https://blog.csdn.net/fuxiaoxiaoyue/article/details/83782167
GPU的并行运算与CUDA的简介
https://www.cnblogs.com/klausage/p/10323245.html
CUDA学习3-Grid&Block
https://blog.csdn.net/qq_41598072/article/details/81086500
CUDA 高性能并行计算入门
https://blog.csdn.net/cyhbrilliant/article/details/79434090
CUDA内存分配
https://blog.csdn.net/chengde6896383/article/details/81023653
CUDA中的常量内存__constant__
https://blog.csdn.net/dcrmg/article/details/54895834
CUDA: 使用shared memory
https://blog.csdn.net/a130737/article/details/44203617?utm_source=blogxgwz9
【CUDA】学习记录（7）- Global Memory
https://www.jianshu.com/p/3d4c9cc3a777
CUDA——内存层级
https://www.jianshu.com/p/7a8fe1aefd4e
CUDA程序性能分析-矩阵乘法
https://www.cnblogs.com/5long/p/cuda_matrix_performance.html
第四篇：使用 CUBLAS 库给矩阵运算提速
https://www.cnblogs.com/muchen/p/6306768.html
Cuda在gpu和主机之间统一内存(Cuda unified memory between gpu and host)
https://www.it1352.com/587846.html
CUDA入门（四）Visual Profiler nvprof
https://blog.csdn.net/qq_25819827/article/details/52107570
CUDA编程(三): GPU架构了解一下!
https://www.jianshu.com/p/87cf95b1faa0
CUDA之同步函数详解
https://blog.csdn.net/Bruce_0712/article/details/64928442