SoCVista深入浅出谈CUDA
- 格式:pdf
- 大小:627.11 KB
- 文档页数:31
cuda原理CUDA,即计算统一设备架构,是由 NVIDIA 公司推出的并行计算平台和编程模型,它利用了 GPU(图形处理器)的并行计算能力,使得普通计算机也能进行高效的科学计算和深度学习等复杂计算任务。
CUDA 的基本原理包括并行计算模型、内存模型、指令集体系结构等。
1. 并行计算模型CUDA 采用的是一种 SPMD(单指令多数据流)并行计算模式,即在多个线程上执行相同的指令,但是每个线程处理的数据不同。
一个 CUDA 程序可以包含多个线程块,每个线程块包含多个线程,每个线程在 GPU 上执行相同的程序代码,但是每个线程处理的数据不同。
CUDA 使用了一种独特的执行模型,即线程束(或称为warp)。
线程束是一个并行计算的处理单元,每个线程束包含最多32个线程,这些线程被编排成一列,当一个线程束被调度时,这列中的各个线程会在一个时钟周期内执行相同的指令,从而实现高效的并行计算。
2. 内存模型CUDA 的内存模型也是一大特点。
在 CPU 和 GPU 之间进行数据传输通常需要繁琐的数据拷贝操作,而进程间通信机制的开销一般也较大,引入许多额外的性能开销。
为了解决这些问题,CUDA 引入了一种新的内存模型,即主机内存(Host Memory)和设备内存(Device Memory)的统一内存模型。
这种内存模型使得主机和设备都能够访问同一片内存,而且无需复制。
CUDA 系统会根据数据的访问模式,自动将数据放置在合适的内存中。
这样,既能够快速地访问设备内存,又能够方便地管理和共享数据。
3. 指令集体系结构CUDA 的指令集体系结构包含了 CUDA 核函数(Kernel Function)和 CUDA 编译器(NVCC)两部分。
核函数相当于一个通用函数,可以在 GPU 上并行地执行。
需要特别指出的是,CUDA 核函数不同于传统的 C/C++ 函数,它不能直接调用其他 C/C++ 函数或标准库函数。
在核函数中,只能使用 CUDA 编译器提供的限制的函数或者自定义的函数进行计算。
第二章CUDA架构2.1 CUDA的编程模型CUDA(Compute Unified Device Architecture),是一种由NVIDIA推出的并行计算架构,非常适合大规模数据密集型计算。
CUDA使GPU的超高计算性能在数据处理和并行计算等通用计算领域发挥优势。
它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。
随着显卡的发展,GPU越来越强大,在计算上已经超越了通用的CPU。
如此强大的芯片如果只是作为显卡会造成计算能力的浪费,因此NVIDIA推出CUDA,让显卡可以用于图像渲染以外的目的。
CUDA 的GPU编程语言基于标准的C语言,通过在标准C语言的基础上增加一小部分关键字,任何有C语言基础的用户都很容易地开发CUDA的应用程序。
CUDA3.0已经开始支持C++和FORTRAN。
2.1.1主机和设备CUDA编程模型在设计结构上采用了异构编程的模式,将CPU作为宿主(Host),GPU作为设备(Device),在同一个系统中可以有同时存在多个设备,但是只能有一个宿主。
在CUDA程序架构中,主程序由CPU来执行,而当遇到数据并行处理的部分,CUDA就会将程序编译成GPU能执行的程序,并传送到GPU。
CUDA使用被称为块(Block)的单元,每个块都由一些CUDA线程组成,线程是CUDA中最小的处理单元,将这些较小的子问题进一步划分为若干更小的细粒度的问题,我们便可以使用线程来解决这些问题了。
对于一个普通的NVIDIA GPU,其CUDA线程数目通常能达到数千个甚至更多,因此,这样的问题划分模型便可以成倍地提升计算机的运算性能。
GPU是由多个流水多处理器构成的,流水处理器以块(Block)为基本调度单元,因此,对于流水处理器较多的GPU,它一次可以处理的块(Block)更多,从而运算速度更快,时间更短。
而反之对于流水处理器较少的GPU,其运算速度便会较慢。
CUDA C是C语言的一个扩展,它允许程序员定义一种被称为内核函数(Kernel Functions)的C函数,内核函数运行在GPU上,一旦启动,CUDA中的每一个线程都将会同时并行地执行内核函数中的代码。
cuda显存分配方式CUDA显存分配方式一、介绍在使用CUDA进行GPU加速的过程中,显存的分配是一个非常重要的环节。
合理的显存分配可以提高程序的性能,而不当的分配则可能导致显存溢出或者浪费。
本文将介绍几种常见的CUDA显存分配方式,并分析它们的优缺点。
二、全局内存分配全局内存是CUDA中最常用的一种显存分配方式。
全局内存是所有线程共享的,可以在所有线程间进行读写操作。
全局内存的分配是通过在设备端使用cudaMalloc函数来完成的。
需要注意的是,全局内存的分配是在主机端进行的,因此需要在主机端进行相应的内存管理。
三、共享内存分配共享内存是一种比全局内存更高效的显存分配方式。
共享内存是每个线程块独享的,可以在线程块内的线程间进行读写操作。
共享内存的分配是通过在内核函数中使用__shared__关键字来完成的。
需要注意的是,共享内存的分配是在设备端进行的,因此不需要在主机端进行相应的内存管理。
四、常量内存分配常量内存是一种比全局内存更高效的显存分配方式。
常量内存也是每个线程块独享的,可以在线程块内的线程间进行读操作,但不能进行写操作。
常量内存的分配是通过在内核函数中使用__constant__关键字来完成的。
需要注意的是,常量内存的分配是在设备端进行的,因此不需要在主机端进行相应的内存管理。
五、纹理内存分配纹理内存是一种特殊的显存分配方式,它主要用于图像处理等需要对数据进行采样的场景。
纹理内存的分配是通过在内核函数中使用texture关键字来完成的。
纹理内存具有缓存机制,可以提高数据的访问效率。
需要注意的是,纹理内存的分配也是在设备端进行的,因此不需要在主机端进行相应的内存管理。
六、动态内存分配除了以上几种静态的显存分配方式,CUDA还提供了动态内存分配的功能。
动态内存分配是在内核函数中根据需要进行的,可以根据具体的计算需求来动态地分配显存。
动态内存分配主要使用cudaMalloc函数来完成,需要注意的是,动态内存的分配和释放都是在设备端进行的。
cuda中文手册摘要:一、引言二、CUDA的定义与背景三、CUDA架构1.统一内存架构2.线程调度3.通信机制四、CUDA编程模型1.编程语言与编译器2.核函数3.设备变量与设备内存五、CUDA应用案例1.图像处理2.深度学习3.高效能计算六、CUDA发展趋势与展望正文:一、引言随着高性能计算需求的不断增长,GPU并行计算逐渐成为研究热点。
NVIDIA公司推出的CUDA(Compute Unified Device Architecture)技术,为通用并行计算提供了一个强大的平台。
本文将对CUDA技术进行详细介绍,包括其定义、背景、架构、编程模型以及应用案例。
二、CUDA的定义与背景CUDA是一种通用并行计算架构,允许开发者利用NVIDIA GPU进行高性能计算。
CUDA起源于NVIDIA对GPU计算潜力的挖掘,旨在为科研、工程和娱乐等领域提供强大的并行计算能力。
三、CUDA架构CUDA架构主要包括统一内存架构、线程调度和通信机制三部分。
1.统一内存架构:CUDA采用统一内存架构,使得CPU和GPU能够共享内存,简化了并行计算的编程复杂度。
2.线程调度:CUDA支持多种线程调度策略,如Warp、Grid等,可以实现高效并行计算。
3.通信机制:CUDA提供了多种通信机制,如共享内存、全局内存和设备内存等,以满足不同应用场景的需求。
四、CUDA编程模型CUDA编程模型包括编程语言、编译器和核心函数等部分。
1.编程语言与编译器:CUDA支持C、C++和Python等编程语言,通过NVIDIA编译器可以将源代码编译为可执行文件。
2.核函数:CUDA中的核函数是一种特殊的函数,能够在GPU上执行。
开发者可以通过核函数实现并行算法。
3.设备变量与设备内存:CUDA提供了设备变量和设备内存,允许开发者直接在GPU上操作数据。
五、CUDA应用案例CUDA在多个领域都有广泛应用,以下为几个典型应用案例:1.图像处理:CUDA可以高效地处理图像和视频数据,如实时图像渲染、图像识别等。
cuda发展历程-回复CUDA,全称Compute Unified Device Architecture,是由NVIDIA开发的一种并行计算平台和编程模型。
CUDA的发展历程可以追溯到2006年,以下将详细介绍CUDA的发展过程。
2006年,NVIDIA推出第一版的CUDA。
当时,CUDA的主要目标是使用GPU(图形处理器)进行通用计算。
在过去,GPU主要用于图形渲染,但NVIDIA意识到GPU的强大并行计算能力,因此决定为开发者提供编程接口,使其能够利用GPU进行更大范围的计算任务。
第一版的CUDA 主要支持C语言,并提供了一套对开发者友好的API,使其能够方便地进行GPU编程。
2007年,NVIDIA发布了CUDA Toolkit 1.0,这是一个全面的开发工具包,为开发者提供了编译器、调试器、性能分析器等工具,以及一系列的开发库。
这些工具大大简化了开发GPU应用程序的过程,使得更多的开发者可以参与到GPU计算的开发中来。
随着CUDA的不断推出和开发者的参与,越来越多的应用程序开始使用GPU进行加速计算。
2008年,NVIDIA发布了CUDA 2.0版本,并引入了线程块和线程束的概念,使得开发者可以更好地管理和利用GPU上的计算资源。
此外,CUDA 2.0还支持动态并行调度,使得开发者能够更加灵活地控制并行计算的流程。
在接下来的几年里,NVIDIA持续不断地更新和改进CUDA平台。
2010年,CUDA 3.0发布,引入了一种新的内存模型,即统一虚拟寻址(Unified Virtual Addressing,UVA)。
这一功能使得开发者可以更方便地在CPU 和GPU之间共享内存,并且不再需要显示地进行内存拷贝。
UVA的引入大大简化了编程的流程,提高了开发效率。
2012年,NVIDIA发布了CUDA 5.0版本,引入了动态并行调度的新特性,这使得开发者能够更好地响应计算需求的变化。
此外,CUDA 5.0还支持GPU加速的MPI(Message Passing Interface),使得CUDA可以更好地与传统的MPI编程模型结合起来,实现更高效的并行计算。
NVIDIA CUDA计算统一设备架构编程指南版本 2.06 /7 / 2008目录目录 (iii)图表目录 (vi)第1 章简介 (1)1.1 CUDA:可伸缩并行编程模型 (1)1.2 GPU:高度并行化的多线程、众核处理器 (1)1.3 文档结构 (3)第2章编程模型 (4)2.1 线程层次结构 (4)2.2 存储器层次结构 (6)2.3 宿主和设备 (7)2.4 软件栈 (8)2.5 计算能力 (9)第3 章GPU 实现 (10)3.1 具有片上共享存储器的一组SIMT 多处理器 (10)3.2 多个设备 (12)3.3 模式切换 (12)第4 章应用程序编程接口 (13)4.1 C 编程语言的扩展 (13)4.2 语言扩展 (13)4.2.1 函数类型限定符 (13)4.2.1.1 __device__ (13)4.2.1.2 __global__ (14)4.2.1.3 __host__ (14)4.2.1.4 限制 (14)4.2.2 变量类型限定符 (14)4.2.2.1 __device__ (14)4.2.2.2 __constant__ (14)4.2.2.3 __shared__ (15)4.2.2.4 限制 (15)4.2.3 执行配置 (16)4.2.4 内置变量 (16)4.2.4.1 gridDim (16)4.2.4.2 blockIdx (16)4.2.4.3 blockDim (16)4.2.4.4 threadIdx (16)4.2.4.5 warpSize (17)4.2.4.6 限制 (17)4.2.5 使用NVCC 进行编译 (17)4.2.5.1 __noinline__ (17)4.2.5.2 #pragma unroll (17)4.3 通用运行时组件 (18)4.3.1 内置向量类型 (18)4.3.1.1 char1、uchar1、char2、uchar2、char3、uchar3、char4、uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1、int2、uint2、int3、uint3、int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、long4、ulong4、float1、float2、float3、float4、double2 (18)4.3.1.2 dim3 类型 (18)4.3.2 数学函数 (18)4.3.3 计时函数 (18)4.3.4 纹理类型 (19)4.3.4.1 纹理参考声明 (19)4.3.4.2 运行时纹理参考属性 (19)4.3.4.3 来自线性存储器的纹理与来自CUDA 数组的纹理 (20)4.4 设备运行时组件 (20)4.4.1 数学函数 (20)4.4.2 同步函数 (20)4.4.3 纹理函数 (21)4.4.3.1 来自线性存储器的纹理 (21)4.4.3.2 来自CUDA 数组的纹理 (21)4.4.4 原子函数 (21)4.4.5 Warp vote 函数 (22)4.5 宿主运行时组件 (22)4.5.1 一般概念 (22)4.5.1.1 设备 (22)4.5.1.2 存储器 (23)4.5.1.3 OpenGL 互操作性 (23)4.5.1.4 Direct3D 互操作性 (23)4.5.1.5 异步并发执行 (23)4.5.2 运行时API (24)4.5.2.1 初始化 (24)4.5.2.2 设备管理 (24)4.5.2.3 存储器管理 (25)4.5.2.4 流管理 (26)4.5.2.5 事件管理 (26)4.5.2.6 纹理参考管理 (27)4.5.2.7 OpenGL 互操作性 (28)4.5.2.8 Direct3D 互操作性 (28)4.5.2.9 使用设备模拟模式进行调试 (29)4.5.3 驱动程序API (30)4.5.3.1 初始化 (30)4.5.3.2 设备管理 (30)4.5.3.3 上下文管理 (30)4.5.3.4 模块管理 (31)4.5.3.5 执行控制 (31)4.5.3.6 存储器管理 (32)4.5.3.7 流管理 (33)4.5.3.8 事件管理 (33)4.5.3.9 纹理参考管理 (34)4.5.3.10 OpenGL 互操作性 (34)4.5.3.11 Direct3D 互操作性 (35)第5 章性能指南 (36)5.1 指令性能 (36)5.1.1指令吞吐量 (36)5.1.1.1 数学指令 (36)5.1.1.2 控制流指令 (37)5.1.1.3 存储器指令 (37)5.1.1.4 同步指令 (38)5.1.2 存储器带宽 (38)5.1.2.1 全局存储器 (38)5.1.2.2 本地存储器 (44)5.1.2.3 常量存储器 (44)5.1.2.4 纹理存储器 (44)5.1.2.5 共享存储器 (44)5.1.2.6 寄存器 (49)5.2 每个块的线程数量 (50)5.3 宿主和设备间的数据传输 (50)5.4 纹理拾取与全局或常量存储器读取的对比 (51)5.5 整体性能优化策略 (51)第6 章矩阵乘法示例 (53)6.1 概述 (53)6.2 源代码清单 (54)6.3 源代码说明 (55)6.3.1 Mul() (55)6.3.2 Muld() (55)附录A技术规范 (57)A.1 一般规范 (57)A.1.1 计算能力1.0 的规范 (57)A.1.2 计算能力1.1 的规范 (58)A.1.3 计算能力1.2 的规范 (58)A.1.4 计算能力1.3 的规范 (58)A.2 浮点标准 (58)附录B标准数学函数 (60)B.1 通用运行时组件 (60)B.1.1 单精度浮点函数 (60)B.1.2 双精度浮点函数 (62)B.1.3 整型函数 (63)B.2 设备运行时组件 (63)B.2.1 单精度浮点函数 (64)B.2.2 双精度浮点函数 (65)B.2.3 整型函数 (65)附录C原子函数 (66)C.1 数学函数 (66)C.1.1 atomicAdd() (66)C.1.2 atomicSub() (66)C.1.3 atomicExch() (66)C.1.4 atomicMin() (66)C.1.5 atomicMax() (67)C.1.6 atomicInc() (67)C.1.7 atomicDec() (67)C.1.8 atomicCAS() (67)C.2 位逻辑函数 (67)C.2.1 atomicAnd() (67)C.2.2 atomicOr() (68)C.2.3 atomicXor() (68)附录D纹理拾取 (69)D.1 最近点采样 (69)D.2 线性过滤 (70)D.3 表查找 (70)录表目录图表目图1-1. CPU 和GPU 的每秒浮点运算次数和存储器带宽...... ......... ..... .............. (2)图1-2. GPU 中的更多晶体管用于数据处理...... .............. .......... .............. .............. (2)图2-1. 线程块网格.......................................... ....... ..................... (6)图2-2. 存储器层次结构................................. .............. . (7)图2-3. 异构编程............................................... ............. .............. .............. (8)图2-4. CUDA软件栈................ .............. .............. ..................... .............. .. .. (9)图3-1. 硬件模型................................................................... .............. . (11)图4-1. 库上下文管理......................................................... .............. ....... ......... (31)图5-1. 存储器合并后的存储器访问模式示例................... .............. .............. .............. .. ........ ..40 图5-2. 未为计算能力是1.0 或1.1 的设备进行存储器合并的全局存储器访问模式示例. (41)图5-3. 未为计算能力是1.0 或1.1 的设备进行存储器合并的全局存储器访问模式示例... .42 图5-4. 计算能力为1.2 或更高的设备的全局存储器访问示例..... .............. ............ .. .. (43)图5-5. 无存储体冲突的共享存储器访问模式示例... .............. .............. .............. .............. (46)图5-6. 无存储体冲突的共享存储器访问模式示例... .............. .............. .............. .............. (47)图5-7. 有存储体冲突的共享存储器访问模式示例......... .............. .............. .............. (48)图5-8. 使用广播机制的共享存储器读取访问模式示例... .............. .............. .............. (49)图6-1. 矩阵乘法........................................................... .............. .............. (53)第 1 章简介1.1 CUDA:可伸缩并行编程模型多核CPU 和众核(manycore)GPU 的出现意味着主流处理器芯片已进入并行时代。
cuda .c后缀-回复CUDA(Compute Unified Device Architecture)是一种并行计算平台和编程模型,由NVIDIA公司开发。
它允许开发人员利用GPU(图形处理器)的强大计算能力,加速并行计算任务的执行速度。
在本文中,我们将逐步回答关于CUDA的相关问题。
第一部分:什么是CUDA?CUDA是一种并行计算平台和编程模型,它利用GPU的并行计算能力来加速计算任务。
传统上,GPU主要用于图形渲染,而CUDA将其转变为通用计算设备。
CUDA利用GPU内部的大量并行计算单元,高效地执行大规模的并行计算任务。
第二部分:为什么使用CUDA?CUDA相比于传统的CPU计算具有很多优势。
首先,GPU具有数以千计的并行计算单元,远远超过CPU的数量,因此可以同时执行更多的计算任务。
其次,GPU具有高带宽的内存访问,可以更快地从内存中读取数据。
最重要的是,CUDA提供了一组强大的编程工具和库,使开发人员可以方便地利用GPU的计算能力。
第三部分:CUDA的工作原理是什么?CUDA的工作原理可以概括为以下几个步骤。
首先,开发人员使用CUDA 编程语言(通常是C)编写并行计算的代码,并使用CUDA工具链将其编译成机器代码。
然后,开发人员将数据从主机(CPU)内存复制到设备(GPU)内存中。
接下来,开发人员将并行计算的任务分解成多个线程块,每个线程块包含多个线程。
每个线程独立地执行计算任务,并访问设备内存中的数据。
最后,开发人员将计算结果从设备内存复制回主机内存,并进行后续处理。
第四部分:如何使用CUDA进行并行计算?要使用CUDA进行并行计算,首先需要安装NVIDIA的显卡驱动和CUDA 开发工具包。
然后,开发人员可以使用CUDA编程语言编写并行计算的代码。
在代码中,需要使用特定的语法和API来定义并行计算任务,并操作设备内存中的数据。
编写完代码后,使用CUDA工具链进行编译,并将生成的可执行文件在GPU上运行。
在深入讨论cudasetdevice参数之前,我们首先要了解CUDA是什么以及它的重要性。
CUDA是NVIDIA推出的一种并行计算评台和编程模型,它允许开发人员使用C语言、C++和Fortran等语言来利用NVIDIA GPU的并行计算能力。
CUDA在高性能计算、科学计算、人工智能、深度学习等领域有着广泛的应用,是一个非常重要的工具。
cudasetdevice参数是CUDA中的一个非常重要的参数,它用于设置当前线程要使用的GPU设备。
在多GPU环境下,通过设置cudasetdevice参数,我们可以指定每个线程要使用的GPU设备,从而实现并行计算的目的。
在许多CUDA程序中,特别是在涉及大规模并行计算或深度学习任务时,我们都需要充分利用多个GPU设备来加速计算过程。
接下来,让我们来深入探讨cudasetdevice参数的使用方法和其重要性。
当我们在编写CUDA程序时,需要考虑到并行计算的特性。
在使用多个GPU设备时,通过设置cudasetdevice参数,我们可以有效地将计算任务分配给不同的GPU设备,从而实现并行计算。
这可以大大提高计算任务的执行效率,特别是对于那些需要大量计算资源的任务而言。
在深度学习训练过程中,通过合理地设置cudasetdevice参数,我们可以将不同的神经网络层分配给不同的GPU设备,从而加速训练过程。
cudasetdevice参数的灵活性也是其重要性的体现。
在实际应用中,我们常常会遇到需要动态地设置GPU设备的场景。
通过在程序中动态设置cudasetdevice参数,我们可以根据实际情况来选择要使用的GPU设备,从而实现资源的灵活分配。
这对于高性能计算环境或云计算评台来说尤为重要,因为在这些环境中,GPU设备的数量和性能往往是动态变化的。
合理地使用cudasetdevice参数可以使我们的计算任务更具适应性和灵活性。
我个人对cudasetdevice参数的理解是,它是多GPU环境下实现并行计算的重要工具之一。
深入浅出谈CUDAHotballCUDA 是 NVIDIA 的 GPGPU 模型,它使用 C 语言为基础,可以直接以大多数人熟悉的 C”CUDA是什么?能吃吗?编者注:NVIDIA的GeFoce 8800GTX发布后,它的通用计算架构CUDA经过一年多的推广后,易了解CUDA,我们征得Hotball的本人同意,发表他最近亲自撰写的本文。
这篇文章的特点是深入浅出,也包含了hotball本人编写一些简单CUDA程序的亲身体验,对于希望了解CUDA的读者来说是非常不错的入门文章,PCINLIFE对本文的发表没有作任何的删减,主要是把一些台湾的词汇转换成大陆的词汇以及作了若干"编者注"的注释。
现代的显示芯片已经具有高度的可程序化能力,由于显示芯片通常具有相当高的内存带宽,以及大量的执行单元,因此开始有利用显示芯片来帮助进行一些计算工作的想法,即 GPGPU。
CUDA 即是NVIDIA 的 GPGP U 模型。
NVIDIA 的新一代显示芯片,包括 GeForce 8 系列及更新的显示芯片都支持 CUDA。
NVIDIA 免费提供 CUDA 的开发工具(包括 Wind ows 版本和 Li nux 版本)、程序范例、文件等等,可以在CUDA Zone 下载。
GPGPU 的优缺点使用显示芯片来进行运算工作,和使用 CPU 相比,主要有几个好处:1.显示芯片通常具有更大的内存带宽。
例如,NVIDIA 的 GeF orce 8800GTX 具有超过50GB/s 的内存带宽,而目前高阶 CPU 的内存带宽则在 10GB/s 左右。
.2.显示芯片具有更大量的执行单元。
例如 G eForce 8800GTX 具有 128 个 "strea mprocessors",频率为 1.35GHz。
CPU 频率通常较高,但是执行单元的数目则要少得多。
3.和高阶 CPU相比,显卡的价格较为低廉。
例如目前一张 G eForce 8800GT 包括512MB 内存的价格,和一颗 2.4GHz 四核心 CP U 的价格相若。
当然,使用显示芯片也有它的一些缺点:1.显示芯片的运算单元数量很多,因此对于不能高度并行化的工作,所能带来的帮助就不大。
2.显示芯片目前通常只支持 32 bits 浮点数,且多半不能完全支持 IEEE 754 规格,有些运算的精确度可能较低。
目前许多显示芯片并没有分开的整数运算单元,因此整数运算的效率较差。
3.显示芯片通常不具有分支预测等复杂的流程控制单元,因此对于具有高度分支的程序,效率会比较差。
4.目前 GPGP U 的程序模型仍不成熟,也还没有公认的标准。
例如 NVIDIA和AMD/ATI 就有各自不同的程序模型。
整体来说,显示芯片的性质类似 str eam processor,适合一次进行大量相同的工作。
CPU 则比较有弹性,能同时进行变化较多的工作。
CUDA 架构CUDA 是 NVIDI A 的 GPGPU 模型,它使用 C 语言为基础,可以直接以大多数人熟悉的 C 语言,写出在显示芯片上执行的程序,而不需要去学习特定的显示芯片的指令或是特殊的结构。
在 CUDA 的架构下,一个程序分为两个部份:host 端和 devi ce 端。
Host 端是指在 CPU 上执行的部份,而 de vice 端则是在显示芯片上执行的部份。
Device 端的程序又称为 "kern el"。
通常 hos t 端程序会将数据准备好后,复制到显卡的内存中,再由显示芯片执行 de vice 端程序,完成后再由 ho st 端程序将结果从显卡的内存中取回。
由于 CPU 存取显卡内存时只能透过 PCI Express 接口,因此速度较慢(PCI Express x16 的理论带宽是双向各 4G B/s),因此不能太常进行这类动作,以免降低效率。
在 CUDA 架构下,显示芯片执行时的最小单位是thread。
数个 threa d 可以组成一个block。
一个 bl ock 中的 thre ad 能存取同一块共享的内存,而且可以快速进行同步的动作。
每一个 bloc k 所能包含的 thre ad 数目是有限的。
不过,执行相同程序的 blo ck,可以组成grid。
不同 b lock 中的 t hread 无法存取同一个共享的内存,因此无法直接互通或进行同步。
因此,不同 b lock 中的 t hread 能合作的程度是比较低的。
不过,利用这个模式,可以让程序不用担心显示芯片实际上能同时执行的 thre ad 数目限制。
例如,一个具有很少量执行单元的显示芯片,可能会把各个 bl ock 中的 th read 顺序执行,而非同时执行。
不同的 gri d 则可以执行不同的程序(即 ker nel)。
Grid、block 和 thre ad 的关系,如下图所示:每个 thre ad 都有自己的一份 register和 lo cal me mory 的空间。
同一个 bl ock 中的每个thread 则有共享的一份 share memory。
此外,所有的 thre ad(包括不同 b lock 的 thre ad)都共享一份 g lobal memory、constant memory、和 text ure memory。
不同的 gri d 则有各自的 gl obal memory、constant memory 和 te xture memory。
这些不同的内存的差别,会在之后讨论。
执行模式由于显示芯片大量并行计算的特性,它处理一些问题的方式,和一般 CPU 是不同的。
主要的特点包括:1.内存存取 la tency 的问题:CPU 通常使用 ca che 来减少存取主内存的次数,以避免内存 la tency 影响到执行效率。
显示芯片则多半没有 ca che(或很小),而利用并行化执行的方式来隐藏内存的 la tency(即,当第一个 thre ad 需要等待内存读取结果时,则开始执行第二个 t hread,依此类推)。
2.分支指令的问题:CPU 通常利用分支预测等方式来减少分支指令造成的 p ipelinebubble。
显示芯片则多半使用类似处理内存 la tency 的方式。
不过,通常显示芯片处理分支的效率会比较差。
因此,最适合利用 CUDA 处理的问题,是可以大量并行化的问题,才能有效隐藏内存的latency,并有效利用显示芯片上的大量执行单元。
使用 CUDA 时,同时有上千个 t hread 在执行是很正常的。
因此,如果不能大量并行化的问题,使用 CUDA 就没办法达到最好的效率了。
CUDA Toolkit的安装目前 NVIDI A 提供的 CUDA Toolkit(可从这里下载)支持 Wind ows (32 bits 及 64 bits 版本)及许多不同的 Li nux 版本。
CUDA Toolkit 需要配合 C/C++ compiler。
在 Wi ndows 下,目前只支持 V isual Studio 7.x 及Visual St udio 8(包括免费的Visual Studio C++ 2005 Ex press)。
Visual S tudio 6 和 gcc 在Windows 下是不支援的。
在 Li nux 下则只支援 gcc。
这里简单介绍一下在 Wi ndows 下设定并使用 CUDA 的方式。
下载及安装在 Win dows 下,CUDA Toolkit 和 CU DA SDK 都是由安装程序的形式安装的。
CUDA Toolkit 包括 CUDA 的基本工具,而 CUDA SDK 则包括许多范例程序以及链接库。
基本上要写CUDA 的程序,只需要安装 CUDA Toolkit 即可。
不过 CUDA SDK 仍值得安装,因为里面的许多范例程序和链接库都相当有用。
CUDA Toolkit 安装完后,预设会安装在 C:\CUDA 目录里。
其中包括几个目录:•bin -- 工具程序及动态链接库•doc -- 文件•include -- header 檔•lib -- 链接库档案•open64 -- 基于 Op en64 的 CUDA compiler•src -- 一些原始码安装程序也会设定一些环境变量,包括:•CUDA_BIN_PATH -- 工具程序的目录,默认为 C:\CUDA\b in•CUDA_INC_PATH -- header 文件的目录,默认为 C:\CUDA\in c•CUDA_LIB_PATH -- 链接库文件的目录,默认为 C:\CUDA\li b在 Visual Studio 中使用 CUDACUDA 的主要工具是 nvcc,它会执行所需要的程序,将 CUDA 程序代码编译成执行档(或object 檔) 。
在 Visu al Studio 下,我们透过设定custom build tool 的方式,让 Visu al Studio 会自动执行 nvcc。
这里以Visual Studio 2005 为例:1.首先,建立一个 W in32 Console 模式的 pro ject(在 A pplication Settings 中记得勾选Empty project),并新增一个档案,例如 ma in.cu。
2.在 ma in.cu 上右键单击,并选择Properties。
点选General,确定Tool的部份是选择Custom Build Tool。
3.选择Custom Build Step,在 C ommand Line 使用以下设定:o Release 模式:"$(CUDA_BIN_PATH)\nvcc.exe" -ccbin "$(VCIn stallDir)bin" -c-DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT-I"$(CUDA_INC_PATH)" -o $(ConfigurationName)\$(InputName).obj $(InputFileName) o Debug 模式:"$(CUDA_BIN_PATH)\nvcc.exe" -ccb in "$(VCIn stallDir)bin" -c-D_DEBUG -D WIN32 -D_CONSOLE -D_MBCS -Xco mpiler/EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I"$(CUDA_INC_PATH)" -o$(ConfigurationName)\$(InputName).obj $(InputFileName)4.如果想要使用软件仿真的模式,可以新增两个额外的设定:o EmuRelease 模式:"$(CUDA_BIN_PATH)\nvcc.exe" -ccb in "$(VCInstallDir)bin"-deviceemu -c -DWIN32 -D_C ONSOLE -D_MBCS -Xcompiler/EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT -I"$(CUDA_INC_PATH)" -o$(ConfigurationName)\$(InputName).obj $(InputFileName)o EmuDebug 模式:"$(CUDA_BIN_PATH)\nvcc.exe" -ccb in "$(VCInstallDir)bin"-deviceemu -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xc ompiler/EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I"$(CUDA_INC_PATH)" -o$(ConfigurationName)\$(InputName).obj $(InputFileName)5.对所有的配置文件,在Custom Build Step的Outputs中加入$(ConfigurationName)\$(InputName).obj。