В одном файле запишите функцию точки входа nestedLoop который принимает два векторных входов A,B размера 32x512. Функция имеет два вложенных for- циклы разной длины итерации, по одному для работы вдоль столбца и по одному для работы вдоль строки. Первый вложенный цикл вычисляет сумму двух векторных входов, а второй вложенный цикл масштабирует сумму в три раза.

function [C] = nestedLoop(A, B)
    G = zeros(32, 512);
    C = zeros(32, 512);        
   
    coder.gpu.kernelfun();
    % This nested loop will be fused
    for i = 1:32
       for j = 1:512
           G(i,j) = A(1,j) + B(1,j);
       end
    end

    coder.gpu.nokernel();  
    for i = 1:32
       for j = 1:512
           C(i,j) = G(i,j) * 3;
       end
    end    
end

Используйте codegen функция для генерации функции MEX CUDA.

cfg = coder.gpuConfig('mex');
cfg.GenerateReport = true;
codegen -config cfg -args {ones(1,512,'double'),ones(1,512,'double')} nestedLoop

GPU Coder создает два ядра: nestedLoop_kernel1 для выполнения расчетов G(i,j) = A(1,j) + B(1,j); первого вложенного цикла и nestedLoop_kernel2 ядро для выполнения расчетов C(i,j) = G(i,j) * 3; второго вложенного цикла. Второе ядро создается для внутреннего цикла второго вложенного цикла. The noKernel прагма применима только к циклу, который немедленно следует за оператором. Показаны фрагменты сгенерированных ядер.

static __global__ __launch_bounds__(512, 1) void nestedLoop_kernel1(const real_T
  B[512], const real_T A[512], real_T G[16384])
{
  uint32_T threadId;
  ...
  if (i < 32) {
    G[i + (j << 5)] = A[j] + B[j];
  }
}
static __global__ __launch_bounds__(512, 1) void nestedLoop_kernel2(real_T G
  [16384], int32_T i, real_T C[16384])
{
  uint32_T threadId;
  ...;
  if (j < 512) {
    C[i + (j << 5)] = G[i + (j << 5)] * 3.0;
  }

Фрагмент основной функции показывает, что генератор кода сплавил первый вложенный цикл, как обозначено параметрами запуска ядра. Как упоминалось ранее, внешний контур второго вложенного цикла является тем, который не сопоставлен с ядром. Следовательно, генератор кода помещает for-loop оператор непосредственно перед вызовом второго ядра CUDA nestedLoop_kernel2.

void nestedLoop(const real_T A[512], const real_T B[512], real_T C[16384])
{
  int32_T i;
  ...
  //  These two loops will be fused
  cudaMemcpy(gpu_B, (void *)&B[0], 4096UL, cudaMemcpyHostToDevice);
  cudaMemcpy(gpu_A, (void *)&A[0], 4096UL, cudaMemcpyHostToDevice);
  nestedLoop_kernel1<<<dim3(32U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_B, *gpu_A, *
    gpu_G);
  for (i = 0; i < 32; i++) {
    nestedLoop_kernel2<<<dim3(1U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_G, i,
      *gpu_C);
    C_dirtyOnGpu = true;
  }
...
  cudaFree(*gpu_C);
}