CUDA课程-CS-139G-1:Introduction to Massively Parallel Computing

size_t是为了软件的兼容性和可移植性而产生的,使应用程序不依赖于具体的系统实现

size_t是针对系统定制的一种数据类型，一般是整形，因为C/C＋＋标准只定义一最低的位数，  
  而不是必需的固定位数。而且在内存里，对数的高位对齐存储还是低位对齐存储各系统都不一样。为了提高代码的可移植性，就有必要定议这样的数据类型。一般这 种类型都会定义到它具体占几位内存等。当然，有些是编译器或系统已经给定义好的。具体要查看技手册。

 

cudaMalloc( void ** devPtr, size_t size) float data_f; size_t size; cudaMalloc((void**) &data_f, size);

 

===CS-139G-1:Introduction to Massively Parallel Computing===
一，课程目标：大规模平行运算的目的是要达到：
1，高性能；2，函数性和可维护性；3，与新生代的兼容性
为此需要掌握：
1，并行编程的原则和模式；
2，处理器的架构特征和局限性；
3，API编程，工具以及技巧。
二，教学人员
三，Web资源
http://stanford-cs193g-sp2010.googlecode.com
四，评分标准
五，个人小灶日
六，学术诚信
七，课程设备
八，文字和注释
九，时间安排（10 weeks）
十，摩尔定律（阐释）
十一，摩尔定律（直观图）
十二，电源功耗的购买表现力
十三，连续功效按比例已经结束
1，处理器频率按比例发展--不超过10G
2, 能量功耗不能继续上升--烧芯片
3，可以继续增加晶体管集成度--摩尔定律
十四，怎样使用晶体管
1，指令的平行（无序执行，臆测；功耗限制使得机会的丢失）
2，数据的平行
（向量单元，SIMD 执行）SIMD-single instruction, mutiple data
 增加:SSE,AVX,Cell SPE,Clearspeed,GPU
 3,线程的平行
 增加多线程，多核，超多核
 intel 双核，Amd,Sun,STL Nvidia Fermi;
 十五，为什么需要并行处理？
 默然的变革和潜在上升的计算，带宽
 十六，新的摩尔定律
 1，计算机不是变快二十变宽了
 2，必须要重新考虑用平行算法
 3，数据平行的计算是最具扩展性的解决方案
  （否则还要求更多的核心，而且数据总比核心多，计算总围绕在数据周围）
  十七，传统多核心
  1，多个核心支持同一个线程
  2，片上存储空间在处理器周围（缓存，RAM,或二者）
  3，共享存储空间（外部DRAM）
 十八，传统超多核心
 1，多个核心支持多个硬件线程
 2，片上存储空间在处理器周围（缓存，RAM,或二者）
 3，共享存储空间（外部DRAM）
 十九，走进GPU
 1, 规模效益
 2，规模并行
 二十，GPU的革命
 1，高吞吐量的计算
 2, 高记忆体带宽
 3，高可用性
 二十一，图形管线课程
 1，吞吐量是最重要的（在架构时间内完成每一个像素；可伸缩性）
 2，非常迅速的创建，运行，释放多个线程
 3，多线程来隐藏延迟
 二十二，为什么和CPU不一样？
 1，不同的目的引发不同的设计
 2，CPU将延迟熟练最小化到1个线程（大容量的缓存和精通控制逻辑）
 3，CPU将所有线程吞吐量最大化
 （运行的线程受资源限制；多线程掩盖了延迟；共享控制逻辑来存取多个线程）
 
 二十三，NVIDIA GPU架构
 Fermi GF100;
 二十四，多核处理器的SM处理单元（Streaming Multiprocessor）
 1，每个SM有32个cuda处理核心，一共512个
 2，8倍64位浮点峰值能力（50%的32位浮点数峰值）
 3，直接加载和存储能力（按位的线性序列和高带宽）
 4，64k的高速片上RAM(由软硬件管理，多个CUDA核心共享，可线程间通信)
 二十四，关键架构观点
 1，SIMT(single instruction, mutiple data)执行
   32个线程一组，成为warps
   一个warps共享指令单元，（IU）
   硬件自动处理分歧
 2，硬件多线程
   硬件资源分配和线程期限
   硬件依赖线程来隐藏延迟
 3，线程拥有所有需要运行的资源
   warp不等待就可以运行
   上下文切换很自由

=============================================

 走进CUDA
 
 1，可伸缩的并行编程模型
 2，为熟悉C/C++环境而进行最小化的扩展
 3，非均匀的串并行计算

 

 

 

CUDA:可伸缩的并行编程

1，在C/C++中增加最小化的简单抽象

   使得程序员可以专注于并行算法；而不用另外学习一门并行编程语言

2，提供和硬件平台的直接映射

   很合适GPU架构；同样和多核CPU也能良好的映射

3，数以百计地核与数以千计的并行线程来衡量

  GPU本身的线程是轻量级的--创建和切换都很自由

  GPU需要数以千计的线程进行全力运用

 

CUDA中并行抽象的关键点：

==============

1，并发线程的等级

  1，有许多线程构成的并行化kernel:所有进程都在执行同一连续的程序

  2，线程按组构成线程块：同一个线程块内地线程可以协作

  3，线程和线程块都具有独一无二的标识号

==================

2，轻量级的同步基元

3，合作线程的共享存储模型

 

CUDA不像一些平面式的多重处理器：

1，全局同步起来消耗资源大

2，全局的存储空间存取起来代价高

3，PRAM(并行随机存储机器模型)

CUDA不像一些分布式的处理器：

1，分布式的计算是不同的设定

2，大批量同步并行模型，MPI

 

非均匀的计算:

多核CPU<->超多核GPU

 

==============================

 

CUDA的C

 

1，原则就是：为了达到扩大性能而需进行的最小化扩展的设定

2，函数的标识符

__global__ void my_kernel() { }
__device__ float my_device_func() { }

3，变量的标识符（内置）

__constant__ float my_constant_array[32];
__shared__ float my_shared_array[32];

4，执行时代配置

dim3 grid_dim(100, 50); // 5000 thread blocks
dim3 block_dim(4, 8, 8); // 256 threads per block
my_kernel <<< grid_dim, block_dim >>> (...); // Launch kernel

5，在GPU端内建的有效变量和功能函数：

dim3 gridDim; // Grid dimension
dim3 blockDim; // Block dimension
dim3 blockIdx; // Block index
dim3 threadIdx; // Thread index
void __syncthreads(); // Thread synchronization

 

========

例子：向量相加：

 

#include <stdio.h> #include <stdlib.h> #include <time.h> #include <cuda_runtime.h> #include <cutil.h> #include <iostream> using namespace std; #define VEC_SIZE 512 //kernel function;run on GPU; //"__global__" is mark for code that run on GPU __global__ void vecAdd(float* d_A,float* d_B,float* d_C) { int index=threadIdx.x; d_C[index]=d_A[index]+d_B[index]; } int main() { //got the size of the data in DDR;float data size is 4 bit size_t size=VEC_SIZE*sizeof(float); //为本地分配内存;在DDR分配；返回指向该空间的一个指针。 float* h_A=(float*)malloc(size); float* h_B=(float*)malloc(size); float* h_C=(float*)malloc(size); //初始化 for (int i=0;i<VEC_SIZE;++i) { h_A[i]= rand() %10; h_B[i]= rand() %10; } int deviceCount; //get the GPU's compute capability cudaGetDeviceCount(&deviceCount); int device; //the CUDA Runtime data type, is a struct. cudaDeviceProp* getcudainfo; printf("%d/n", deviceCount); for (device = 0; device < deviceCount; ++device) { //get the device date type(a struct) cudaDeviceProp deviceProp; //get the property content,return to struct cudaGetDeviceProperties(&deviceProp, device); if (device == 0) { if (deviceProp.major == 9999 && deviceProp.minor == 9999) printf("There is no device supporting CUDA./n"); else if (deviceCount == 1) { getcudainfo=&deviceProp; printf("There is 1 device supporting CUDA/n"); printf ("%d/n/n",getcudainfo->sharedMemPerBlock ); //get the max threads number per block } else printf("There are %d devices supporting CUDA/n", deviceCount); } } float* devPtr; cudaMalloc((void**)&devPtr, 256 * sizeof(float)); //4*256=1024 bit //将本地内存的中的数据复制到设备中 float* d_A; cudaMalloc((void**)&d_A,size); cudaMemcpy(d_A,h_A,size,cudaMemcpyHostToDevice); float* d_B; cudaMalloc((void**)&d_B,size); cudaMemcpy(d_B,h_B,size,cudaMemcpyHostToDevice); //分配存放结果的空间 float* d_C; cudaMalloc((void**)&d_C,size); //定义512个线程 dim3 dimblock(512); vecAdd<<<1,dimblock>>>(d_A,d_B,d_C); //讲计算结果复制回主存中 cudaMemcpy(h_C,d_C,size,cudaMemcpyDeviceToHost); //输出计算结果 for (int j=0;j<VEC_SIZE;++j) { printf("%f/t",h_C[j]); } //释放主机和设备内存 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); free(h_A); free(h_B); free(h_C); return 0; }

 

应用程序的加速

 

1，用GeForce 8800 GTX和2.2GHz的皓龙处理器248比较

2，只要kernel能够占据足够的并行线程，正常kernel可达到10倍加速

3，如果功能函数的数据需求和控制流和GPU与优化过的应用程序相搭配，可达到25到400倍的加速。

 

最终观点：

 

1，并行化的硬件将为人通用；

2，GPU系列产品及是大规模并行超多核新处理器：容易获得和全面可编程

3，并行化和可伸缩性是成功关键点

4，给许多重要研究挑战提供了便利，同样也给教育意义的挑战性提供了便利。