logo标签列表

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

c++cuda
3

我注意到一种奇怪的现象，可以使用三角括号符号在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);

I get:

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

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

这段话涉及IT技术方面的内容，需要更具体的上下文才能进行精确翻译。
1个回答
1
我能理解这可能有些令人困惑，但是不论你认为发生了什么，Copy 从未在GPU上运行。 CopyKernel 在设备上被调用了三次，但所有的启动都是在主机上进行的。具体如下。
第一个需要理解的见解是如何使CUDA runtime API中的内核编译和启动实际工作。当nvcc编译您的 CopyKernel 和一个基于runtime 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);

ie。启动内核的启动参数被推送到驱动程序上，并调用作为f提供的主机函数。如果f恰好是内核包装器函数（即CopyKernel），则将通过包装器包含的API调用进行内核启动，否则不会。如果f恰好是一个包含运行时API内核调用的主机函数（即Copy），那么该主机代码也会执行相同的操作，最终导致内核启动，只是在调用堆栈中更深处。

这就是您可以将CopyKernel或Copy作为参数提供给TestFunctionPointerLaunch，并且它仍然可以正常工作的方式。从技术上讲，这是未定义的行为，因为CUDA运行时API内部工作的方式是故意模糊的，实现细节可能随时间而变化。但现在它能够工作。

原因是

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

由于Copy是主机函数，nvcc可以在编译时检测到无法编译的问题--只有__global__函数可以被启动，编译器会执行此检查。但是当您传递函数指针时，编译器无法应用该检查。生成的代码可能适用于主机函数或主机内核包装函数，因为运行时支持代码不会（也可能无法）发出能够对函数指针进行内省并确定函数指针不会调用内核的代码。因此，语言规范要求被跳过，事情偶然发生了。我强烈建议不要尝试依赖这种行为。
请问为什么cudaLaunch使用指向CopyKernel的指针？因为函数CopyKernel在那里定义，而不是应该在GPU上运行的设备函数。根据您的第一个代码片段，似乎函数CopyKernel和__device_stub__Z10CopyKernelPKiPi相互调用，从而导致递归。因此，程序没有调用设备函数。 - Virux
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接