第二章
cuda代码写在.cu/.cuh里面
cuda 7.0 / 9.0开始,NVCC就支持c++11 / 14里面绝大部分的语言特性了.
Dim3
__host__ __device__ dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) {}
__host__ __device__ dim3(uint3 v): x(v.x), y(v.y), z(v.z) {}
Single Instruction, Multiple Data (SIMD)
用SIMD也不是一直是好的.
Streaming Multiprocessor (SM)一般每个有128个single precision CUDA cores(也就是一个线程)和对应的cache.
Block会被分成Warps, Warp是32个线程的集合(都在一个block里面).所有的32线程必须都跑同一组命令集.
一个SM里的Warps是同时跑的.
如果你想用一个Warp做不同的事儿,会按顺序做,也叫Warp Divergence.
Device Memory 也叫 Global Memory, 也就是GPU的RAM.从Device Memory拿比从真的RAM快.
Global Memory也没那么快,是很多GPU程序的bottleneck.
第三章
第四章
Latency: 硬件导致的延迟
Thoughput: 吞吐量
- CPU: 低延迟, 低吞吐量
- CPU clock: 3GHz
- main Memory latency: ~ 100+ns
- arithmetic instruction latency: ~1+ns
- GPU:高延迟,高吞吐量
- 1GHz
- 300+ns
- 10+ns
GPU非常的IOlimited,所以对与IO要谨慎处理.
- Registers:最快的,只有线程才能用,生命周期和线程一样.
- Local Memory: 150倍慢(比register和shared memory来说).
- shared memory:当没有bank conflicts或者从同一个地址读的时候,可以和register一样快. 对于一个block里面的所有线程都可见.和block一样的生命周期.
- global memory: 150倍慢(比register和shared memory)
global memory和GPU核不在一块.
- 有最大的容量
- GPU有.5到24GB的global memory,一般是~2GB.
- 延迟大概是~300ns 在kepler上.
shared memory
- 在SM上灰常快的memory
- 和L1 cache是一样的
- ~5ns的延迟
- 最大大约~48KB
shared memory语法
可以静态的分配shared memory,或者动态的分配
-
static allocation
__shared__ float data[1024]在kernel里面声明
-
dynamic allocation syntax
Host:
kernel<<<grid_dim, block_dim, numByteShMem>>>(args);
Device:
extern shared float s[];
还有些别的..
一个常见的pattern
Bank conflicts
任何序列的GPU变成都会对于导致表现下降.
Registers
大概比shared memory快10x.
每一个SM里大概有1万个registers.
一般kernel里声明的stack变量就是存储在registers.
Local Memory
是任何在stack上不能塞进register里的东西.
local memory只能给thread用.
L1 Cache
每一个SM有它自己的L1 Cache.
L2 Cache
被所有的SM共享
L3 Cache
比L2慢点但也大.
Constant Memory Constant Cache
In host code:
cudaMemcpyToSymbol(foo, h_src, sizeof(int)*1024);
Texture Memory
这个东西非常复杂,而且对于一般的计算来说只是有点用.
有用的特性:
- (没看懂) 2D or 3D data locality for caching purposes through "CUDA arrays". Goes into special texture cache.
- 在一维/二维/三维的array插值快.
- 把int型转化为统一的("unitized")浮点数.
常用的场景:
- 用texture cache和cuda array来读输入数据,来利用空间缓存(spatial caching).
- 利用numerical texture capabilities.
- 和OpenGL以及一般的computer graphics交互.