主要内容

GPU内存分配和最小化

离散和管理模式

GPU Coder™为您提供两种不同的内存分配(malloc)可在CUDA中使用的模式®编程模型,cudaMalloc而且cudaMallocManagedcudaMallocAPI适用于传统的CPU分离,GPU全局内存。cudaMallocManaged适用于统一内存

从程序员的角度来看,传统的计算机架构要求在CPU和GPU内存空间之间分配和共享数据。应用程序需要管理这两个内存空间之间的数据传输,这增加了复杂性。统一内存创建一个托管内存池,由CPU和GPU共享。CPU和GPU都可以通过一个指针访问托管内存。统一内存试图通过将数据迁移到需要它的设备来优化内存性能,同时对程序隐藏迁移细节。尽管统一内存简化了编程模型,但当GPU上写入的数据在CPU上被访问时,它需要设备同步调用。GPU Coder插入这些同步调用。根据NVIDIA®当使用CUDA 8.0时,统一内存可以提供显著的性能优势,或者当针对嵌入式硬件时,如NVIDIA Tegra®

若要更改GPU Coder应用程序中的内存分配模式,请使用Malloc模式下拉框更多设置->GPU编码器.在使用命令行界面时,请使用MallocMode生成配置属性并将其设置为任意一个“离散”“统一”

内存极小化

GPU Coder分析CPU和GPU分区之间的数据依赖关系,并进行优化,使CPU和GPU分区之间的数据依赖关系最小化cudaMemcpy生成的代码中的函数调用。分析还确定了数据必须在CPU和GPU之间复制的最小位置集cudaMemcpy

例如,函数喷火具有在CPU上按顺序处理数据和在GPU上并行处理数据的代码段。

function [out] = foo(input1,input2)…% CPU工作input1 =…input2 =…tmp1 =…tmp2 = ... ... % GPU工作内核1(gpuInput1, gpuTmp1);kernel2(gpuInput2, gpuTmp1, gpuTmp2);kernel3(gpuTmp1, gpuTmp2, gpuOut);…% CPU工作…=结束

一个未优化的CUDA实现可能有多个cudaMemcpy函数调用来传递所有输入gpuInput1, gpuInput2,和暂时的结果gpuTmp1, gpuTmp2在内核调用之间。因为中间结果gpuTmp1, gpuTmp2不使用GPU外部,它们可以存储在GPU内存内,从而减少cudaMemcpy函数调用。这些优化提高了生成代码的整体性能。优化后的实现为:

gpuInput1 = input1;gpuInput2 = input2;kernel1<<< >>>(gpuInput1, gpuTmp1);kernel2<<< >>>(gpuInput2, gpuTmp1, gpuTmp2);kernel3<<< >>>(gpuTmp1, gpuTmp2, gpuOut);out = gpuOut;

消除冗余cudaMemcpy调用,GPU Coder分析给定变量的所有使用和定义,并使用状态标志来执行最小化。这个表中显示了原始代码的示例以及生成的代码的外观。

原始代码 优化生成的代码
A(:) = ... ... for i = 1:N gB = kernel1(gA);gA = kernel2(gB);if (somcondition) gC = kernel3(gA, gB);... ... = C;
A(:) =…A_isDirtyOnCpu = true;i = 1:N if (A_isDirtyOnCpu) gA = A;A_isDirtyOnCpu = false;end gB = kernel1(gA);gA = kernel2(gB);if (somcondition) gC = kernel3(gA, gB);C_isDirtyOnGpu = true;if (C_isDirtyOnGpu) C = gC;C_isDirtyOnGpu = false;结束,结束;

_isDirtyOnCpuflag告诉GPU Coder内存优化例程,给定变量被声明并在CPU或GPU上使用。