超大规模并行程序设计.ppt

Second Generation Parallel Thread Execution ISA 64-bit memory address space - with ECC Unified address space with full C++ support Optimized execution of OpenCL and DirectCompute Full IEEE 754-2008 single and double precision floating point numbers NVIDIA GigaThreadTM Engine Multi-kernel execution 10x faster application context switching Concurrent kernel execution Out-of order thread block execution Two streaming transfer engines Work in a pipelined, overlapped manner Each could saturate the PCIe interface Performance 参考书目 NVIDIA, CUDA Programming Guide, NVidia, 2008, 2009. 张舒、褚艳利、 赵开勇、张钰勃, GPU高性能运算之CUDA, 2009. /. T. Mattson, et al., Patterns for Parallel Programming, Addison Wesley, 2005. 参考文献 Special issues on GPU computing IEEE, Proceedings of IEEE, Vol. 96, No. 5, May, 2008. ACM, Queue, Vol. 6, No. 2, March/April, 2008. Elsevier, Journal of Parallel and Distributed Computing, Vol. 68, No. 10, October, 2008. * CUDA函数定义 __global__ 定义kernel函数 必须返回void __device__ 函数 不能用运算符取地址, 不支持递归调用, 不支持静态变量(static variable), 不支持可变长度参数函数调用 Executed on the: Only callable from the: __device__ float DeviceFunc() device device __global__ void KernelFunc() device host __host__ float HostFunc() host host CUDA数学函数 pow, sqrt, cbrt, hypot, exp, exp2, expm1, log, log2, log10, log1p, sin, cos, tan, asin, acos, atan, atan2, sinh, cosh, tanh, asinh, acosh, atanh, ceil, floor, trunc, round, etc. 只支持标量运算 许多函数有一个快速、较不精确的对应版本 以”__”为前缀，如__sin() 编译开关-use_fast_math强制生成该版本的目标码 实例2: 矩阵相乘 矩阵数据类型 – 不属于CUDA! 单精度浮点数 width ? height个元素 矩阵元素在elements中 1-D数组存放矩阵数据 Row-major storage typedef struct { int width; int height; float* elements; } Matrix; A B C WM.width = N.heightI M.height M.width N.width 实例2: 矩阵相乘 C = A ? B of size WIDTH x WIDTH 一个线程处理一个矩阵元素 简化: 假定 WIDTH x WIDTH 512 只需要一个thread block 线程载入A的一行和B的一列 A和B的一对相应元素作一次乘法和一次加法 A B C WIDTH WIDTH WIDTH WIDTH CUDA Implementation – Host Side // Matrix multiplication on the device void Mul(const Matrix A, const Matrix B, Matri