CUDA在使用函数指针时启动主机函数作为内核

CUDA launches host function as kernel when using function pointers

本文关键字：函数 主机 内核 启动 指针 CUDA      更新时间：2023-10-16

我注意到一个奇怪的现象，它允许您使用CUDA中的三角括号符号来启动主机函数。为了测试这一点，我编写了一个简单的内核，在两个整数数组之间复制数据。请注意，我正在特斯拉K40上运行所有这些代码，并使用-gencode arch=compute_35，code=sm_35:进行编译

#ifndef HOST_LAUNCH_H
#define HOST_LAUNCH_H
using namespace std;
// Assumes input and output are both length 32
__global__ void CopyKernel(const int* input, int* output) {
  size_t global_idx = blockIdx.x * blockDim.x + threadIdx.x;
  output[global_idx] = input[global_idx];
}
__host__ void Copy(const int* input, int* output) {
  int* d_input = 0;
  int* d_output = 0;
  cudaMalloc((void**)&d_input, 32 * sizeof(int));
  cudaMalloc((void**)&d_output, 32 * sizeof(int));
  cudaMemcpy(d_input, input, 32 * sizeof(int), cudaMemcpyHostToDevice);
  CopyKernel<<<1,32>>>(d_input, d_output);
  cudaMemcpy(output, d_output, 32 * sizeof(int), cudaMemcpyDeviceToHost);
  cudaFree(d_input);
  cudaFree(d_output);
}
#endif

然后我写了以下单元测试：

#include "host_launch.h"
#include <assert.h>
using namespace std;
__host__ void TestKernelLaunch() {
  int input[32];
  int output[32];
  for(int i = 0; i < 32; i++) {
    input[i] = i;
    output[i] = 0;
  }
  int* d_input = 0;
  int* d_output = 0;
  cudaMalloc((void**)&d_input, 32 * sizeof(int));
  cudaMalloc((void**)&d_output, 32 * sizeof(int));
  cudaMemcpy(d_input, input, 32 * sizeof(int), cudaMemcpyHostToDevice);
  for(int i = 0; i < 32; i++) {
    assert(output[i] == 0);
  }
  CopyKernel<<<1,32>>>(d_input, d_output);
  cudaMemcpy(output, d_output, 32 * sizeof(int), cudaMemcpyDeviceToHost);
  for(int i = 0; i < 32; i++) {
    assert(output[i] == i);
  }
  cudaFree(d_input);
  cudaFree(d_output);
}
__host__ void TestHostLaunch() {
  int input[32];
  int output[32];
  for(int i = 0; i < 32; i++) {
    input[i] = i + 1;
    output[i] = 0;
  }
  int* d_input = 0;
  int* d_output = 0;
  cudaMalloc((void**)&d_input, 32 * sizeof(int));
  cudaMalloc((void**)&d_output, 32 * sizeof(int));
  cudaMemcpy(d_input, input, 32 * sizeof(int), cudaMemcpyHostToDevice);
  for(int i = 0; i < 32; i++) {
    assert(output[i] == 0);
  }
  //Copy<<<1,32>>>(d_input, d_output);
  cudaMemcpy(output, d_output, 32 * sizeof(int), cudaMemcpyDeviceToHost);
  for(int i = 0; i < 32; i++) {
    assert(output[i] == i + 1);
  }
  cudaFree(d_input);
  cudaFree(d_output);
}
__host__ void TestFunctionPointerLaunch(void (*f)(const int*, int*)) {
  int input[32];
  int output[32];
  for(int i = 0; i < 32; i++) {
    input[i] = i + 2;
    output[i] = 0;
  }
  int* d_input = 0;
  int* d_output = 0;
  cudaMalloc((void**)&d_input, 32 * sizeof(int));
  cudaMalloc((void**)&d_output, 32 * sizeof(int));
  cudaMemcpy(d_input, input, 32 * sizeof(int), cudaMemcpyHostToDevice);
  for(int i = 0; i < 32; i++) {
    assert(output[i] == 0);
  }
  f<<<1,32>>>(d_input, d_output);
  cudaMemcpy(output, d_output, 32 * sizeof(int), cudaMemcpyDeviceToHost);
  for(int i = 0; i < 32; i++) {
    assert(output[i] == i + 2);
  }
  cudaFree(d_input);
  cudaFree(d_output);
}
int main() {
  TestKernelLaunch();
  TestFunctionPointerLaunch(CopyKernel);
  TestFunctionPointerLaunch(Copy);
}

如果我取消注释行：

//Copy<<<1,32>>>(d_input, d_output);

我得到：

host_launch_unittest.cu(49): error: a host function call cannot be configured

但等效操作是使用执行的

f<<<1,32>>>(d_input, d_output);

在TestFunctionPointerLaunch中，它传递所有断言。我只是想知道GPU到底在做什么，让这个主机功能的启动行为正确。我编写这些测试是为了隔离行为，但也发现它适用于更复杂的内核/主机函数。此外，我决定对这些进行计时，看看它们是否以某种方式被编译为等效操作：

#include "host_launch.h"
#include <iostream>
#include <assert.h>
using namespace std;
__host__ float MeanCopyTime(const int copy_count, void (*f)(const int*, int*)) {
  int input[32 * copy_count];
  int output[32 * copy_count];
  for(int i = 0; i < 32 * copy_count; i++) {
    input[i] = i;
    output[i] = 0;
  }
  int* d_input = 0;
  int* d_output = 0;
  cudaMalloc((void**)&d_input, 32 * copy_count * sizeof(int));
  cudaMalloc((void**)&d_output, 32 * copy_count * sizeof(int));
  cudaMemcpy(d_input, input, 32 * copy_count * sizeof(int), cudaMemcpyHostToDevice);
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start);
  for(int i = 0; i < copy_count; i++)
    f<<<1,32>>>(d_input + i * 32, d_output + i * 32);
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  float msecs = 0;
  cudaEventElapsedTime(&msecs, start, stop);
  cudaMemcpy(output, d_output, 32 * copy_count * sizeof(int), cudaMemcpyDeviceToHost);
  cudaFree(d_input);
  cudaFree(d_output);
  for(int i = 0; i < 32 * copy_count; i++) {
    assert(output[i] == i);
  }
  return msecs / copy_count;
}
int main() {
  int copy_count = 10000;
  cout << endl;
  cout << "Average Kernel Launch Time: " << MeanCopyTime(copy_count, CopyKernel) << endl;
  cout << "Average Host Function Launch Time: " << MeanCopyTime(copy_count, Copy) << endl;
  cout << endl;
}

对于我的体系结构，这返回：

Average Kernel Launch Time: 0.00420756
Average Host Function Launch Time: 0.169097

如果您对这里发生的事情有任何想法，我们将不胜感激。

我明白为什么这可能有点令人困惑，但尽管你可能认为发生了什么，Copy从未在GPU上运行。CopyKernel在设备上被调用了三次，但所有的启动都是在主机上启动的。以下是操作方法。

所需的第一个见解是阐明内核是如何编译的，以及它们的启动在CUDA运行时API中是如何实际工作的。当nvcc编译CopyKernel和该内核的运行时API式启动时，将发出一对主机函数，如下所示：

void __device_stub__Z10CopyKernelPKiPi(const int *__par0, int *__par1)
{
    if (cudaSetupArgument((void *)(char *)&__par0, sizeof(__par0), (size_t)0Ui64) != cudaSuccess) return;
    if (cudaSetupArgument((void *)(char *)&__par1, sizeof(__par1), (size_t)8Ui64) != cudaSuccess) return;
    {
       volatile static char *__f; 
       __f = ((char *)((void ( *)(const int *, int *))CopyKernel)); 
       (void)cudaLaunch(((char *)((void ( *)(const int *, int *))CopyKernel)));
    };
}
void CopyKernel( const int *__cuda_0,int *__cuda_1)
{
    __device_stub__Z10CopyKernelPKiPi( __cuda_0,__cuda_1);
}

这些为必要的API调用提供了一个包装器，以便将内核参数推送到CUDA驱动程序并启动内核。您会注意到，内核的执行配置没有在这些函数中处理。相反，每当预处理器遇到CopyKernel<<< >>>()调用时，就会发出这样的代码：

(cudaConfigureCall(1, 32)) ? (void)0 : (CopyKernel)(d_input, d_output);

即。内核启动配置被推送到驱动程序，然后调用包装器函数，其中参数被推送到驱动器并启动内核。

那么TestFunctionPointerLaunch中会发生什么呢？基本上是一样的。此代码

f<<<1,32>>>(d_input, d_output);

由CUDA前端预处理器编译

(cudaConfigureCall(1, 32)) ? (void)0 : f(d_input, d_output);

即。内核启动的启动参数被推送到驱动程序上，并调用作为f提供的主机函数。如果f恰好是内核包装器函数（即CopyKernel），那么内核启动将通过包装器包含的API调用来实现，否则不会。如果f恰好是一个主机函数，它本身包含一个运行时API内核调用（即Copy），那么该主机代码也会做同样的事情，并且最终会导致内核启动，只是在调用堆栈的下游。

这就是您可以提供CopyKernel或Copy作为TestFunctionPointerLaunch的参数的方式，它仍然可以工作。从技术上讲，这是一种未定义的行为，因为内核在CUDA运行时API内部启动工作的方式是故意不透明的，实现细节可能会随着时间的推移而变化。但现在它起作用了。

的原因

Copy<<<1,32>>>(d_input, d_output);

不编译，是因为Copy是一个主机函数，nvcc可以在编译时检测到这一点——在语言规范中，只有__global__函数可以启动，编译器强制执行此检查。

但是，当您传递函数指针时，编译器无法应用该检查。生成的代码恰好与宿主函数或宿主内核包装函数一起工作，因为运行时支持代码不会（也可能不会）发出可以对函数指针执行内省并确定函数指针不会调用内核的代码。因此，语言规范要求被跳过，事情就不小心发生了。

我强烈建议不要试图依赖这种行为。

相关文章：