c++ - 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在设备上被调用三次，但所有的启动都是在主机上启动的。方法如下。

首先需要了解内核是如何编译的，以及它们的启动是如何在 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);
}

这些提供了一个围绕必要的 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__功能可以启动并且编译器强制执行此检查。

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

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