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