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).
|
|