热门问题
时间线
聊天
视角
CUDA
NVIDIA所推出的一種軟硬體整合技術 来自维基百科，自由的百科全书
Remove ads
CUDA（Compute Unified Devices Architecture，统一计算架构[1]）是由英伟达（NVIDIA）所推出的一种软硬体整合技术，是该公司对于GPGPU的正式名称。透过这个技术，使用者可利用NVIDIA的GPU进行图像处理之外的运算，亦是首次可以利用GPU作为C-编译器的开发环境。CUDA 开发套件（CUDA Toolkit ）只能将自家的CUDA C-语言（对OpenCL只有链接的功能[2]），也就是执行于GPU的部分编译成PTX中间语言或是特定NVIDIA GPU架构的机器码（NVIDIA 官方称为 "device code"）；而执行于中央处理器部分的C / C++程式码（NVIDIA 官方称为 "host code"）仍依赖于外部的编译器，如Microsoft Windows下需要Microsoft Visual Studio；Linux下则主要依赖于GCC。[3][4][5]
 | 此条目需要更新。 (2021年3月6日) |
Remove ads
简介
.PNG)
1.将主记忆体的资料传入GPU记忆体
2. CPU指令驱动GPU
3. GPU平行运算
4. 将运算结果自GPU记忆体传回主记忆体
GPU不仅用于进行图形渲染，而且用于物理运算（物理效果如碎片、烟、火、流体）如PhysX和Bullet。进一步的，GPU可以用在计算生物学与密码学等领域的非图形应用上。在NVIDIA收购AGEIA后，NVIDIA取得相关的物理加速技术，即是PhysX物理引擎。配合CUDA技术，显示卡可以模拟成一颗PhysX物理加速晶片[6]。目前，全系列的GeForce 8显示核心都支援CUDA。而NVIDIA亦不会再推出任何的物理加速卡，显示卡将会取代相关产品。
而使用CUDA技术，GPU可以用来进行通用处理（不仅仅是图形）；这种方法被称为GPGPU。与CPU不同的是，GPU以较慢速度并行大量执行绪，而非快速执行单一执行绪。以GeForce 8800 GTX为例，其核心拥有128个内处理器。利用CUDA技术，就可以将那些内处理器做为线程处理器，以解决数据密集的计算。而各个内处理器能够交换、同步和共享数据。GeForce 8800 GTX的运算能力可达到520GFlops，如果建设SLI系统，就可以达到1TFlops。[7]
目前，已有软体厂商利用CUDA技术，研发出Adobe Premiere Pro的插件。通过插件，使用者就可以利用显示核心去加速H.264/MPEG-4 AVC的编码速度。速度是单纯利用CPU作软体加速的7倍左右。
虽然CUDA底层是以C/C++为主，并以使用“NVCC”——NVIDIA基于LLVM的C/C++编译器介面来进行编译，但工程师也可以使用编译器指令（如OpenACC）以及多种程式设计语言扩展对CUDA平台进行操作。如Fortran工程师可以使用“CUDA Fortran”，或PGI公司的PGI CUDA Fortran 编译器进行编译。除了此之外CUDA平台还支援其它计算介面，如Khronos Group的OpenCL，Microsoft的DirectCompute，以及C++AMP。也可以透过其他语言如 Python，Perl，Java，Ruby，Lua，Haskell，MATLAB，IDL及Mathematica 的介面间接调用CUDA。
CUDA最初的CUDA软体发展包（SDK）于2007年2月15日公布，同时支持Microsoft Windows和Linux。而后在第二版中加入对Mac OS X的支持(但于CUDA Toolkit 10.2起放弃对macOS的支援)，取代2008年2月14日发布的测试版。所有G8x系列及以后的NVIDIA GPUs皆支援CUDA技术，包括GeForce，Quadro和Tesla系列。CUDA与大多数标准作业系统相容。Nvidia声明：根据二进位相容性，基于G8x系列开发的程式无需修改即可在未来所有的Nvidia显卡上运行。
Remove ads
优点
在GPUs（GPGPU）上使用图形APIs进行传统通用计算，CUDA技术有下列几个优点：[8]
- 分散读取——代码可以从记忆体的任意位址读取
- 统一虚拟记忆体（Unified Memory, 从 CUDA 6.0 开始）—— 将所有 CPU 和 GPU 的内存置于统一管理的虚拟记忆体空间下。
- 共用记忆体（Global Memory）—— 存取快速的区域，使之在多个执行绪间共用，有效频宽比纹理记忆体（Texture Memory）更大。
- 与GPU之间更快的下载与回读
- 全面支持整型与位操作，包括整型纹理查找
限制
- 不支持纹理渲染（CUDA 3.2及以后版本通过在CUDA阵列中引入“表面写操作”——底层的不透明资料结构——来进行处理）
- 受系统主线的频宽和延迟的影响，主机与设备记忆体之间资料复制可能会导致性能下降（通过GPU的DMA引擎处理，非同步记忆体传输可在一定范围内缓解此现象）
- 当执行绪总数为数千时，执行绪应按至少32个一组来运行才能获得最佳效果。如果每组中的32个执行绪使用相同的执行路径，则程式分支不会显著影响效果；在处理本质上不同的任务时，单指令流多数据流执行模型将成为一个瓶颈（如在光线追踪演算法中遍历一个空间分割的资料结构）
- 只有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舍入模式（舍位与舍入到最近的偶数），这些在每条指令的基础上指定，而非控制字码；除法/平方根的精度比单精确度略低。
显卡的受支持情况
'*' – 仅限 OEM 产品
Remove ads
应用
利用CUDA技术，配合适当的软体（例如MediaCoder[21]、Freemake Video Converter），就可以利用显示核心进行高清视频编码加速。视频解码方面，同样可以利用CUDA技术实现。此前，NVIDIA的显示核心本身已集成PureVideo单元。可是，实现相关加速功能的一个微软API－DXVA，偶尔会有加速失效问题。所以利用CoreAVC配合CUDA，变相在显示核心上实现软体解码，解决兼容性问题[22]。另外，配合适当的引擎，显示核心就可以计算光线跟踪。NVIDIA就放出了自家的Optix实时光线跟踪引擎，透过CUDA技术利用GPU计算光线跟踪。[23]
支援的产品
示例
下面将示范以最底层的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 做一次加法
}
Remove ads
下列的范例是以相较于 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);
}
以下以pycuda （页面存档备份，存于互联网档案馆），间接调用 CUDA 做列向量的阿达玛乘积(也就是元素间乘积而非内积)
import pycuda.driver as drv #CUDA drivers
import pycuda.autoinit # 自動初始化CUDA
import numpy # 矩陣運算
# 讀取並編譯GPU執行的程式碼，以 CUDA-C 寫成
mod = drv.SourceModule("""
__global__ void multiply_them(float *c, float *a, float *b)
{
int i = threadIdx.x;
c[i] = a[i] * b[i];
}
""")
# 獲取GPU程式碼中的 multiply_them 函數
multiply_them = mod.get_function("multiply_them")
#生成兩個常態隨機分布的浮點數ndarray，shape 為(400, )
h_a = numpy.random.randn (400).astype(numpy.float32)
h_b = numpy.random.randn (400).astype(numpy.float32)
# 儲存結果的 h_c 列向量，其 shape 和向量 a 相同但值為零
h_c = numpy.zeros_like (a)
'''
執行GPU的函數:
注意這裡h_a, h_b 和 h_c 是在CPU記憶體的 python 變數
以下按照原來 GPU 程式碼的區域變數順序，指定哪些是從CPU傳入GPU (drv.In)；哪些是從GPU傳入CPU (drv.Out)，
簡單來說，自 GPU 傳出到 h_c； h_a 傳入 GPU；h_b 傳入 GPU。
(400,1,1) 代表一個block裡有的thread數量為 400 x 1 x 1
'''
multiply_them(
drv.Out(h_c), drv.In(h_a), drv.In(h_b),
block=(400,1,1))
#印出結果
print (d_c)
也可以用pycublas 间接调用 CUDA ，来计算矩阵乘法
import numpy
from pycublas import CUBLASMatrix
# 以 numpy 定義矩陣並傳入 CUBLASMatrix
A = CUBLASMatrix(numpy.mat([[1,2,3],[4,5,6]],numpy.float32))
B = CUBLASMatrix(numpy.mat([[2,3],[4,5],[6,7]],numpy.float32))
# 以 CUBLASMatrix 做矩陣乘法
C = A*B
# 將運算結果轉回 numpy 並印出
print(C.np_mat())
Remove ads
相关条目
参考文献
外部链接
Wikiwand - on
Seamless Wikipedia browsing. On steroids.
Remove ads