CUDA之大规模GPU并行系列教程(二)——大规模并行计算GPU与CUDA简介
“The number of transistors on an integrated circuit doubles every two years.” – Gordon E. Moore著名的摩尔定律。
下面是摩尔定律的可视化分析图:
串行性能的扩展已经结束:
- 不能继续再扩展处理器 (没有10GHz的芯片)
- 不能继续再增加能源的消耗(不能把芯片融化)
- 能够继续增加晶体管的密度(摩尔定律)
怎么使用晶体管???
- Instruction-level parallelism
- 乱序执行(out-of-order execution),推测(speculation)
- 在能量有限的世界,有消失的机会
- Data-level parallelism
- 向量单元(vector unit),单指令多数据(SIMD execution)
- 应用增加(SSE,AVX,Cell SPE,Clearspeed,GPU)
- Thread-level parallelism
- 应用增加(多线程(multithreading),multicore(多核cpu),manycore(GPU微处理器))
GPU与CPU性能对比:
新的摩尔定律:
- 计算机不再变得更快,只是容量更大
- 你必须重构并行化的算法
- 数据并行计算是最有扩展能力的解决方案
而不是仅为双核,4核,8核,16核重构代码,你总会有比处理器数更多的计算数据。
通常的多核芯片
通常的mangcore芯片
GPU
GPU的发展
GPU从图形流水线上得到了宝贵的经验:
- 吞吐量是最重要的
必须在单位帧时间内绘制每个像素
- 快速的创建,运行,休眠大量的线程
在increament()kernel函数上测量14.8 Gthread/s
- 使用多线程来隐藏延迟
如果有一百个线程在预备状态(ready)等待调度运行,有一个线程停止或延缓也没有关系。
GPU与CPU的区别
- 不同的目的生成不同的设计
GPU工作量是高度的并行
CPU必须兼顾并行和串行
- CPU:每个线程最小化的延迟
大量片上缓存(on-chip caches)
复杂的逻辑控制单元(control unit)
- GPU: 所有线程最大化的吞吐量
线程的运行被资源限制---》提供大量的资源(寄存器(registers),带宽(bandwidth),共享内存(share memory),constant memory)
多线程能够隐藏延迟---》忽略缓存
共享逻辑控制能够访问多线程
SM: stream multiprocessor
CUDA
- 可扩展的并行编程模型
- C/C++家族最小化的扩展集合
- 异构串行-并行计算(heteogeneous serial-parallel computing)
cuda 作用
cuda:可扩展的并行编程
- 对C/C++增加最小程度的抽象
使程序员集中精力在并行算法上,而不是在并行语言的结构机制上
- 提供硬件的直接映射
适合GPU体系结构
也很好的映射到多核CPU
- GPU线程是轻量级的
- GPU需要大量的线程才能充分运用
扩展语法:
1 //Philosophy: provide minimal set of extensions necessary to //expose power 2 // Function qualifiers: 3 __global__ void my_kernel() { } 4 __device__ float my_device_func() { } 5 //Variable qualifiers: 6 __constant__ float my_constant_array[32]; 7 __shared__ float my_shared_array[32]; 8 //Execution configuration: 9 dim3 grid_dim(100, 50); // 5000 thread blocks 10 dim3 block_dim(4, 8, 8); // 256 threads per block 11 my_kernel <<< grid_dim, block_dim >>> (...); // Launch kernel 12 // Built-in variables and functions valid in device code: 13 dim3 gridDim; // Grid dimension 14 dim3 blockDim; // Block dimension 15 dim3 blockIdx; // Block index 16 dim3 threadIdx; // Thread index 17 void __syncthreads(); // Thread synchronization
例如:向量加法
1 //NVIDIA Corporation 2 //Example: vector_addition 3 // compute vector sum c = a + b 4 // each thread performs one pair-wise addition 5 6 __global__ void vector_add(float* A, float* B, float* C) 7 { 8 int i = threadIdx.x + blockDim.x * blockIdx.x; 9 C[i] = A[i] + B[i]; 10 } 11 12 int main() 13 { 14 // allocate and initialize host (CPU) memory 15 float *h_A = …, *h_B = …; 16 // allocate device (GPU) memory 17 float *d_A, *d_B, *d_C; 18 cudaMalloc( (void**) &d_A, N * sizeof(float)); 19 cudaMalloc( (void**) &d_B, N * sizeof(float)); 20 cudaMalloc( (void**) &d_C, N * sizeof(float)); 21 // copy host memory to device 22 cudaMemcpy( d_A, h_A, N * sizeof(float), 23 cudaMemcpyHostToDevice) ); 24 cudaMemcpy( d_B, h_B, N * sizeof(float), 25 cudaMemcpyHostToDevice) ); 26 // launch N/256 blocks of 256 threads each 27 vector_add<<<N/256, 256>>>(d_A, d_B, d_C); 28 29 }
在UIUC提供了一些工程项目
so, 这周就这样,各位加油吧!!!