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

Question

我注意到一个奇怪的现象，它允许您在 CUDA 中使用三尖括号表示法来启动主机函数。为了测试这一点，我编写了一个简单的内核，在两个整数数组之间复制数据。请注意，我在 Tesla 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

再次感谢您对这里发生的事情的任何想法。

Answer 1

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

首先需要了解内核是如何编译的，以及它们的启动是如何在 CUDA 运行时 API 中实际工作的。当 nvcc 编译您的 CopyKernel 并为该内核启动运行时 API 样式时，会发出一对 host 函数，如下所示：

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);
}

这些为将内核参数推送到 CUDA 驱动程序并启动内核所需的 API 调用提供了一个包装器。您会注意到内核的执行配置不在这些函数中处理。相反，只要预处理器遇到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__ 函数可以启动，编译器会强制执行此检查。

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

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