CUDA编程——zero copy

零复制

  zero copy(零复制)是一种特殊形式的内存映射,它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问(direct memory access,DMA)方式来访问主机的锁页内存。
  


锁页主机内存

  现代操作系统都支持虚拟内存,操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存(硬盘等大容量存储)上,从而腾出空间存放将要调入内存的信息。这样,系统好像为用户提供了一个比实际内存大得多的存储器,称为虚拟存储器。
  锁页就是将内存页面标记为不可被操作系统换出的内存。所以设备驱动程序给这些外设编程时,可以使用页面的物理地址直接访问内存(DMA),从而避免从外存到内存的复制操作。CPU 仍然可以访问上述锁页内存,但是此内存是不能移动或换页到磁盘上的。CUDA 中把锁页内存称为pinned host memory 或者page-locked host memory。


锁页主机内存的优势

  使用锁页内存(page-locked host memory)有一些优势:

  • 锁页内存和GPU内存之间的拷贝可以和内核程序同时执行,也就是异步并发执行。
  • 在一些设备上锁页内存的地址可以从主机地址空间映射到CUDA 地址空间,免去了拷贝开销。
  • 在拥有前线总端的系统上,如果主机内存被分配为锁页内存,主机内存和GPU 内存带宽可以达到更高,如果主机内存被分配为Write-Combining Memory,带宽会进一步提升。

然而锁页主机存储器是稀缺资源,所以锁页内存分配得多的话,分配会失败。另外由于减少了系统可分页的物理存储器数量,分配太多的分页锁定内存会降低系统的整体性能


使用锁页主机内存

  在GPU 上分配的内存默认都是锁页内存,这只是因为GPU 不支持将内存交换到磁盘上。在主机上分配的内存默认都是可分页,如果需要分配锁页内存,则需要使用cudaMallocHost() 或者cudaHostAlloc()。释放时需要使用cudaFreeHost() 释放这一块内存。调用常规的C函数释放,可能会崩溃或者出现一些不常见的错误。也可以通过函数cudaHostRegister() 把可分页内存标记为锁页内存。

__host__ ​cudaError_t cudaMallocHost ( void** ptr, size_t size )

__host__ ​cudaError_t cudaHostAlloc ( void** pHost, size_t size, unsigned int  flags )

__host__ ​cudaError_t cudaFreeHost ( void* ptr )

cudaHostAlloc() 多了一个可选形参flags ,功能更强大。flags 的值可以取如下值。

#define cudaHostAllocDefault 0x00
Default page-locked allocation flag

#define cudaHostAllocMapped 0x02
Map allocation into device space

#define cudaHostAllocPortable 0x01
Pinned memory accessible by all CUDA contexts

#define cudaHostAllocWriteCombined 0x04
Write-combined memory

cudaHostRegister() 函数用于把已经的存在的可分页内存注册为分页锁定的。

__host__ ​cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int  flags )

flags 是一个可选形参,可以取如下值。

#define cudaHostRegisterDefault 0x00
Default host memory registration flag

#define cudaHostRegisterIoMemory 0x04
Memory-mapped I/O space

#define cudaHostRegisterMapped 0x02
Map registered memory into device space

#define cudaHostRegisterPortable 0x01
Pinned memory accessible by all CUDA contexts

下面分别介绍这些flags 的作用。

Portable Memory

  一块锁页内存可被系统中的所有设备使用(一个系统中有多个CUDA设备时)。 启用这个特性需要在调用cudaHostAlloc() 时使用cudaHostAllocPortable 选项,或者在调用cudaHostRegister() 使用cudaHostRegisterPortable 选项。
  

Write-Combining Memory

  默认情况下,锁页主机存储是可缓存的。可以在调用cudaHostAlloc() 时传入cudaHostAllocWriteCombined 标签使其被分配为写结合的(Write-Combining Memory)。写结合存储不使用L1 和L2 cache,所以程序的其它部分就有更多的缓存可用。此外,写结合内存通过PCI-E 传输数据时不会被监视(snoop),这能够获得高达40%的传输加速。 从主机读取写结合存储非常慢(因为没有使用L1、L2cache),所以写结合存储应当只用于那些主机只写的存储。
  

Mapped Memory

  一块锁页内存可以在调用cudaHostAlloc() 分配时传入cudaHostAllocMapped 标签或者在使用cudaHostRegister() 注册时使用cudaHostRegisterMapped 标签,把锁页内存地址映射到设备地址空间。这样,这块存储会有两个地址:一个是从cudaHostAlloc() 或malloc() 返回的在主机内存地址空间上;另一个在设备存储器上,可以通过cudaHostGetDevicePointer() 取得。内核函数可以使用这个指针访问这块存储。 cudaHostAlloc() 返回的地址指针一个的例外情况是,主机和设备使用统一地址空间(Unified Virtual Address Space)。
内核直接存取主机内存有很多优势:

  • 无需在设备上分配内存,也无需在主机内存和设备内存之间拷贝数据。数据传输是在内核需要的时候隐式进行的。
  • 无须使用流(cuda stream)就可以并发数据传输和内核执行;数据传输和内核执行自动并发执行。

因为映射的锁页主机内存是主机和设备之间共享的,所以在使用cuda stream 或者cuda event 时必须对内存读写同步;避免潜在的写后读,读后写或者写后写等多线程同步问题。
  为了能够对任何映射的锁页主机内存解引用设备指针,必须在调用任何cuda 运行时函数前调用cudaSetDeviceFlags(),并传入cudaDeviceMapHost 标签。否则,cudaHostGetDevicePointer() 将会返回错误。
  如果设备不支持被映射分页锁定存储,cudaHostGetDevicePointer() 将会返回错误。程序员可以检查canMapHostMemory 属性,如果设备支持映射锁页主机内存,将会返回1。

注意:使用映射锁页主机内存看,原子操作将不再保证原子性。cudaHostRegisterIoMemory 是cudaHostRegister() 特有的选项,可以把主机内存映射到IO 地址空间。


参考文献

[1]https://en.wikipedia.org/wiki/CUDA_Pinned_memory
[2] Cook, Shane (2013). CUDA Programming: A Developer’s Guide to Parallel Computing with GPUs (1st ed.). Morgan Kaufmann Publishers Inc. pp. 334–335. ISBN 9780124159334.
  

### 图像修复技术 RePaint 的实现方法 RePaint 是一种基于去噪扩散概率模型(Denoising Diffusion Probabilistic Models, DDPM)的图像修复方法[^1]。它通过条件生成的方式,在已知区域的基础上完成缺失部分的补全,而无需依赖特定分布的 mask,从而提高了泛化能力。 #### 技术背景 RePaint 利用了扩散模型的核心思想——通过对噪声逐步建模并逆向采样来生成高质量的内容。相比传统的方法,其优势在于能够处理标准和极端掩码下的多种场景,并在人脸和其他复杂图像数据集上表现出 SOTA 性能。 以下是关于如何实现或使用 RePaint 技术的一些指导: --- ### 工具与环境准备 为了实现 RePaint 方法,需要以下基础工具和技术栈: - **编程语言**: Python - **深度学习框架**: PyTorch 或 TensorFlow(PyTorch 更常见于扩散模型的研究) - **硬件支持**: GPU 加速推荐 NVIDIA CUDA 支持的显卡 - **开源库**: `torch`, `numpy`, `matplotlib` 等常用科学计算包 安装必要的依赖项可以通过 pip 完成: ```bash pip install torch torchvision numpy matplotlib scikit-image ``` --- ### 数据预处理 图像修复任务的数据通常由两部分组成:原始图像及其对应的掩码。对于 RePaint 而言,由于不严格依赖特定分布的 mask,因此可以灵活设计掩码形式[^1]。 1. 准备训练数据集(如 CelebA-HQ、Places2 等),并将每张图片分割为输入图像和对应掩码。 2. 掩码可以随机生成,也可以手动标注,具体取决于应用场景需求[^2]。 代码示例(生成随机矩形掩码): ```python import numpy as np import cv2 def generate_random_mask(image_shape, min_size=30, max_size=80): height, width = image_shape[:2] top = np.random.randint(0, height - min_size) left = np.random.randint(0, width - min_size) size = np.random.randint(min_size, max_size) mask = np.zeros((height, width), dtype=np.uint8) mask[top:top+size, left:left+size] = 255 return mask # 测试函数 image = np.ones((256, 256, 3)) * 255 # 白色背景 mask = generate_random_mask(image.shape) masked_image = image.copy() masked_image[mask == 255] = 0 # 将掩码区域置零 cv2.imshow('Masked Image', masked_image) cv2.waitKey(0) cv2.destroyAllWindows() ``` --- ### 模型架构 RePaint 基于扩散模型的思想构建,主要包括以下几个模块: 1. **前向过程 (Forward Process)**: 对输入图像逐渐增加高斯噪声,直到变为纯噪声。 2. **反向过程 (Reverse Process)**: 学习从噪声中恢复原图的过程。 3. **条件约束**: 在反向过程中引入已知区域的信息作为条件。 核心网络结构通常是 U-Net 变体,用于估计当前时间步的噪声残差。 代码片段(U-Net 骨干定义): ```python import torch.nn as nn class UNet(nn.Module): def __init__(self, in_channels=3, out_channels=3): super().__init__() self.encoder = nn.Sequential( nn.Conv2d(in_channels, 64, kernel_size=3, padding=1), nn.ReLU(), nn.Conv2d(64, 128, kernel_size=3, stride=2, padding=1), nn.ReLU() ) self.decoder = nn.Sequential( nn.ConvTranspose2d(128, 64, kernel_size=4, stride=2, padding=1), nn.ReLU(), nn.Conv2d(64, out_channels, kernel_size=3, padding=1), nn.Tanh() ) def forward(self, x): encoded = self.encoder(x) decoded = self.decoder(encoded) return decoded ``` --- ### 训练流程 1. 初始化扩散模型参数,包括时间步数 \( T \) 和噪声调度器。 2. 构造损失函数,通常采用均方误差(MSE)衡量预测噪声与真实噪声之间的差异。 3. 进行端到端优化,更新模型权重以最小化损失。 伪代码描述如下: ```python from tqdm import tqdm model = UNet().cuda() # 移动至 GPU optimizer = torch.optim.Adam(model.parameters(), lr=1e-4) criterion = nn.MSELoss() for epoch in range(num_epochs): model.train() progress_bar = tqdm(train_loader, desc=f'Epoch {epoch}') for batch_images, masks in progress_bar: optimizer.zero_grad() # 添加噪声 timesteps = torch.randint(0, T, (batch_images.size(0),)).long().cuda() noise = torch.randn_like(batch_images).cuda() noisy_images = q_sample(batch_images.cuda(), timesteps, noise=noise) # 条件输入 condition = batch_images.clone() condition[masks.bool()] = 0 # 模型推理 predicted_noise = model(noisy_images, timesteps, condition) # 计算损失 loss = criterion(predicted_noise, noise) loss.backward() optimizer.step() progress_bar.set_postfix({'Loss': loss.item()}) ``` 其中,\( q\_sample \) 表示前向扩散过程的具体实现。 --- ### 应用实例 假设我们已经完成了模型训练,下面是如何应用该模型进行实际图像修复的一个简单例子: ```python import torch def inpaint_with_repaint(model, corrupted_image, mask, device='cuda'): """ 使用 RePaint 模型对损坏图像进行修复。 :param model: 训练好的扩散模型 :param corrupted_image: 输入的带遮罩的图像 (Tensor 形状 [C,H,W]) :param mask: 缺失区域的二值掩码 (形状同上) :return: 修复后的完整图像 """ model.eval() with torch.no_grad(): input_tensor = corrupted_image.unsqueeze(0).to(device) mask_tensor = mask.unsqueeze(0).to(device) condition = input_tensor.clone() condition[mask_tensor.bool()] = 0 sampled_image = diffusion_reverse_process(model, condition, num_steps=T) return sampled_image.squeeze().cpu().clamp(-1, 1).permute(1, 2, 0).numpy() / 2 + 0.5 ``` 调用此函数即可获得修复结果。 --- ### 结论 综上所述,RePaint 提供了一种强大的无监督图像修复方案,尤其适合处理复杂的缺失模式[^1]。其实现涉及扩散模型的基础理论以及高效的训练策略。希望以上内容可以帮助您快速入门并实践这一先进技术。 ---
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值