配色: 字号:
GPU与CUDA简介
2012-08-08 | 阅:  转:  |  分享 
  
GPU与CUDA简介

上海超级计算中心研发部徐磊

Outline

?GPU简介

?GPU(Tesla)系统架构

?CUDAArchitecture

?CUDAProgrammingModel

?基于NVIDIAGPU的应用结果展示

?CUDA实践

一、GPU简介

GraphicProcessingUnit(GPU)

?用于个人计算机、工作站和游戏机的专用图像显示设

?NVIDIA和AMD(ATI)是主要制造商

GPU与CPU硬件架构的对比

Cache

ALU

Control

ALU

ALU

ALU

DRAM

CPU

DRAM

GPU

?CPU:面向通用计算;大量的晶体管用于Cache和控制电路

?GPU:面向计算密集型和大量数据并行化的计算;大量的晶体

管用于计算单元

GPU的强大运算能力

?浮点运算性能比较

?存储器带宽比较

CUDA成功案例

广泛应用于生命科学、机械、石油、金融、数学、天文和通信等行业

二、GPU(Tesla)系统架构

Processor1xTeslaT10

Numberofcores240

CoreClock1.296GHz

On-boardmemory4.0GB

Memorybandwidth102GB/secpeak

MemoryI/O512-bit,800MHzGDDR3

Single/Double

Precisionfloating

pointperformance

(peak)

933/78GFlops

SystemI/OPCIex16Gen2

Typicalpower160W

TeslaC1060ComputingProcessor

GPU(Tesla)系统架构

系统组成:

?10TPC(ThreadProcessorCluster)

?3SM(StreamMultiprocessor)/TPC

?8SP(StreamProcessor)/SM

TeslaStreamingMultiprocessor(SM)

?SM包含8个SP

–支持IEEE75432-bitfloatingpoint

–支持32-bitand64-bitinteger

–包含16K32-bitregisters

?SM包含2个SFU单元(SpecialFunctionUnits)

?SM包含1个DP(DoublePrecisionUnit)

–支持IEEE75464-bitfloatingpoint

?16KB共享内存

–Concurrentthreadssharedata

–Lowlatencyload/store

三、CUDAARCHITECTURE

?GPU上的通用编程模型

-单指令、多数据执行模式(SIMD)

?所有线程执行同一段代码

?大量并行计算资源处理不同数据

CUDA--COMPUTEUNIFIEDDEVICEARCHITECTURE

...

...

CPUSerialCode

CPUSerialCode

GPUParallelKernel

GPUParallelKernel

CforCUDAandOpenCL

编译器生成的中间代码

OpenCL

CforCUDA

PTX

GPU

适合熟悉C

的开发人员

使用

倾向于使用底层

API的开发人员

使用

CUDA软件体系

应用程序可以使用CUDALibraries、CUDARuntimeAPI和

CUDADriverAPI从3个层次使用GPU资源

四、CUDAPROGRAMMINGMODEL

CUDA设备与线程

?计算设备(device)

–作为CPU(host)的协处理器

–有独立的存储设备(devicememory)

–同时启动大量线程

?计算密集部分使用大量线程并行的kernel,在GPU上

执行

?GPU与CPU线程的区别

–GPU的线程非常轻量,线程切换~1cycle,而CPU需要~1000cycle

–GPU上的线程数>1000时才能有效利用GPU的计算能力

Device

Grid1

Block

(0,0)

Block

(1,0)

Block

(2,0)

Block

(0,1)

Block

(2,1)

Block

(1,1)

Host

serialcode

serialcode

Kernel2

Kernel1

Grid2

Block(1,1)

Thread

(0,0)

Thread

(1,0)

Thread

(2,0)

Thread

(3,0)

Thread

(4,0)

Thread

(0,1)

Thread

(1,1)

Thread

(2,1)

Thread

(3,1)

Thread

(4,1)

Thread

(0,2)

Thread

(1,2)

Thread

(2,2)

Thread

(3,2)

Thread

(4,2)

CUDA执行模型

?Thread:并行的基本单位

?Threadblock:互相合作的线程



–CooperativeThreadArray(CTA)

–允许彼此同步

–通过快速共享内存交换数据

–以1维、2维或3维组织

–最多包含512个线程

?Grid:一组threadblock

–以1维或2维组织

–共享全局内存

?Kernel:在GPU上执行的核心程序

–Onekernel?onegrid

软硬件映射

Block被分配到一个SM中

CUDA存储器模型

?R/Wper-threadregisters

–1-cyclelatency

?R/Wper-threadlocalmemory

–Slow–如果register存不下将被放入global

memory

?R/Wper-blocksharedmemory

–1-cyclelatency

–使用__shared__定义

?R/Wper-gridglobalmemory

–~500-cyclelatency

–使用__device__定义

?Readonlyper-gridconstantandtexture

memories

–~500-cyclelatency,但是拥有缓存

GPUGlobalMemory分配

?cudaMalloc()

z分配显存中的globalmemory

z两个参数

?对象数组指针和数组尺寸

?cudaFree()

z释放显存中的globalmemory

?对象数组指针

intblk_sz=64;

floatMd;

intsize=blk_szblk_szsizeof(float);

cudaMalloc((void)&Md,size);



cudaFree(Md);

Host–Device数据交换

?cudaMemcpy()

–Memorydatatransfer

–Requiresfourparameters

?Pointertodestination

?Pointertosource

?Numberofbytescopied

?Typeoftransfer

–HosttoHost,HosttoDevice,

DevicetoHost,DevicetoDevice

cudaMemcpy(Md,M.elements,size,

cudaMemcpyHostToDevice);

cudaMemcpy(M.elements,Md,size,

cudaMemcpyDeviceToHost);

CUDA函数定义

?__global__定义kernel函数

–必须返回void

?__device__函数

–不能用&运算符取地址,不支持递归调用,不支持静态变量(static

variable),不支持可变长度参数函数调用

Executedonthe:Onlycallablefromthe:

__device__floatDeviceFunc()devicedevice

__global__voidKernelFunc()devicehost

__host__floatHostFunc()hosthost

Kernel函数调用

?调用时必须给出线程配置方式

__global__voidKernelFunc(...);

dim3DimGrid(100,50);//5000threadblocks

dim3DimBlock(4,8,8);//256threadsperblock

size_tSharedMemBytes=64;//64bytesofsharedmemoryperblock

KernelFunc<<>>(...);

CUDAC示例:向量相加VecAdd

//main.c文件,计算向量A+B,结果

//保存于向量C中

#definesize100

intmain(intargc,charargv[])

{

intbytes=sizesizeof(float)

floatA=(float)malloc(bytes);//Host端分配内存

floatB=(float)malloc(bytes);

floatC=(float)malloc(bytes);

inti=0;

floatd_A,d_B,d_C;

cudaMalloc((void)&d_A,bytes);//device端分配内存

cudaMalloc((void)&d_B,bytes);

cudaMalloc((void)&d_C,bytes);

for(i=0;i
{

A[i]=i/3.0;

B[i]=i/7.0;

}

cudaMemcpy(d_A,A,bytes,cudaMemcpyHostToDevice);//内存拷贝

cudaMemcpy(d_B,B,bytes,cudaMemcpyHostToDevice);

VecAdd<<<1,100>>>(d_A,d_B,d_C);//lanuchkernel

cudaMemcpy(C,d_C,bytes,cudaMemcpyDeviceToHost);//获取结果

return0;

}

//kernel.cu文件

__global__voidVecAdd(floatd_A,floatd_B,floatd_C)

{

inttid=threadIdx.x;//内建变量,获取线程ID

d_C[tid]=d_A[tid]+d_B[tid];

}

CUDAC程序的编译

CUDACC

CPUCompiler

CforCUDA

Kernels

CUDAobject

files

RestofC

Application

CPUobject

files

CPU-GPU

Executable

NVCC

CforCUDA

Application

Linker

CombinedCPU-GPUCode

CUDA程序的编译

?使用nvcc编译工具

nvcc.cu[-oexcutable]

?调试选项:-g(debug)、-deviceemu(CPU模拟GPU)

五、基于NVIDIAGPU的应用结果展示

GPU上应用案例

?Amber11(著名的分子动力学软件,用于蛋白质、核酸、糖等生物

大分子的计算模拟)

–PMEMD(“ParticleMeshEwaldMolecularDynamics,”):sander的GPU版本,

针对NvidiaCUDA进行优化

?Gromacs(生物分子体系的分子动力学程序包):

–Gromacs4.5-beta1支持GPU

–GPU加速采用SimBios的openMM2.0库

?NAMD(用于在大规模并行计算机上快速模拟大分子体系的并行分

子动力学代码)

–NAMD2.7b2支持GPU

GPU上应用案例

?AutoDock:

–计算机辅助药物设计软件,2010/4开源

–运行方式:

AUTODOCKHOME/autodock4–p

?天体多体问题astrophysics

–Gadget-2

–颗粒部分移植到GPU上,气体模拟部分SPH没有GPU版本

?LAMMPS(大规模原子分子并行模拟器)

–目前支持L-J和Gay-Berne势能函数

–和Gadget-2类似,对颗粒间作用力计算在GPU上加速

AMBER

?测试平台

–GPU:TeslaC1060;峰值性能:933GFlops;显存:4GB

–CPU:AMDBarcelonaOpteron1.9GHz,16cpucores64GB;intel11.1,

mpichIBversion

0

2

4

6

8

10

12

DHFRNVE=

23,558atoms

DHFRNPT=

23,558atoms

FactorIXNVE=

90,906atoms

CelluloseNVE

=408,609

atoms

ns/day

ExplicitSolventPMEBencmark

CPU

GPU

$AMBERHOME/bin/pmemd.cuda-O-imdin-omdout-pprmtop-cinpcrd-rrestrt-xmdcrd

AMBER

0

5

10

15

20

25

30

35

Myoglobin=2492atoms

ImplicitSolventGB

CPU

GPU

0

0.2

0.4

0.6

0.8

1

1.2

Nucleosome=25095atoms

ImplicitSolventGB

CPU

GPU

?ExplicitSolventGBbenchmark:加速比1.2~2.25

?ImplicitSolventGBbenchmark:加速比6.7,25

Gromacs

?测试平台:

–GPU:TeslaC1060;峰值性能:933GFlops;显存:4GB

–CPU:2xIntelquadcoreE5462

?以下3个案例来自于nVIdia官方网站:加速比3.5,5.2,22倍

NAMD

?测试平台

–GPU:nVidia1060;峰值性能:933GFlops;显存:4GB

–CPU:AMDBarcelonaOpteron1.9GHz,1cpucore;64GB;intel11.1,

mpichIBversion

?运行方式:

–$NAMDHOME/namd2+idlepoll

?加速比:

–CPU:17.2551/step

–GPU:3.19981/step

–加速比:5.39

Gadget-2的GPU版本g2x

?测试平台

–GPU:nVidia1060;峰值性能:933GFlops;显存:4GB

–CPU:AMDAthlon644600+X2

–Np=30

3

–Np=32

3

–Np=64

3

总计算

时间(sec)

force_treeevaluate

_shortrange(sec)

总计算

时间(sec)

force_treeevaluate

_shortrange(sec)

总计算时

间(sec)

force_treeevaluate

_shortrange(sec)

CPU71.470.999.398.7937.6934.8

GPU18.117.522.722.132.828.3

加速比

3.944.054.374.4729.6833.03

LAMMPS

?SandiaNL开发的分子动力学并行计算软件

–有很好的并行可扩展性(10000核)

?2009SupercomputingConference宣布向GPU平台移植

–支持单节点多GPU卡并行运算

六、CUDA实践

----MonteCarlo模拟用于欧式期权定价

欧式期权定价

欧式看涨期权的支付函数为

}0,max{),(KSKSf

Tt

?=

---K为到期执行价格

---T为到期时间

---S

t

为标的资产在t时刻的价格

在有效市场,金融资产的价格服从随机游走模型,标的资产

价格变化所遵循的过程可以写作:

tttt

dWSdtSdSσμ+=

(0,1)

tt

dWdtdWdt

N

μ

σ

ε

ε

??

??

??=

飘移率

波动率

与相关的维纳过程,

为服从随机变量

?

?

?

?

?

?

Δ+Δ

?

?

?

?

?

?

?

?

?=Δ+ttrtSttSσε

σ

2

exp)()(

2

–通过推导可以得到资产的价格为

t

r

Δ??

??

时间步长度

无风险利率

–最终得到欧式看涨期权的价格为

)],([KSfEe

T

rT



?

–蒙特卡罗模拟可以通过独立随机采样来逼近式这个价格

内层循环

交由GPU的thread进行计算

外层循环

使用n个

Threads

进行计算

GPU上具体实现

?数据准备

doubled_payoffSum;

intbytes=NPATH_PER_ITERsizeof(double);

cutilSafeCall(cudaMalloc((void)&d_payoffSum,bytes));

cutilSafeCall(cudaMemset(d_payoffSum,0.0,bytes));

doubleh_payoffSum;

h_payoffSum=(double)malloc(bytes);

intnum_bytes=sizeof(double)NNUM_PER_ITER;

doubledev_rand_num;

cutilSafeCall(cudaMalloc((void)&dev_rand_num,num_bytes));

curandStatestate;

cutilSafeCall(cudaMalloc((void)&state,sizeof(curandState)NNUM_PER_ITER));

?使用GPU计算

init_kernel<<>>(state);

for(inti=0;i
{

gen_kernel<<>>(dev_rand_num,state);

mc_kernel<<>>(dev_rand_num,d_payoffSum);

}

?传回结果

cutilSafeCall(cudaMemcpy(h_payoffSum,

d_payoffSum,

NPATH_PER_ITERsizeof(double),

cudaMemcpyDeviceToHost)

);

计算结果比较

0

200

400

600

800

1000

1200

1400

CPU串行16核CPU并行使用单块GPU计算

时间(

秒)

计算时间(PATH=1.7810

7



14.69X

43.76X

–测试环境

–CPUAMD8347HE1.9GHz

–GPUTeslaS2050

–Nvcc、gcc编译器

–Mvapich通讯库

–魔方单节点测试MPI并行

cuRand库的使用

?CPU端的随机数生成器函数不能在GPU端调用

?NVIDIA于2010年8月提供cuRand库,使用GPU生成随机数

?cuRand生成随机数示例

curandStatestate;

curand_init(1234,tid,0,&state);

doublerand_num=curand_normal_double(&state);

性能提示

?最好使用不同的kernel生成state和随机数。

?调用curand_init非常耗时,建议state生成后直接保存于

globalmemory中,供生成新的随机数来使用。

?如果需要生成的随机数为N,需要测试确定需要initialize的state

的数目

随机数生成示例

__global__voidinit_kernel(curandStatestate)

{

inttid=threadIdx.x+blockDim.xblockIdx.x;

curand_init(1234,tid,0,&state[tid]);

}

__global__voidgen_kernel(doubledev_rand_num,curandStatestate)

{

inttid=threadIdx.x+blockDim.xblockIdx.x;

curandStatelocal_state=state[tid];

dev_rand_num[tid]=curand_normal_double(&local_state);

state[tid]=local_state;

}

补充材料

激动人心的CUDA4.0(2011/2/28)

TheNVIDIACUDA4.0Toolkitwasdesignedtomakeparallel

programmingeasier,andenablemoredeveloperstoporttheir

applicationstoGPUs.Thishasresultedinthreemainfeatures:

–NVIDIAGPUDirect?2.0Technology

–UnifiedVirtualAddressing(UVA)

–ThrustC++TemplatePerformancePrimitivesLibraries









TheCUDA4.0architecturereleaseincludesanumberofotherkeyfeaturesand

capabilities,including:

–MPIIntegrationwithCUDAApplications--ModifiedMPIimplementations

automaticallymovedatafromandtotheGPUmemoryoverInfinibandwhenan

applicationdoesanMPIsendorreceivecall.(GPUDirect?1.0Technology)

–Multi-threadSharingofGPUs--MultipleCPUhostthreadscansharecontexts

onasingleGPU,makingiteasiertoshareasingleGPUbymulti-threaded

applications.

–Multi-GPUSharingbySingleCPUThread--AsingleCPUhostthreadcanaccess

allGPUsinasystem.Developerscaneasilycoordinateworkacrossmultiple

GPUsfortaskssuchas"halo"exchangeinapplications.

–NewNPPImageandComputerVisionLibrary--Arichsetofimage

transformationoperationsthatenablerapiddevelopmentofimagingandcomputer

visionapplications.

–NewandImprovedCapabilities

–AutoperformanceanalysisintheVisualProfiler

–Newfeaturesincuda-gdbandaddedsupportforMacOS

–AddedsupportforC++featureslikenew/deleteandvirtualfunctions

–NewGPUbinarydisassembler

关于CUDALinux开发环境

?Eclipse+CDT+Fixstarscudaplugin+cuda-gdb

谢谢!

衍生物及期权定价

?利用随机微分函数描述衍生物价格模型

–漂移率

–波动率

–白噪声

?假设漂移和波动率与执行价格线性相关,有以下解析解

?用公式()对期权价格进行模拟,取平均值,可以得到收益。

()()

tttt

dWtSdttSadS,,σ+=

()

tt

StSσσ=,

()

tt

aStSa=,

t

dW

()()

()

()

2

2

1

2

2

1

1

0

t

t

dWdta

tt

Wta

t

eSS

eSS

σσ

σσ

+?

?

+?

=

=

激动人心的CUDA4.0(2011/2/28)

TheNVIDIACUDA4.0Toolkitwasdesignedtomakeparallel

programmingeasier,andenablemoredeveloperstoporttheir

applicationstoGPUs.Thishasresultedinthreemainfeatures:

–NVIDIAGPUDirect?2.0Technology--Offerssupportforpeer-

to-peercommunicationamongGPUswithinasingleserveror

workstation.Thisenableseasierandfastermulti-GPU

programmingandapplicationperformance.

–UnifiedVirtualAddressing(UVA)--Providesasinglemerged-

memoryaddressspaceforthemainsystemmemoryandtheGPU

memories,enablingquickerandeasierparallelprogramming.

–ThrustC++TemplatePerformancePrimitivesLibraries--

ProvidesacollectionofpowerfulopensourceC++parallel

algorithmsanddatastructuresthateaseprogrammingforC++

developers.WithThrust,routinessuchasparallelsortingare

5Xto100XfasterthanwithStandardTemplateLibrary(STL)and

ThreadingBuildingBlocks(TBB).

献花(0)
+1
(本文系yangshiquan...首藏)