Exponenta Banner

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

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

GPU Coder™ предоставляет вам доступ к двум различным режимам (malloc) выделения памяти, доступным в модели программирования CUDA®, cudaMalloc и cudaMallocManaged. API cudaMalloc применим к традиционно отдельному центральному процессору и глобальным памятям графического процессора. 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
Поддержка