分类导航

程序问答发布时间：2022-06-01 发布网站：大佬教程 code.js-code.com

大佬教程收集整理的这篇文章主要介绍了与 CUDA、Clang 和 LLVM IR 斗争，并得到：CUDA 失败：'无效的设备功能'，大佬教程大佬觉得挺不错的，现在分享给大家，也给大家做个参考。

如何解决与 CUDA、Clang 和 LLVM IR 斗争，并得到：CUDA 失败：'无效的设备功能'？

开发过程中遇到与 CUDA、Clang 和 LLVM IR 斗争，并得到：CUDA 失败：'无效的设备功能'的问题如何解决？下面主要结合日常开发的经验，给出你关于与 CUDA、Clang 和 LLVM IR 斗争，并得到：CUDA 失败：'无效的设备功能'的解决方法建议，希望对你解决与 CUDA、Clang 和 LLVM IR 斗争，并得到：CUDA 失败：'无效的设备功能'有所启发或帮助；

我正在尝试在配备 V100 GPU、CUDA 10.1 和 LLVM 11（从源代码构建）的 PowerPC 系统（RHel 7.6，无根访问权限）上使用 LLVM 传递优化 CUDA 代码。此外，我在一个简单的 C++ 代码上测试了 clang、lli、 和 opt，一切正常。

经过几天的搜索、阅读和反复试验，我设法编译了一个简单的 CUDA 源代码。代码就是著名的axpy：

#include <iostream>

#define cudacheckerror()                                       \
  {                                                            \
    cudaError_t e = cudaGetLastError();                        \
    if (e != cudasuccess) {                                    \
      printf("Cuda failure %s:%d: '%s'\n",__file__,__liNE__,\
             cudaGetErrorString(E));                           \
      exit(EXIT_FAILURE);                                      \
    }                                                          \
  }

__global__ voID axpy(float a,float* x,float* y) {
  Y[threadIDx.x] = a * x[threadIDx.x];
}

int main(int argc,char* argv[]) {
  const int kDataLen = 4;

  float a = 2.0f;
  float host_x[kDataLen] = {1.0f,2.0f,3.0f,4.0f};
  float host_Y[kDataLen];

  // copy input data to device.
  float* device_x;
  float* device_y;
  cudamalloc(&device_x,kDataLen * sizeof(float));
  cudamalloc(&device_y,kDataLen * sizeof(float));
  cudamemcpy(device_x,host_x,kDataLen * sizeof(float),cudamemcpyHostToDevicE);
  
  // Launch the kernel.
  axpy<<<1,kDataLen>>>(a,device_x,device_y);
  cudacheckerror();
  
  // copy output data to host.
  cudaDeviceSynchronize();
  cudamemcpy(host_y,device_y,cudamemcpyDeviCEToHost);
 
  // Print the results.
  for (int i = 0; i < kDataLen; ++i) {
    std::cout << "Y[" << i << "] = " << host_Y[i] << "\n";
  }

  cudaDevicereset();
  return 0;
}

我在Makefile中写了这样的编译步骤（我知道，它的风格可以改进！）：我还没有添加任何 LLVM 通行证。

BIN_file=axpy
SRC_file=$(BIN_filE).cu

main: $(BIN_filE)

$(BIN_filE).ll: $(SRC_filE)
    clang++ -stdlib=libc++ -Wall $(BIN_filE).cu --cuda-gpu-arch=sm_70 -S -c -emit-llvm

$(BIN_filE)-cuda-nvptx64-nvIDia-cuda-sm_70.ll: $(SRC_filE)
    clang++ -stdlib=libc++ -Wall $(BIN_filE).cu --cuda-gpu-arch=sm_70 -S -c -emit-llvm

$(BIN_filE).ptx: $(BIN_filE)-cuda-nvptx64-nvIDia-cuda-sm_70.ll
    llc -march=nvptx64 $(BIN_filE)-cuda-nvptx64-nvIDia-cuda-sm_70.ll -o $(BIN_filE).ptx

$(BIN_filE)_dlink.o: $(BIN_filE).ptx
    ptxas -m64 --gpu-name=sm_70 $(BIN_filE).ptx -o $(BIN_filE).ptx.o
    fatbinary --64 --create $(BIN_filE).fatbin --image=profile=sm_70,file=$(BIN_filE).ptx.o
    nvcc $(BIN_filE).fatbin -arch=sm_70 -dlink -o $(BIN_filE)_dlink.o -rdc=true

# For the host code:
$(BIN_filE).o: $(BIN_filE).ll
    llc -mcpu=ppc64 $(BIN_filE).ll -o $(BIN_filE).s
    clang++ -c $(BIN_filE).s -o $(BIN_filE).o

# link both object files together with a linker:
$(BIN_filE): $(BIN_filE).o $(BIN_filE)_dlink.o
    nvcc $(BIN_filE).o $(BIN_filE)_dlink.o -o $(BIN_filE) -arch=sm_70 -lc++

clean:
    rm *.ll *.s *.ptx *.ptx.o *.fatbin $(BIN_filE) $(BIN_filE).o $(BIN_filE)_dlink.o

似乎所有步骤都运行顺利，没有任何警告，但运行可执行文件后，出现错误： Cuda failure axpy.cu:33: 'invalID device function'

我还用以下命令替换了最后一个链接器命令，它运行正常，但出现相同的错误。

clang++ -stdlib=libc++ $(BIN_filE).o $(BIN_filE)_dlink.o -o $(BIN_filE) -lcuda -lcudart -lcudadevrt \
        -L/path-to-gcc-lib/

我非常感谢任何帮助。值得一提的是，我需要利用 CUDA 的一些现代功能，尤其是所有合作组，所以我认为最近的 LLVM 版本可能会起作用。想知道更改 LLVM 版本是否有帮助。

-- 编辑：

cuobjdump axpy.fatbin -ptx -sass 的输出：

Fatbin elf code:
================
arch = sm_70
code version = [1,7]
producer = <unkNown>
host = linux
compile_size = 64bit

code for sm_70
        Function : _Z4axpyfPfS_
.headerflags    @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
   /*0000*/           MOV R1,c[0x0][0x28] ;                       /* 0x00000a0000017a02 */
                                                                   /* 0x000fd00000000f00 */
   /*0010*/      @!PT SHFl.IDX PT,RZ,RZ ;                /* 0x000000fffffff389 */
                                                                   /* 0x000fe200000e00ff */
   /*0020*/           IADD3 R1,R1,-0x18,RZ ;                    /* 0xffffffe801017810 */
                                                                   /* 0x000fe20007ffe0ff */
   /*0030*/           IMAD.MOV.U32 R9,c[0x0][0x16c] ;     /* 0x00005b00ff097624 */
                                                                   /* 0x000fe200078e00ff */
   /*0040*/           MOV R8,c[0x0][0x168] ;                      /* 0x00005a0000087a02 */
                                                                   /* 0x000fe40000000f00 */
   /*0050*/           IADD3 R2,P0,c[0x0][0x20],RZ ;         /* 0x0000080001027a10 */
                                                                   /* 0x000fc80007f1e0ff */
   /*0060*/           IADD3.X R3,c[0x0][0x24],!PT ;  /* 0x00000900ff037a10 */
                                                                   /* 0x000fd000007fe4ff */
   /*0070*/           ST.E.64.SYS [R2+0x8],R8 ;                   /* 0x0000000802007385 */
                                                                   /* 0x0001e8000010eb08 */
   /*0080*/           LD.E.64.SYS R4,[R2+0x8] ;                   /* 0x0000000802047980 */
                                                                   /* 0x000ea2000010eb00 */
   /*0090*/           IMAD.MOV.U32 R10,c[0x0][0x170] ;    /* 0x00005c00ff0a7624 */
                                                                   /* 0x000fe200078e00ff */
   /*00a0*/           MOV R11,c[0x0][0x174] ;                     /* 0x00005d00000b7a02 */
                                                                   /* 0x000fe20000000f00 */
   /*00b0*/           IMAD.MOV.U32 R15,c[0x0][0x160] ;    /* 0x00005800ff0f7624 */
                                                                   /* 0x000fe200078e00ff */
   /*00c0*/           S2R R13,SR_TID.X ;                          /* 0x00000000000d7919 */
                                                                   /* 0x000eac0000002100 */
   /*00d0*/           ST.E.64.SYS [R2+0x10],R10 ;                 /* 0x0000001002007385 */
                                                                   /* 0x000fe8000010eb0a */
   /*00e0*/           LD.E.64.SYS R6,[R2+0x10] ;                  /* 0x0000001002067980 */
                                                                   /* 0x000ee8000010eb00 */
   /*00f0*/           ST.E.SYS [R2],R15 ;                         /* 0x0000000002007385 */
                                                                   /* 0x000fe8000010e90f */
   /*0100*/           LD.E.SYS R0,[R2] ;                          /* 0x0000000002007980 */
                                                                   /* 0x000e22000010e900 */
   /*0110*/           IMAD.WIDE.U32 R4,R13,0x4,R4 ;             /* 0x000000040d047825 */
                                                                   /* 0x004fd400078e0004 */
   /*0120*/           LD.E.SYS R5,[R4] ;                          /* 0x0000000004057980 */
                                                                   /* 0x000e22000010e900 */
   /*0130*/           IMAD.WIDE.U32 R6,R6 ;             /* 0x000000040d067825 */
                                                                   /* 0x008fe400078e0006 */
   /*0140*/           FMul R9,R0,R5 ;                            /* 0x0000000500097220 */
                                                                   /* 0x001fd00000400000 */
   /*0150*/           ST.E.SYS [R6],R9 ;                          /* 0x0000000006007385 */
                                                                   /* 0x000fe2000010e909 */
   /*0160*/           EXIT ;                                       /* 0x000000000000794d */
                                                                   /* 0x000fea0003800000 */
   /*0170*/           BRA 0x170;                                   /* 0xfffffff000007947 */
                                                                   /* 0x000fc0000383ffff */
    .......................

输出以点结束。

解决方法

该问题与 PowerPC 架构无关。我需要使用 fatbin 将 -Xclang -fcuda-include-gpubinary -Xclang axpy.fatbin 文件传递给主机端编译命令以复制整个编译行为。

这是更正后的 Makefile：

BIN_FILE=axpy
SRC_FILE=$(BIN_FILE).cu

main: $(BIN_FILE)

# Host Side
$(BIN_FILE).ll: $(SRC_FILE) $(BIN_FILE).fatbin
    clang++ -stdlib=libc++ -Wall -Werror $(BIN_FILE).cu -march=ppc64le --cuda-host-only -relocatable-pch \
        -Xclang -fcuda-include-gpubinary -Xclang $(BIN_FILE).fatbin -S -g -c -emit-llvm

$(BIN_FILE).o: $(BIN_FILE).ll
    llc -march=ppc64le $(BIN_FILE).ll -o $(BIN_FILE).s
    clang++ -c -Wall $(BIN_FILE).s -o $(BIN_FILE).o

# GPU Side
$(BIN_FILE)-cuda-nvptx64-nvidia-cuda-sm_70.ll: $(SRC_FILE)
    clang++ -x cuda -stdlib=libc++ -Wall -Werror $(BIN_FILE).cu --cuda-device-only \
        --cuda-gpu-arch=sm_70 -S -g -emit-llvm

$(BIN_FILE).ptx: $(BIN_FILE)-cuda-nvptx64-nvidia-cuda-sm_70.ll
    llc -march=nvptx64 -mcpu=sm_70 -mattr=+ptx64 $(BIN_FILE)-cuda-nvptx64-nvidia-cuda-sm_70.ll -o $(BIN_FILE).ptx

$(BIN_FILE).ptx.o: $(BIN_FILE).ptx
    ptxas -m64 --gpu-name=sm_70 $(BIN_FILE).ptx -o $(BIN_FILE).ptx.o

$(BIN_FILE).fatbin: $(BIN_FILE).ptx.o
    fatbinary --64 --create $(BIN_FILE).fatbin --image=profile=sm_70,file=$(BIN_FILE).ptx.o \
        --image=profile=compute_70,file=$(BIN_FILE).ptx -link

$(BIN_FILE)_dlink.o: $(BIN_FILE).fatbin
    nvcc $(BIN_FILE).fatbin -gencode arch=compute_70,code=sm_70 \
        -dlink -o $(BIN_FILE)_dlink.o -lcudart -lcudart_static -lcudadevrt

# Link both object files together (either nvcc or clang works herE):
$(BIN_FILE): $(BIN_FILE).o $(BIN_FILE)_dlink.o
    nvcc $(BIN_FILE).o $(BIN_FILE)_dlink.o -o $(BIN_FILE) -arch=sm_70 -lc++