CUDA内核函数从外部文件调用“ __device__ __host__”函数

问题描述

我正在尝试将CUDA和C ++混合使用。我遇到以下错误：

main.cpp：定义“ main（）”。调用“ gpu_main（）”和“ add_test（）”
|
|-> add_func.cu：定义“ gpu_main（）”和“ __global__ void add（）”作为内核。“ add（）”将调用“ add_test（）”
|
|-> basic_add.cu：定义“ __host__ __device__ int add_test（int a，int b）“

我以这种方式编译代码：

nvcc basic_add.cu -c
nvcc -rdc=true add_func.cu -c
g++ main.cpp -c
g++ -o main main.o basic_add.o add_func.o -lcudart -L/usr/local/cuda/lib64

在第二步，它给了我这个错误：

add_func.cu（14）：错误：从计算机中调用 host 函数（“ add_test”） 不允许全局功能（“添加”）

add_func.cu（14）：错误：设备中的标识符“ add_test”未定义 代码

有人知道如何解决此问题吗？还是我不应该从外部文件调用主机和设备功能？谢谢。

---

代码如下（仅供参考）：

• basic_add.h：

#ifndef BASIC_ADD_H_
#define BASIC_ADD_H_

int add_test( int a,int b );

#endif

• basic_add.cu：

__host__ __device__ int add_test(int a,int b)
{
    return a + b;
}

• add_func.h

#ifndef ADD_FUNC_H_
#define ADD_FUNC_H_

#include <iostream>
#include <math.h>
#include "basic_add.h"

int gpu_main(void);

#endif

• add_func.cu

#include "add_func.h"

// Kernel function to add the elements of two arrays
__global__
void add(int n,float *x,float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  printf("gridDim %d,blockDim %d,blockIdx %d,threadIdx %d\n",gridDim.x,blockDim.x,blockIdx.x,threadIdx.x);
  for (int i = index; i < n; i += stride)
  {
    y[i] = add_test(x[i],y[i]);
    printf("blockIdx %d,threadIdx %d,%d\n",threadIdx.x,i);
    break;
  }
}

int gpu_main(void)
{
  int N = 1<<10;
  float *x,*y;

  // Allocate Unified Memory – accessible from CPU or GPU
  cudaMallocManaged(&x,N*sizeof(float));
  cudaMallocManaged(&y,N*sizeof(float));

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the GPU
  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks,blockSize>>>(N,x,y);

  // Wait for GPU to finish before accessing on host
  cudaDeviceSynchronize();

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError,fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  cudaFree(x);
  cudaFree(y);
  
  return 0;
}

• main.cpp：

#include <iostream>
#include <math.h>
#include "add_func.h"
#include "basic_add.h"

int main(void)
{
  gpu_main();
  int a = add_test(1,2);
  std::cout << a << std::endl;
  
  return 0;
}

解决方法

您在这里尝试做的一件事是行不通的。如果您想要一个用__host__ __device__装饰的函数，首先您应该在任何地方都以相同的方式装饰它（即，在声明它的头文件中），这样的函数就不能直接从.cpp文件，除非您用.cpp编译该nvcc文件并通过-x cu作为编译命令行开关，所以您最好将其放入{{1} }从我的角度来看。

您也没有正确链接可重定位的设备代码，但这是可以修复的。

如果您想拥有一个.cu函数，该函数可通过{.1} __host__ __device__，那么我唯一的建议就是为其提供包装。

以下是我所能找到的最接近的信息：

.cpp