CV-CUDA: NVIDIA 官方出品高性能图像处理加速库

引言

随着短视频 APP、视频会议平台以及 VR/AR 等技术的发展,视频与图像已逐渐成为全球互联网流量的主要组成部分。包含我们平时接触到的这些视频图像,也有很多是被 AI 和计算机视觉(CV)算法处理并增强过的。然而,随着社交媒体和视频分享服务的快速增长,作为 AI 图像算法基础的视频图像处理部分,也早已成为计算流程中不可忽视的成本和瓶颈。

CV-CUDA 应用场景

我们先带大家简单回顾一下图像处理的一些常见的例子,以更好地理解 CV-CUDA(Computer Vision – Compute Unified Device Architecture)应用场景。

AI 算法图像背景模糊化

AI 背景模糊(CPU 前后处理方案)

图像背景模糊化通常会被应用于视频会议,美图修图等场景。在这些场景中,我们通常希望 AI 算法可以把主体之外的背景部分模糊化,这样可以保护用户隐私,美化图像等。图像背景模糊化的流程大体可以分为 3 个过程:

  • 前处理:通常包含了对图像做 Resize、Padding、Image2Tensor 等操作;
  • DNN 网络:可以是一些常见 segmentation network,比如 Unet 等;
  • 后处理:通常包括 Tensor2Mask、Crop、Resize、Denoise 等操作。

在传统的图像处理流程中,前处理和后处理部分通常都是使用 CPU 进行操作,这导致整个图像背景模糊化流程中,有 90% 的工作时间消耗在前后处理部分,因而成为了整个算法流水线的瓶颈。若能把前后处理妥善利用 GPU 加速,这将能大幅提升整体的计算性能。

AI背景模糊(GPU前后处理方案)

当我们把前后处理部分都放到 GPU 上后,我们就可以对整个 pipeline 进行端到端的 GPU 加速。经过测试,在单个 GP U上,相比于传统图像处理方式,把整个 pipeline 移植到 GPU 后,可以获得 20 倍以上的吞吐率提升。这无疑会大大的节省计算成本。

AI 算法图像分类

AI 图像分类

图像分类是最常见的 AI 图像算法之一,通常可以用于物体识别,以图搜图等场景,几乎是所有 AI 图像算法的基础。图像分类的 pipeline 大体可以分为 2 个部分:前处理部分和 DNN 部分。其中前处理部分,在训练和推理过程中最常见的 4 种操作包括:图片解码、Resize、Padding、Normalize。DNN 部分已经有了 GPU 的加速,而前处理部分通常都会使用 CPU 上的库函数进行处理。如果能够把前处理部分也移植到 GPU 上,那么一方面可以释放 CPU 资源,另一方面也可以进一步提升 GPU 利用率,从而可以对整个 pipeline 进行加速。

当前图像处理主流方案

针对前后处理部分,我们总结一下目前已有的一些主流应用方案:

(1)目前应用最广泛的图像处理库是 OpenCV。前后处理中的操作,往往是依赖于任务类型的,这导致了前后处理相关的操作种类繁多,数量庞杂。利用 OpenCV 确实可以实现大部分的图像处理操作,不过,这些操作大部分是用 CPU 实现的,缺少对应的 GPU 加速版本。而部分 CPU 操作虽然存在对应的 GPU 版本,但是这些 GPU 版本的实现可能存在一些问题,包括:

a. 部分算子的 CPU 和 GPU 结果无法对齐。在某些场景下,我们的算法团队使用了 CPU 的算子进行训练,而在推理阶段负责部署的同学考虑到性能问题,决定使用 GPU 算子,然而在测试过程中,发现 CPU 和 GPU 算子结果无法对齐,这可能会导致端到端的处理结果出现精度异常。

b. 部分算子 GPU 性能比 CPU 性能还弱。除了某些算法本身就不太适合 GPU 化之外,我们还发现部分 GPU 实现并不是很高效。比如,有些算子会在运行过程中,临时申请显存空间,导致增加了额外的耗时。

c. 在真实的部署 pipeline 中,可能有些算子只有 CPU 版本,而有些算子既有 CPU 版本又有 GPU 版本。如果 pipeline 里既有 CPU 算子又有 GPU 算子,那么会存在 GPU 和 CPU 同步,以及 H2D,D2H 内存拷贝问题,从而引入额外的耗时。

(2)使用 PyTorch 框架进行模型训练引入的 torchvision 图像处理库。在算法开发阶段,出于训练考虑,算法工程师可能会优先使用 torchvision 来完成图像处理操作。然而在模型部署阶段,负责模型部署的工程师通常会使用 C++ 作为开发语言,没有办法使用 torchvision 完成图像处理操作,那么可能会使用其他的 C++ 图像处理库,比如 OpenCV,然而这就涉及到,torchvision 和 OpenCV 结果对齐的问题。此外,与使用 OpenCV 的情况类似,torchvision 中的算子也有不支持 GPU 加速,或者 GPU 加速效果不佳的情况。

CV-CUDA 特点

如上所述,传统的图像预处理操作一般在 CPU 上进行,一方面会占用大量的 CPU 资源,使得 CPU和 GPU 的负载不均衡;另一方面由于基于 CPU 的图像加速库不支持 batch 操作,导致预处理的效率低下。为了解决当前主流的图像处理库所存在的一些问题,NVIDIA 和字节跳动的机器学习团队[1] 联合开发了基于 GPU 的图像处理加速库 CV-CUDA,并拥有以下特点:

Batch

支持 batch 操作,可以充分利用 GPU 高并发、高吞吐的并行加速特性,提升计算效率和吞吐率。

Variable Shape

支持同一 batch 中图片尺寸各不相同,保证了使用上的灵活性。此外,在对图片进行处理时,可以对每张图片指定不同的参数。例如调用 RotateVarShape 算子时,可以对 batch 中每张图片指定不同的旋转角度。

C/C++/Python 接口

在部署机器学习算法时需要对齐训练和推理流程。一般来说,训练时利用 python 进行快速验证,推理时利用 C++ 进行高性能部署,然而一些图像处理库仅支持 python,这给部署带来了极大的不便。如果在训练和推理采用不同的图像处理库,则又需要在推理端重新实现一遍逻辑,过程会非常繁琐。

CV-CUDA 提供了 C、C++ 和 Python 接口,可以同时服务于训练和推理场景。从训练迁移到推理场景时,也可免去繁琐的对齐流程,提高部署效率。

独立算子设计

CV-CUDA 作为基础图像处理库,采用了独立算子设计,不需要预先定义流水线。独立算子的设计具有更高的灵活性,使调试变得更加的容易,而且可以使其与其他的图像处理交互,或者将其集成在用户自己的图像处理上层框架中。

结果对齐 OpenCV

不同的图像处理库由于对一些算子的实现方式不一致导致计算结果难以对齐。例如常见的 Resize 操作,OpenCV、OpenCV-gpu 以及 torchvision 的实现方式都不一样,计算结果存在差异。因此如果在训练时用 OpenCV CPU 版本而推理时若要采用 GPU 版本或其他图像处理库,就会面临结果存在误差的问题。

在设计之初,我们考虑到当前图像处理库中,很多用户习惯使用 OpenCV 的 CPU 版本,因此在设计算子时,不管是函数参数还是图像处理结果上,尽可能对齐 OpenCV CPU 版本的算子。当用户从 OpenCV 迁移到 CV-CUDA 时,只需做少许改动便可使用,且图片处理结果和 OpenCV 一致,不需要重新训练模型。

易用性

CV-CUDA 提供了 Image、ImageBatchVarShape 等结构体,方便用户的使用。同时还提供了 Allocator 类,用户可以自定义显存分配策略(例如用户可以设计显存池分配策略来提高显存分配速度),方便上层框架集成和管理资源。目前 CV-CUDA 提供了 PyTorch、OpenCV 和 Pillow 的数据转化接口,方便用户进行算子替换和进行不同图像库之间的混用。

针对不同 GPU 架构的性能高度优化

CV-CUDA 可以支持 Volta、Turing、Ampere 等 GPU 架构,并针对不同架构 GPU 的特点,在 CUDA kernel 层面进行了性能上的高度优化,可在云服务场景中规模化部署。

算子数量及其性能

支持的算子

CV-CUDA 目前提供了 25 种算子,包括常用的 CvtColor、Resize、Crop、PadStack、Normalize 等。当前算子能够覆盖大部分应用场景,后续将会支持更多算子。

CV-CUDA 支持算子清单

性能对比

在本文开头的背景模糊算法里,采用 CV-CUDA 替代 OpenCV 和 TorchVision 的前后处理后,整个推理流程的吞吐率提升了 20 倍以[2] 上。下图展示了在同一个计算节点上(2x Intel Xeon Platinum 8168 CPUs , 1x NVIDIA A100 GPU),以 30fps 的帧率处理 1080p 视频,采用不同的 CV 库所能支持的最大的并行流数。测试采用了 4 个进程,每个进程 batchSize 为 64。

AI背景模糊(2x Intel Xeon Platinum 8168 CPUs , 1x NVIDIA A100 GPU)

其中涉及到的前处理操作有:

Resize (Downscale)、Padding、Convert Data Type、Normalize及Image to Tensor

涉及到的后处理操作有:Tensor to Mask、Convert Data Type、Crop、Resize (Upscale)、Bilateral Filter (Denoise)、Gaussian Blur 及Composite

对于单个算子的性能,我们也做了性能测试,下图的测试场景选用的图片大小为 480*360,CPU选择为 Intel® Core™ i9-7900X,BatchSize 大小为 1,进程数为 1。

CV-CUDA性能对比(CPU: Intel(R) Core(TM) i9-7900X CPU @ 3.30GHz)

性能优化

为了使 CV-CUDA 能够更加高效的运行在 GPU 上,我们采取了一系列的优化手段。

  • kernel 融合

采用了大量的 kernel 融合策略,减少了 kernel launch 和 global memory 的访问时间。

  • 访存优化

采用了合并访存,向量化读写,shared memory 等策略,提高了数据读写的效率。

  • 异步处理

CV-CUDA 中所有算子均采用异步处理的方式,可以减少同步带来的等待耗时。

  • 高效计算

采用了 fast math、warp/block reduce、table lookup 等优化手段,可以有效提升计算效率。

  • 预分配显存

CV-CUDA 采用了预分配显存策略,并且提供了 Allocator 类,帮助使用者自定义显存分配策略或者可采取默认的显存分配策略。算子所需要的 buffer 和图片显存会在初始化阶段分配好,而在执行阶段不会再进行耗时的显存分配操作。

框架和API

CV-CUDA 整体架构

整个 CV-CUDA 库包含了以下几个组成部分

a. CV-CUDA 核心模块

核心模块包含了 C/C++ 和 Python API、NVCV 模块,Operator 算子模块以及 CV-CUDA Tools。

b. CV-CUDA Interop 模块

这个模块包含了和其他图像处理库以及推理框架的交互接口,目前支持 OpenCV、Pytorch 和 Pillow,后续将陆续加入其他图像处理库的交互接口。

c. CV-CUDA Tools/Tests

包含一些单元测试模块和工具函数

CV-CUDA 整体架构

CV-CUDA 核心模块

CV-CUDA 核心模块

核心模块包含以下几个部分

a. NVCV 核心支持库

包含 Image 抽象类、Memory workspace 类、Allocator 和 Batch/VarShapeBatch 类

b. Operators CV 算子

包含了各种独立算子(resize、filter等等)

c. CV-CUDA Tools

包含了用于开发算子所需的各种功能函数

API

ImageBatchVarShape

Resize

代码实战

以图片分类为例,以下代码展示了如何利用 CV-CUDA 对图片进行预处理,以及如何和 Pytorch 进行交互。

Pipeline

import torch
from torchvision import models
from torchvision.io.image import read_file, decode_jpeg
import numpy as np

# Import CV-CUDA module
import nvcv

"""
Image Classification python sample

The image classification sample uses Resnet50 based model trained on Imagenet
The sample app pipeline includes preprocessing, inference and post process stages
which takes as input a batch of images and returns the TopN classification results
of each image.

This sample gives an overview of the interoperability of pytorch with CVCUDA
tensors and operators
"""

# Set the image and labels file
filename = "./assets/tabby_tiger_cat.jpg"
labelsfile = "./models/imagenet-classes.txt"

# Read the input imagea file
data = read_file(filename)  # raw data is on CPU

# NvJpeg can be used to decode the image
# to the necessary color format on the device
inputImage = decode_jpeg(data, device="cuda")  # decoded image in on GPU
imageWidth = inputImage.shape[2]
imageHeight = inputImage.shape[1]

# decode_jpeg is currently limited to batchSize 1
# and Planar format (CHW)

# A torch tensor can be wrapped into a CV-CUDA Object using the "as_tensor"
# function in the specified layout. The datatype and dimensions are derived
# directly from the torch tensor.
nvcvInputTensor = nvcv.as_tensor(inputImage, "CHW")

# The input image is now ready to be used

# The Reformat operator can be used to convert CHW format to NHWC
# for the rest of the preprocessing operations

nvcvInterleavedTensor = nvcvInputTensor.reformat("NHWC")

"""
Preprocessing includes the following sequence of operations.
Resize -> DataType Convert(U8->F32) -> Normalize(Apply mean and std deviation)
-> Interleaved to Planar
"""

# Model settings
layerHeight = 224
layerWidth = 224
batchSize = 1

# Resize
# Resize to the input network dimensions
nvcvResizeTensor = nvcvInterleavedTensor.resize(
    (batchSize, layerWidth, layerHeight, 3), nvcv.Interp.CUBIC
)

# Convert to the data type and range of values needed by the input layer
# i.e uint8->float. A Scale is applied to normalize the values in the range 0-1
nvcvConvertTensor = nvcvResizeTensor.convertto(np.float32, scale=1 / 255)

"""
The input to the network needs to be normalized based on the mean and
std deviation value to standardize the input data.
"""

# Create a torch tensor to store the mean and standard deviation values for R,G,B
scale = [0.485, 0.456, 0.406]
std = [0.229, 0.224, 0.225]
scaleTensor = torch.Tensor(scale)
stdTensor = torch.Tensor(std)

# Reshape the the number of channels. The R,G,B values scale and offset will be
# applied to every color plane respectively across the batch
scaleTensor = torch.reshape(scaleTensor, (1, 1, 1, 3)).cuda()
stdTensor = torch.reshape(stdTensor, (1, 1, 1, 3)).cuda()

# Wrap the torch tensor in a CV-CUDA Tensor
nvcvScaleTensor = nvcv.as_tensor(scaleTensor, "NHWC")
nvcvBaseTensor = nvcv.as_tensor(stdTensor, "NHWC")

# Apply the normalize operator and indicate the scale values are std deviation
# i.e scale = 1/stddev
nvcvNormTensor = nvcvConvertTensor.normalize(
    nvcvBaseTensor, nvcvScaleTensor, nvcv.NormalizeFlags.SCALE_IS_STDDEV
)

# The final stage in the preprocess pipeline includes converting the RGB buffer
# into a planar buffer
nvcvPreprocessedTensor = nvcvNormTensor.reformat("NCHW")

# Inference uses pytorch to run a resnet50 model on the preprocessed input and outputs
# the classification scores for 1000 classes
# Load Resnet model pretrained on Imagenet
resnet50 = models.resnet50(pretrained=True)
resnet50.to("cuda")
resnet50.eval()

# Run inference on the preprocessed input
torchPreprocessedTensor = torch.as_tensor(nvcvPreprocessedTensor.cuda(), device="cuda")
inferOutput = resnet50(torchPreprocessedTensor)

"""
Postprocessing function normalizes the classification score from the network and sorts
the scores to get the TopN classification scores.
"""
# top results to print out
topN = 5

# Read and parse the classes
with open(labelsfile, "r") as f:
    classes = [line.strip() for line in f.readlines()]

# Apply softmax to Normalize scores between 0-1
scores = torch.nn.functional.softmax(inferOutput, dim=1)[0]

# Sort output scores in descending order
_, indices = torch.sort(inferOutput, descending=True)

# Display Top N Results
[
    print("Class : ", classes[idx], " Score : ", scores[idx].item())
    for idx in indices[0][:topN]
]

写在最后

如果您也对人工智能和计算机视觉全栈领域感兴趣,强烈推荐您关注有料、有趣、有爱的公众号『CVHub』,每日为大家带来精品原创、多领域、有深度的前沿科技论文解读及工业成熟解决方案!欢迎添加小编微信号:cv_huber,一起探讨更多有趣的话题!

本文正在参加 人工智能创作者扶持计划

  • 0
    点赞
  • 10
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
CUDA是一种并行计算平台和编程模,可以利用GPU的并行计算能力来加速像处理等任务。下面是使用CUDA加速图像处理的步骤: 1. 首先,需要安装CUDA工具包和相应的驱动程序。CUDA工具包包含了编译器、和具,用于开发和运行CUDA程序。 2. 在像处理代码中,需要包含CUDA的头文件,并使用`global__`关键字定义一个CUDA核函数。CUDA核函数将在GPU上并行执行。 3. 在CUDA核函数中,可以使用CUDA提供线程和块的概念来管理并行计算。可以`threadIdx``blockIdx`等变量来获取当前程和块的引。 4. 在CUDA核函数中,可以使用CUDA提供并行计算指令和函数来进行图像处理操作。例如,可以使用CUDA提供的并行循环来遍历图像像素,并使用GPU的并行计算能力来加速像处理算法。 5. 在主机代码中,需要调用CUDA核函数来启动GPU上的并行计算。可以使用`cudaMalloc`函数来分配GPU内存,使用`cudaMemcpy`函数来在主机和设备之间传输数据,使用`cudaFree`函数来释放GPU内存。 下面是一个使用CUDA加速图像处理的示例代码: ```cpp #include <cuda_runtime.h> #include <device_launch_parameters.h> __global__ void imageProcessingKernel(unsigned char* image, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { // 图像处理操作 // ... } } void imageProcessing(unsigned char* image, int width, int height) { // 在设备上分配内存 unsigned char* dev_image; cudaMalloc((void**)&dev_image, width * height * sizeof(unsigned char)); // 将图像数据从主机内存复制到设备内存 cudaMemcpy(dev_image, image, width * height * sizeof(unsigned char), cudaMemcpyHostToDevice); // 定义线程块和线程数量 dim3 blockSize(16, 16); dim3 gridSize((width + blockSize.x - 1) / blockSize.x, (height + blockSize.y - 1) / blockSize.y); // 调用CUDA核函数启动并行计算 imageProcessingKernel<<<gridSize, blockSize>>>(dev_image, width, height); // 将处理后的图像数据从设备内存复制到主机内存 cudaMemcpy(image, dev_image, width * height * sizeof(unsigned char), cudaMemcpyDeviceToHost); // 释放设备内存 cudaFree(dev_image); } ``` 请注意,上述代码只是一个示例,具体的图像处理操作需要根据实际需求进行编写。另外,还需要在编译时链接CUDA

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值