Exponenta Banner

Выделение памяти графического процессора и минимизация

Дискретные и режимы автоматического управления

GPU Coder™ предоставляет вам доступ к двум различным выделениям памяти (malloc) режимы, доступные в модели программирования CUDA®, cudaMalloc и cudaMallocManaged. cudaMalloc API применим к традиционно отдельному центральному процессору и глобальным памятям графического процессора. cudaMallocManaged применимо к Unified Memory.

С точки зрения программиста традиционная архитектура ЭВМ требует, чтобы данные были выделены и совместно использованы пространствами памяти центрального процессора и графического процессора. Потребность в приложениях, чтобы справиться с передачами данных между этими двумя пространствами памяти добавляет к увеличенной сложности. Объединенная память создает пул управляемой памяти, совместно использованной центральным процессором и графическим процессором. Управляемая память доступна и для центрального процессора и для графического процессора через один указатель. Объединенная память пытается оптимизировать производительность памяти путем миграции данных на устройство, которому нужен он, одновременно скрывая детали миграции от программы. Хотя объединенная память упрощает модель программирования, она требует синхронизирующих устройством вызовов, когда к данным, записанным на графическом процессоре, получают доступ на центральном процессоре. GPU Coder вставляет эти вызовы синхронизации. Согласно NVIDIA®, объединенная память может обеспечить значительные выигрыши в производительности, когда при помощи CUDA 8.0, или когда предназначение встроило оборудование как NVIDIA Tegra®.

Чтобы изменить режим выделения памяти в приложении GPU Coder, используйте Malloc Mode выпадающее поле под More Settings->GPU Coder. При использовании интерфейса командной строки используйте MallocMode свойство конфигурации сборки и набор это к любому 'discrete' или 'unified'.

Минимизация памяти

GPU Coder анализирует зависимость по данным между разделами центрального процессора и графического процессора и выполняет оптимизацию, чтобы минимизировать количество cudaMemcpy вызовы функции в сгенерированном коде. Анализ также определяет минимальный набор местоположений, где данные должны быть скопированы между центральным процессором и графическим процессором при помощи cudaMemcpy.

Например, функциональный foo имеет разделы кода, которые обрабатывают данные последовательно на центральном процессоре и параллельно на графическом процессоре.

function [out] = foo(input1,input2)
	   …
     % CPU work
			input1 = …
			input2 = …
			tmp1 = …
			tmp2 = …
   	…
     % GPU work
			kernel1(gpuInput1, gpuTmp1);
       kernel2(gpuInput2, gpuTmp1, gpuTmp2);
       kernel3(gpuTmp1, gpuTmp2, gpuOut);

   	…
     % CPU work
       … = out

end

Неоптимизированная реализация CUDA может потенциально иметь несколько cudaMemcpy вызовы функции передать все входные параметры gpuInput1,gpuInput2, и временные результаты gpuTmp1,gpuTmp2 между вызовами ядра. Поскольку промежуточное звено заканчивается gpuTmp1,gpuTmp2 не используются вне графического процессора, они могут храниться в памяти графического процессора, приводящей к меньшему количеству 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 (somecondition)
      gC = kernel3(gA, gB);
   end
   …
end
…
… = C;

A(:) = …
A_isDirtyOnCpu = true;
…
for i = 1:N
   if (A_isDirtyOnCpu)
      gA = A;
      A_isDirtyOnCpu = false;
   end
   gB = kernel1(gA);
   gA = kernel2(gB);
   if (somecondition)
      gC = kernel3(gA, gB);
      C_isDirtyOnGpu = true;
   end
   …
end
…
if (C_isDirtyOnGpu)
   C = gC;
   C_isDirtyOnGpu = false;
end
… = C;

_isDirtyOnCpu флаг говорит оптимизацию памяти GPU Coder о стандартных программах, где данная переменная объявляется и использовала или на центральном процессоре или на затем графическом процессоре.

Документация GPU Coder
Поддержка