В этом примере показано, как использовать nokernel
прагма в функции и препятствует тому, чтобы генератор кода генерировал ядра CUDA для операторов в цикле
В одном файле запишите функцию точки входа nestedLoop
это признает, что два вектора вводят A,B
из размера 32x512
. Функция имеет два, вложил for
- циклы различных длин итерации, один для работы вдоль столбца и один для работы вдоль строки. Первый вложенный цикл вычисляет сумму двух векторных входных параметров, в то время как второй вложенный цикл масштабирует сумму на коэффициент три.
Используйте codegen
функция, чтобы сгенерировать MEX-функцию CUDA.
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);
}