В одном файле запишите функцию точки входа 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; из второго вложенного цикла. Второе ядро создается для внутреннего цикла второго вложенного цикла. 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);
}