无法 link 从 CUBIN 中间表示生成的 CUDA 目标文件

Question

来自 NVIDIA documentation，当生成 PTX、CUBIN 或 FATBIN 时，主机代码将从文件中丢弃。现在我有了主机代码 (main.cu) 和设备代码 (shared.cu)。当使用 nvcc 选项 nvcc -c main.cu shared.cu 或什至使用 nvcc -dc main.cu shared.cu 将每个文件编译为 *.o 并使用选项 nvcc -link main.o shared.o 链接它们时，我可以生成可执行文件。但是，当 shared.cu 被编译为 shared.cubin 并进一步编译为 *.o 时，链接将失败并出现错误 tmpxft_00001253_00000000-4_main.cudafe1.cpp:(.text+0x150): undefined reference to <KERNEL FUNCTION>

这里我想知道 shared.cu 只包含设备代码，即使删除主机代码为什么链接会失败。

源代码文件是main.cu

#include <stdio.h>
#include <cuda_runtime_api.h>
#include <cuda_runtime.h>
#include <cuda.h>
#include "shared.h"
 int main()
{
        int a[5]={1,2,3,4,5};
        int b[5]={1,1,1,1,1};
        int c[5];
        int i;

        int *dev_a;
        int *dev_b;
        int *dev_c;

        cudaMalloc( (void**)&dev_a, 5*sizeof(int) );
        cudaMalloc( (void**)&dev_b, 5*sizeof(int) );
        cudaMalloc( (void**)&dev_c, 5*sizeof(int) );

        cudaMemcpy(dev_a, a , 5 * sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(dev_b, b , 5 * sizeof(int), cudaMemcpyHostToDevice);

        add<<<1,5>>>(dev_a,dev_b,dev_c);

        cudaMemcpy(&c,dev_c,5*sizeof(int),cudaMemcpyDeviceToHost);

        for(i = 0; i < 5; i++ )
        {
                printf("a[%d] + b[%d] = %d\n",i,i,c[i]);
        }
        cudaFree( dev_a);
        cudaFree( dev_b);
        cudaFree( dev_c);
        return 0;
}

shared.cu

#include<stdio.h>

__global__  void add(int *dev_a, int *dev_b, int *dev_c){

        //allocate shared memory
        __shared__ int a_shared[5];
        __shared__ int b_shared[5];
        __shared__ int c_shared[5];
        {
                //get data in shared memory
                a_shared[threadIdx.x]=dev_a[threadIdx.x];
                __syncthreads();

                b_shared[threadIdx.x]=dev_b[threadIdx.x];
                __syncthreads();

                //perform the addition in the shared memory space
                c_shared[threadIdx.x]= a_shared[threadIdx.x] + b_shared[threadIdx.x];
                __syncthreads();

                //shift data back to global memory
                dev_c[threadIdx.x]=c_shared[threadIdx.x];
                __syncthreads();
        }
}

shared.h

#ifndef header
#define header

extern __global__  void add(int *dev_a, int *dev_b, int *dev_c);

#endif

Answer 1

我相信您假设 "device code only" 文件（例如您的 shared.cu）不包含主机代码。这实际上是不正确的。

内核函数生成主机和设备代码构造，这些构造由 CUDA 预处理器 (cudafe) 生成并分开。参考the documentation.

中的CUDA编译轨迹

请注意，主机和设备代码最初是分开的，随后创建 .cudafe1.stub.c 文件，然后将其传递到主机端（即与 cubin 路径分开，有效地启用 linking).

如图所示，.cudafe1.stub.c文件并没有成为cubin的一部分，而是进入主机端处理流，最终成为fatbinary文件的一部分。

如果你只处理到 cubin，你将丢弃这个。cudafe1.stub.c，这是最终 link 创建可执行 fat 二进制文件所必需的。

因此，由于存根文件中缺少引用，尝试仅使用 cubin 创建胖二进制文件 link 将失败，该引用在您的问题的错误输出中指示。

无法 link 从 CUBIN 中间表示生成的 CUDA 目标文件

Unable to link CUDA object files generated from the CUBIN intermediate representation

cuda

linker

nvidia

nvcc

ptx