CUDA并行计算架构及编程
cuda使用教程

cuda使用教程CUDA(Compute Unified Device Architecture)是一种用于并行计算的平台和编程模型,可以利用GPU(Graphics Processing Unit,图形处理器)的强大计算能力来加速各种应用程序。
本文将为读者介绍如何使用CUDA进行并行计算,并提供一些基本的教程和示例。
要使用CUDA进行并行计算,我们需要一个支持CUDA的显卡。
大多数NVIDIA的显卡都支持CUDA,可以到NVIDIA官方网站查看显卡的兼容性列表。
另外,我们还需要安装NVIDIA的CUDA Toolkit,这是一个开发和运行CUDA程序的软件包。
安装完CUDA Toolkit后,我们就可以开始编写CUDA程序了。
CUDA 程序主要由两部分组成:主机代码(Host Code)和设备代码(Device Code)。
主机代码运行在CPU上,用于控制和管理CUDA设备;设备代码运行在GPU上,用于实际的并行计算。
在CUDA中,我们使用C/C++语言来编写主机代码,使用CUDA C/C++扩展来编写设备代码。
CUDA C/C++扩展是一种特殊的语法,用于描述并行计算的任务和数据的分配。
通过在设备代码中定义特定的函数(称为内核函数),我们可以在GPU上并行地执行这些函数。
下面是一个简单的示例,展示了如何使用CUDA计算两个向量的和:```c++#include <stdio.h>__global__ void vectorAdd(int* a, int* b, int* c, int n) { int tid = blockIdx.x * blockDim.x + threadIdx.x;if (tid < n) {c[tid] = a[tid] + b[tid];}}int main() {int n = 1000;int *a, *b, *c; // Host arraysint *d_a, *d_b, *d_c; // Device arrays// Allocate memory on hosta = (int*)malloc(n * sizeof(int));b = (int*)malloc(n * sizeof(int));c = (int*)malloc(n * sizeof(int));// Initialize host arraysfor (int i = 0; i < n; i++) {a[i] = i;b[i] = i;}// Allocate memory on devicecudaMalloc((void**)&d_a, n * sizeof(int));cudaMalloc((void**)&d_b, n * sizeof(int));cudaMalloc((void**)&d_c, n * sizeof(int));// Copy host arrays to devicecudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);// Launch kernelint block_size = 256;int grid_size = (n + block_size - 1) / block_size;vectorAdd<<<grid_size, block_size>>>(d_a, d_b, d_c, n); // Copy result back to hostcudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);// Print resultfor (int i = 0; i < n; i++) {printf("%d ", c[i]);}// Free memoryfree(a);free(b);free(c);cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return 0;}```在这个示例中,我们首先定义了一个内核函数`vectorAdd`,用于计算两个向量的和。
cuda原理

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架构

第二章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是NVIDIA推出的一种并行计算平台和编程模型。
它允许开发人员利用GPU的高度并行性能执行复杂的计算任务,从而加速应用程序的运行速度。
下面将从CUDA的概述、架构和应用场景三个方面来介绍CUDA显卡。
首先,CUDA概述。
CUDA是Compute Unified Device Architecture的缩写,意为统一计算架构。
它最早是为GPU而设计的,而不是传统的CPU。
CUDA架构将GPU的大规模并行计算能力与CPU的强大控制能力结合起来,使得开发人员可以充分利用GPU的性能优势。
CUDA利用GPU的多个核心进行并行计算,从而更高效地完成任务。
其次,CUDA架构。
CUDA显卡的核心是由大量的计算单元组成的。
每个计算单元都包含ALU(算术逻辑单元)和存储单元。
这些计算单元可以同时处理大量的数据,并且具备高度的并行计算能力。
CUDA采用了SIMD(Single Instruction, Multiple Data)执行模型,即同一指令可以同时作用于多个数据。
这使得程序可以在GPU上并行地执行计算任务,从而大幅提高应用程序的性能。
最后,CUDA的应用场景。
由于CUDA使用GPU进行并行计算,它在很多领域都有很好的应用前景。
首先,科学计算领域。
CUDA可以加速大规模数据的计算、模拟和分析,例如气象学、物理学、生物学等。
其次,人工智能领域。
CUDA可以加速机器学习、深度学习和神经网络的训练和推理过程,提高模型的训练速度和预测性能。
再次,图形渲染领域。
CUDA可以加速三维渲染、光线追踪和图像处理等图形计算任务,提供更快速、更真实、更高品质的图像效果。
综上所述,CUDA显卡作为NVIDIA推出的一种并行计算平台和编程模型,在科学计算、人工智能和图形渲染等领域都有广泛应用。
通过充分发挥GPU的并行计算能力,CUDA可以大幅提高应用程序的运行速度,提供更好的用户体验。
2024版CUDA编程入门极简教程

行划分,每个线程处理一部分数据;任务并行是将任务划分为多个子任
务,每个线程执行一个子任务。
02
共享内存与全局内存
CUDA提供共享内存和全局内存两种存储空间。共享内存位于处理器内
部,访问速度较快,可用于线程间通信;全局内存位于处理器外部,访
问速度较慢,用于存储大量数据。
03
异步执行与流
CUDA支持异步执行,即CPU和GPU可以同时执行不同的任务。通过创
2023
PART 02
CUDA环境搭建与配置
REPORTING
安装CUDA工具包
下载CUDA工具包
01
访问NVIDIA官网,下载适用于您的操作系统的CUDA工具包。
安装CUDA工具包
02
按照安装向导的指示,完成CUDA工具包的安装。
验证安装
03
安装完成后,可以通过运行CUDA自带的示例程序来验证算,每个线 程处理一个子任务。计算完成后, 将结果从设备内存传输回主机内 存,并进行必要的后处理操作。
2023
PART 05
CUDA优化策略与技巧
REPORTING
优化内存访问模式
合并内存访问
通过确保线程访问连续的内存地址,最大化内 存带宽利用率。
使用共享内存
利用CUDA的共享内存来减少全局内存访问, 提高数据重用。
避免不必要的内存访问
精心设计算法和数据结构,减少不必要的内存读写操作。
减少全局内存访问延迟
使用纹理内存和常量内存
利用CUDA的特殊内存类型,如纹理内存和常量内存,来加速数 据访问。
数据预取和缓存
通过预取数据到缓存或寄存器中,减少全局内存访问次数。
展望未来发展趋势
CUDA与深度学习
cuda工作原理

cuda工作原理CUDA (Compute Unified Device Architecture) is a parallel computing platform and programming model created by NVIDIA. It allows software developers to use a CUDA-enabled graphics processing unit (GPU) for general purpose processing, which can greatly accelerate computationally intensive tasks. CUDA works by offloading the computationally intensive parts of an application to the GPU, while the rest of the code runs on the CPU.CUDA(计算统一设备架构)是由NVIDIA创建的并行计算平台和编程模型。
它允许软件开发人员将基于CUDA的图形处理单元(GPU)用于通用目的处理,从而可以大大加速需要大量计算的任务。
CUDA通过将应用程序的计算密集部分卸载到GPU上来运行,而其余部分的代码则在CPU上运行。
From a technical perspective, CUDA works by allowing developers to write C, C++ or Fortran code that can be compiled and executed on the GPU. This means that developers can use the familiar syntax and structure of these programming languages to harness the power of the GPU for parallel processing. The CUDA programming model alsoprovides a set of libraries, tools, and extensions that allow developers to optimize their code for GPU execution.技术上来看,CUDA允许开发人员编写能够在GPU上编译和执行的C、C++或Fortran代码。
CUDA编程之快速入门
CUDA编程之快速⼊门CUDA(Compute Unified Device Architecture)的中⽂全称为计算统⼀设备架构。
做图像视觉领域的同学多多少少都会接触到CUDA,毕竟要做性能速度优化,CUDA是个很重要的⼯具,CUDA是做视觉的同学难以绕过的⼀个坑,必须踩⼀踩才踏实。
CUDA编程真的是⼊门容易精通难,具有计算机体系结构和C语⾔编程知识储备的同学上⼿CUDA编程应该难度不会很⼤。
本⽂章将通过以下五个⽅⾯帮助⼤家⽐较全⾯地了解CUDA编程最重要的知识点,做到快速⼊门:1. GPU架构特点2. CUDA线程模型3. CUDA内存模型4. CUDA编程模型5. CUDA应⽤⼩例⼦1. GPU架构特点⾸先我们先谈⼀谈串⾏计算和并⾏计算。
我们知道,⾼性能计算的关键利⽤多核处理器进⾏并⾏计算。
当我们求解⼀个计算机程序任务时,我们很⾃然的想法就是将该任务分解成⼀系列⼩任务,把这些⼩任务⼀⼀完成。
在串⾏计算时,我们的想法就是让我们的处理器每次处理⼀个计算任务,处理完⼀个计算任务后再计算下⼀个任务,直到所有⼩任务都完成了,那么这个⼤的程序任务也就完成了。
如下图所⽰,就是我们怎么⽤串⾏编程思想求解问题的步骤。
但是串⾏计算的缺点⾮常明显,如果我们拥有多核处理器,我们可以利⽤多核处理器同时处理多个任务时,⽽且这些⼩任务并没有关联关系(不需要相互依赖,⽐如我的计算任务不需要⽤到你的计算结果),那我们为什么还要使⽤串⾏编程呢?为了进⼀步加快⼤任务的计算速度,我们可以把⼀些独⽴的模块分配到不同的处理器上进⾏同时计算(这就是并⾏),最后再将这些结果进⾏整合,完成⼀次任务计算。
下图就是将⼀个⼤的计算任务分解为⼩任务,然后将独⽴的⼩任务分配到不同处理器进⾏并⾏计算,最后再通过串⾏程序把结果汇总完成这次的总的计算任务。
所以,⼀个程序可不可以进⾏并⾏计算,关键就在于我们要分析出该程序可以拆分出哪⼏个执⾏模块,这些执⾏模块哪些是独⽴的,哪些⼜是强依赖强耦合的,独⽴的模块我们可以试着设计并⾏计算,充分利⽤多核处理器的优势进⼀步加速我们的计算任务,强耦合模块我们就使⽤串⾏编程,利⽤串⾏+并⾏的编程思路完成⼀次⾼性能计算。
cuda的线程组织结构
cuda的线程组织结构CUDA是一种用于并行计算的编程模型和平台,可以在NVIDIA的GPU 上进行高效的并行计算。
在CUDA中,线程组织结构是非常重要的,它决定了并行计算的方式和效率。
本文将介绍CUDA的线程组织结构,包括线程块、网格和线程索引等概念。
1. 线程块(Thread Block)线程块是CUDA中最基本的并行计算单元,它由一组线程组成。
这些线程可以同时执行相同的指令,但是它们有各自的线程索引。
线程块中的线程可以通过共享内存进行通信和协作。
线程块是在GPU上分配的,一个GPU卡上可以有多个线程块同时执行。
线程块的数量和大小是可以在CUDA程序中进行设置的,通常选择的线程块大小是32的倍数。
2. 网格(Grid)网格是线程块的集合,它是一个三维的结构。
网格中的线程块可以并行地执行,每个线程块可以有不同的线程索引。
网格中的线程块可以通过全局内存进行通信和协作。
网格是在Host端(CPU)上进行设置的,通过指定网格的大小和线程块的大小来控制并行计算的规模。
网格的大小可以是一维、二维或三维的,取决于问题的特性。
3. 线程索引(Thread Index)线程索引是用来标识线程在线程块或网格中的位置的。
在CUDA程序中,可以通过内置的变量来获取线程索引。
线程索引包括三个部分:x、y和z,分别表示线程在网格中的位置。
线程索引可以用来确定线程需要处理的数据的位置。
在访问全局内存或共享内存时,线程索引可以用来计算数据的地址。
同时,线程索引也可以用来实现线程的分工和协作。
4. 线程束(Warp)线程束是一组连续的线程,它是GPU上最小的调度单元。
在一个线程束中的线程会同时执行相同的指令,但是每个线程有自己的线程索引。
线程束的大小是固定的,通常是32。
线程束的概念是为了实现SIMD(Single Instruction Multiple Data)的并行计算。
GPU上的硬件会将一组线程束作为一个调度单元,同时执行相同的指令,以充分利用GPU的并行计算能力。
cuda常用的数学库
cuda常用的数学库CUDA(Compute Unified Device Architecture)是一种并行计算平台和编程模型,用于利用GPU(图形处理器)进行高性能计算。
在CUDA编程中,数学库是非常重要的工具,它提供了丰富的数学函数来支持各种数值计算任务。
本文将介绍CUDA常用的数学库及其功能。
1. CUDA Math Library(cuBLAS):cuBLAS是CUDA的基本线性代数库,提供了丰富的线性代数运算函数,如矩阵乘法、矩阵转置、矩阵求逆等。
它能够充分利用GPU的并行计算能力,加速线性代数运算任务。
2. CUDA Random Number Generation(cuRAND):cuRAND是CUDA的随机数生成库,提供了各种随机数生成函数,如均匀分布、正态分布、泊松分布等。
它能够高效地生成大量随机数,并利用GPU的并行计算能力进行加速。
3. CUDA Fast Fourier Transform(cuFFT):cuFFT是CUDA的快速傅里叶变换库,提供了各种快速傅里叶变换函数,如一维傅里叶变换、二维傅里叶变换、多维傅里叶变换等。
它能够高效地进行信号和图像处理任务,如滤波、频域分析等。
4. CUDA Sparse Linear Algebra Library(cuSPARSE):cuSPARSE 是CUDA的稀疏线性代数库,提供了各种稀疏矩阵运算函数,如稀疏矩阵乘法、稀疏矩阵转置等。
它能够高效地处理大规模稀疏矩阵,节省内存和计算资源。
5. CUDA Performance Primitives(NPP):NPP是CUDA的性能优化库,提供了各种图像和信号处理函数,如图像滤波、图像变换、信号采样等。
它能够高效地进行图像和信号处理任务,加速计算过程。
6. CUDA Math Library(cuBLAS):cuBLAS是CUDA的基本线性代数库,提供了丰富的线性代数运算函数,如矩阵乘法、矩阵转置、矩阵求逆等。
GPU并行计算与CUDA编程02
GPU并行计算与CUDA编程 第2课DATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)本周介绍内容 1. 并行编程的通讯模式 1.1 什么是通讯模式 1.2 常见通讯模式的类型和原来2. GPU硬件模式 2.1 GPU,SM(流处理器),Kernel(核),thread block(线程块),线程3. CUDA编程模型 3.1 CUDA编程模型的优点和缺点 3.2 CUDA编程编程模型的一些原则 3.3 CUDA内存模型 3.4 同步性synchronisation和屏障barrier 3.5 编程模型4. 开始编写CUDA程序 4.1 GPU程序的一般步骤 4.2 第一个GPU程序讲解——并行求平方 DATAGURU专业数据分析社区第一版 讲师 罗韵 (WeChat:LaurenLuoYun)1. 并行编程的通讯模式(Communication Patterns)1.1 什么是通讯模式 1.2 通讯模式的类型和原理DATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)1.1 通讯模式(Communication Patterns) 并行计算:非常多的线程在合作解决一个问题Communication内存:DATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)1.2 常见通信模式 1. 映射Map 2. 聚合gather 3. 分散scatter 4. 模板stencil 5. 转换transpose6. 压缩reduce7. 重排scan/sortDATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)1. 映射Map 输入输入关系:一一对应(one-to-one) 例子:每个元素倍数扩大,y[i]=3*x[i]DATAGURU专业数据分析社区第一版 讲师 罗韵 (WeChat:LaurenLuoYun)2.聚合gatter 输入输出关系:多对一(many-to-one) 例子:每相邻3个元素求平均,y[i]=(x[i-1]+x[i]+x[i+1])/3DATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)3.分散scatter输入输出关系:一对多(one-to-many)DATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)4.模板stencil:以固定的模式读取相邻的内存数值输入输出关系:serveral-to-oneDATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)5.转置transpose输入输出关系:一对一(one-to-one)DATAGURU专业数据分析社区 第一版 讲师 罗韵 (WeChat:LaurenLuoYun)∙ 6.压缩reduce∙输入输出关系:多对一(all-to-one)∙7.重排scan/sort∙输入输出关系:多对多(all-to-all)1 2 3 41 3 6 10ADD2.GPU硬件模式2.1 GPU,SM(流处理器),Kernel(核),thread block(线程块),线程线程块Kernel核: 可以理解为C/C++中的一个函数functionThread Blocks: group of thread blocks to solve a functionThread Block: a group of threads that cooperate to solve a (sub)problem线程块GPU∙SM(stream multiprocessor): 流处理器∙GPU:每个GPU有若干个SM,最少有1个,目前16个算大的,每个SM并行而独立运行simple processormemoryGPU3.CUDA编程模型3.1 CUDA编程模型的优点和缺点3.2 CUDA编程编程模型的一些原则3.3 CUDA内存模型3.4 同步性synchronisation和屏障barrier 3.5 编程模型3.1CUDA编程的优点和后果∙CUDA最大的特点:对线程块将在何处、何时运行不作保证。
- 1、下载文档前请自行甄别文档内容的完整性,平台不提供额外的编辑、内容补充、找答案等附加服务。
- 2、"仅部分预览"的文档,不可在线预览部分如存在完整性等问题,可反馈申请退款(可完整预览的文档不适用该条件!)。
- 3、如文档侵犯您的权益,请联系客服反馈,我们会尽快为您处理(人工客服工作时间:9:00-18:30)。
CUDA并行架构及编程摘要CUDA是一种由NVIDIA推出的并行计算架构,非常适合大规模数据密集型计算。
CUDA使GPU的超高计算性能在数据处理和并行计算等通用计算领域发挥优势,本文讨论了CUDA的计算架构和基于GPU的CUDA C编程语言,CUDA使GPU流处理器阵列的性能得到充分发挥。
极大地提高了并行计算程序的效率。
关键词并行计算,GPU通用计算,CUDAAbstract CUDA is a parallel computing architecture introduced by NVIDIA , it mainly used for large scale data-intensive computing. CUDA makes GPU a high performance in parallel computing ,data processing and other general computing. this paper discusses the CUDA computing architecture and CUDA C programming language based on GPU, CUDA makes GPU stream processor arrays fully used and Greatly improved the efficiency of parallel computing program.Keywords parallel computing ,GPU general purpose computation,CUDA1 引言并行计算是指同时使用多种计算资源解决计算问题的过程。
并行计算科学中主要研究的是空间上的并行问题。
从程序和算法设计的角度来看,并行计算又可分为数据并行和任务并行。
一般来说,GPU更注重于数据并行计算,主要是将一个大任务化解成相同的各个子任务。
早期的GPU研究是通过可编程计算单元为屏幕上的每个像素计算出一个颜色值即渲染问题。
自CUDA C出现后,基于GPU的通用计算已经成为一个新的研究领域。
通常,像素着色器对各种颜色值进行合成并计算出最终的颜色值。
实际上,输入值可以为任意数据,这样不一定非要使用GPU来处理图形,还可以实现某些通用计算。
由于GPU有着很高的计算吞吐量,从而给大规模的数据计算应用提供了一种比CPU更加强大的计算能力。
CUDA是一种由NVIDIA推出的并行计算架构,该架构使GPU能够解决复杂的计算问题。
它包含了CUDA指令集架构以及GPU内部的并行计算引擎[1]。
随着显卡的发展,GPU越来越强大,在计算上已经超越了通用的CPU。
如此强大的芯片如果只是作为显卡会造成计算能力的浪费,因此NVIDIA推出CUDA,让显卡可以用于图像渲染以外的目的。
CUDA的GPU编程语言基于标准的C语言,通过在标准C语言的基础上增加一小部分关键字,任何有C语言基础的用户都很容易地开发CUDA的应用程序。
数以千计的软件开发人员正在使用免费的CUDA软件开发工具来解决各种专业中的问题[2]。
这些解决方案涵盖了石油天然气勘探、产品设计、医学成像以及科学研究等领域。
2 CUDA架构CUDA程序架构分为两部分:主机和设备。
一般而言,主机指的是CPU及其内存,设备指的是GPU[1]。
在CUDA程序架构中,主程序由CPU来执行,而当遇到数据并行处理的部分,CUDA就会将程序编译成GPU能执行的程序,并传送到GPU。
这种函数在CUDA里叫做核函数。
在GPU中要执行的线程,根据最有效的数据共享来创建线程块,其类型不止一维。
在同一个块里的线程,使用同一个共享内存。
每个线程由线程ID标识。
这是线程块中的线程号。
为了帮助基于线程ID的复杂寻址,还可以将线程块指定为任意大小的一维、二维或高维线程阵列,并使用多个索引分量来标识每个线程[2]。
线程块的大小是也有限制的,所以不能把所有的线程都放到同一个块里。
可以用同样维度和大小的块来组成一个网格做批处理。
执行内核的线程被组织成线程块。
而线程块又组成了线程格,如图1。
线程块内一起协作的线程通过一些快速的共享内存有效地共享数据并同步执行,以协调内存访问。
用户可以在核函数中指定同步点。
线程块中的线程在到达此同步点时挂起。
图1 线程被组织为线程格与线程块图2 SM单元的组成在GPU 内部,SM 代表流处理器,即计算核心。
每个SM 中又包含8 个标量流处理器(SP)以及少量的其他计算单元[3],如图2。
实际上SP只是执行单元,并不是完整的处理核心。
处理核心必须包含取指、解码、分发逻辑和执行单元。
隶属同一SM 的8 个SP 共用同一套取指和发射单元,也共用一块共享存储器。
一个线程块必须被分配到一个SM 中,但是一个SM 中同一时刻可以有多个活动线程块等待执行。
这可以更好地利用执行单元的资源,当一个线程块进行同步或者访问显存等高延迟操作时,另一个线程块就可以占用GPU 执行资源。
每个线程有一个私有的本地存储器,即每个标量处理器私有的缓存。
每个线程块有一个共享存储器,块内所有的线程都可以访问。
每个线程网格都可以访问设备内存,即GPU片上的全局存储器。
每个线程均可以读写全局存储器。
本地存储器用于储存设备代码中无法放进寄存器的变量[3]。
CUDA架构的所有这些功能都是为了使GPU不仅能执行传统的图形计算,还能高效的执行通用计算。
3 CUDA C编程3.1 CUDA并行编程运行在GPU上的CUDA并行计算函数称为核函数(kernel)。
一个kernel 函数并不是一个完整的程序,而是整个程序中的一个可以被并行执行的步骤[2]。
内核函数必须通过_global_ 函数类型限定符定义,如_global_ void kernel(void)。
并且只能在主机端代码中调用。
在调用时,必须声明内核函数的执行参数。
现考虑给一个N位矢量的每一位加上一个常数。
为完成此操作设计一个add函数,分别用C代码和CUDA C表示为:#define N 10 #define N 10void add(int *a,int b){ _global_void add(int*a,int b) int index=0; {while(index<N){ int index=blockIdx.x;a[index]=a[index]+b; if(index<N)a[index]=a[index]+b;index+=1; }}在CPU运行的程序中,通过while循环串行的给N位矢量每一位加上常数。
值得注意的是右边的CUDA核函数,在实际运行时,CUDA会产生许多在GPU上执行的线程,每一个线程都会去执行内核这个程序,虽然程序是同一份,但是因为设置了变量blockIdx,这是一个内置变量,在CUDA运行中已经预先定义了这个变量,变量的值是当前执行设备代码线程块的索引,从而取得不同的数据来进行并行计算。
这里考虑的实例是10位矢量的增量运算,当矢量的位数N远远大于10且进行的并不是简单的增量运算时,使用CUDA C就会有明显的差别。
int main(void){int host_a[N];int *device_a;cudaMalloc((void**)&device_a,N * sizeof(int)); //在设备上分配内存for (int i=0; i<N; i++)host_a[i] = i; //为数组赋值cudaMemcpy(device_a, host_a, N * sizeof(int),cudaMemcpyHostToDevice);add<<<N,1>>>(device_a, 5);cudaMemcpy( host_a,device_a,N*sizeof(int),cudaMemcpyDeviceToHost);cudaFree(device_a); //释放内存return 0;}在定义了核函数后,可以在主函数中调用它。
add<<<N,1>>>表示启动了N 个线程块,每个线程块中含一个线程。
CPU 串行代码完成的工作包括在核函数启动前进行数据准备和设备初始化的工作。
其中cudaMalloc()函数用来在GPU 设备上分配内存,第一个参数是一个指针,指向用于保存新分配内存地址的变量,第二个参数指明分配内存的大小,需要注意的是由这个内存分配函数得到的指针是不能在主机内存上使用的。
cudaMemcpy()通过设置参数在CPU和GPU之间传递数据。
主函数先完成数组初始化的工作,再将数据传入GPU,并行计算完成后再将结果传回CPU。
3.2 线程的同步与通信上一节矢量增值运算的例子,通过创建N 个含有一个线程的进程快给矢量的各位各加上一个常数。
上述运算并不需要考虑线程块内线程的同步与通信,但实际中的大多数应用都需要线程之间传递数据。
这也是CUDA 所提供的最重要的创新,它使得工作在GPU 上的线程可以协作解决问题,允许应用程序更加高效的执行[2]。
在同一个 block 中的线程通过同步函数_syncthreads ()来保证块内的线程同步即线程块中的所有线程都执行到同一位置。
为了在硬件上用很小的代价就能实现 _syncthreads() 函数,一个 block 中所有线程的数据都必须交由同一处理核心进行处理[3]。
所以,这导致每个线程块中的线程数量受到处理核心硬件资源的限制。
目前规定每个块里最多只能有 512 个线程[2]。
在同一个块中的线程通过共享存储器交换数据,为了保证线程块中的各个线程能够有效协作,访问共享存储器的延迟必须很小。
所以在 GPU 中,共享存储器与执行单元的物理距离必须很小。
在CUDA C 中同一个块中的线程通过共享变量来实现通信,使用__shared__来定义一个共享变量。
现利用一个大规模矩阵相乘的例子来说明线程块内线程同步与通信及线程之间的互斥。
矩阵A (MxN )与矩阵B (NxK )的乘法公式为[5]:1,...1,0;1,...1,0,10-=-=⨯=∑-=k j m i b a C ti n t it ij (1)为简单起见,这里仅考虑1xN 矩阵与Nx1矩阵之积,于是公式简化为:i Ni i NN b a a a a b b b ∑==⨯12121. (2)可以先让每个进程计算a i b i ,由于要计算累加和每个进程需要保存一个临时变量,每计算完一个积,索引平移blockDim.x * gridDim.x 个单位。
这样线程块内的每个线程都保存着若干个a i b i ,还需要在每个线程块内设置一个共享数组变量__shared__ int var[threadsperblock],这个数组的大小是线程块里线程的数量。