GPU内存分配和最小化

离散和管理模式

GPU编码器™为您提供访问两个不同的内存分配(的malloc)在CUDA模式下可用®编程模型,cudaMalloccudaMallocManagedcudaMallocAPI适用于传统上独立的CPU、GPU全局内存。cudaMallocManaged适用于统一存储

从一个程序员点,一个传统的计算机体系结构要求数据被分配,并且CPU和GPU存储器空间之间共享。为应用程序这两个存储空间之间管理数据传输的需求增加了复杂性增加。统一存储器中创建管理存储器的池,在CPU和GPU之间共享。被管理的存储器是在CPU和通过单个指针GPU两者访问。通过将数据迁移到需要它,同时隐藏从程序迁移详细信息的设备的统一存储器试图优化存储器的性能。尽管统一存储简化了编程模型,它要求设备同步调用时,被在CPU上访问写在GPU上的数据。GPU编码器插入这些同步调用。根据NVIDIA®,统一内存可以提供显著的性能优势时,使用CUDA 8.0,或当目标嵌入式硬件,如的NVIDIA Tegra®

要改变内存分配模式在GPU编码器应用中,从Malloc模式下的下拉框更多设置 - > GPU编码器。当使用命令行接口,使用MallocMode构建配置属性并将其设置为“离散”'统一'

内存最小化

GPU编码器分析CPU和GPU分区和进行优化之间的数据依赖性最小化的数量cudaMemcpy在生成的代码的函数调用。该分析还确定的位置处的最小集合,其中数据必须CPU和GPU之间通过使用复制cudaMemcpy

例如,该功能FOO在CPU上顺序处理数据和在GPU上并行处理数据的代码段。

function [out] = foo(input1,input2)…% CPU工作kernel2 (gpuInput2 gpuTmp1 gpuTmp2);kernel3 (gpuTmp1 gpuTmp2 gpuOut);…% CPU工作…=输出结束

未经优化的CUDA实现都可能有多个cudaMemcpy函数调用来传输所有输入gpuInput1,gpuInput2和临时结果gpuTmp1,gpuTmp2内核调用之间。因为中间结果gpuTmp1,gpuTmp2所述GPU外不使用,可以将它们存储导致较少的GPU存储器内cudaMemcpy函数调用。这些优化提高了生成代码的总体性能。优化后的实现为:

gpuInput1 = INPUT1;gpuInput2 =输入2;kernel1 <<< >>>(gpuInput1,gpuTmp1);kernel2 <<< >>>(gpuInput2,gpuTmp1,gpuTmp2);kernel3 <<< >>>(gpuTmp1,gpuTmp2,gpuOut);OUT = gpuOut;

为了消除多余的cudaMemcpy电话,GPU编码器分析所有用途和特定变量的定义和使用状态标记来执行最小化。原代码,以及生成的代码看起来像在这个表所示的例子。

原始代码 优化生成的代码
A(:) = ......对于i = 1:N GB = kernel1(GA);GA = kernel2(GB);如果(somecondition)GC = kernel3(GA,GB);端...端...... = C;
A(:) = ... A_isDirtyOnCpu = TRUE;...对于i = 1:N如果(A_isDirtyOnCpu)GA = A;A_isDirtyOnCpu = FALSE;结束GB = kernel1(GA);GA = kernel2(GB);如果(somecondition)GC = kernel3(GA,GB);C_isDirtyOnGpu = TRUE;端...端...如果(C_isDirtyOnGpu)C = GC;C_isDirtyOnGpu = FALSE;结束... = C;

_isDirtyOnCpu标志讲述其中所述给定变量声明并且在CPU或GPU然后或者是用例程GPU编码器存储器的优化。