如何针对 cublas_device 编译/链接具有不同计算能力的代码？

Posted 2023-02-16

【中文标题】如何针对 cublas_device 编译/链接具有不同计算能力的代码？

【英文标题】：How to compile/link code with different compute capabilities against cublas_device?

【发布时间】：2013-06-25 14:03:23

【问题描述】：

我正在我的一个内核中使用动态并行（和 cublas），并希望为 sm_20 提供一个备用内核。在 maxentropy_cuda.cu 中，我编写了两个内核并使用 CUDA_ARCH 为架构>=3.5 编译动态并行内核。这很好用。

Makefile 的一部分：

    nvcc $(NVCCFLAGS) -gencode arch=compute_35,code=sm_35 -gencode arch=compute_20,code=sm_20 $(CINCL) -M maxentropy_cuda.cu -o maxentropy_cuda.d 
    nvcc --device-c $(NVCCFLAGS) -gencode arch=compute_35,code=sm_35 -gencode arch=compute_20,code=sm_20 -x cu maxentropy_cuda.cu -o maxentropy_cuda.o 

当我将它链接到另一个文件中的内核时：

    nvcc --cudart static --relocatable-device-code=true -link -gencode arch=compute_35,code=sm_35  -gencode arch=compute_20,code=sm_20 $(LIBPATHS) -o main main.o selgen.o maxentropy.o maxentropy_omp.o maxentropy_cuda.o maxentropy_kernels.o $(OBJINFRA) $(LIBS)  -lcublas_device  -lcudadevrt

我收到以下错误：

nvlink fatal   : could not find compatible device code in /opt/cuda/lib64/libcublas_device.a
make: *** [main] Error 255

当然我不需要 libcublas_device 作为备用内核...

有没有办法在一个二进制文件中同时获得两种计算能力？ （我使用的是 CUDA 5.5）

编辑：示例（尚未测试输出...）：

#include <stdio.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>


__global__ void calc_dev(double* u, double* s, double* r) 
#if __CUDA_ARCH__<350
    printf("should not have been called - compiled for the wrong CUDA architecture");
#else

    cublasHandle_t cnpHandle;
    cublasStatus_t status = cublasCreate(&cnpHandle);

    if (status != CUBLAS_STATUS_SUCCESS)    
        printf("error while initializing cublas\n");
        return;
    

    status = cublasDdot(cnpHandle,5,u,1,s,1,r);
    cudaDeviceSynchronize();
    if (status != CUBLAS_STATUS_SUCCESS)        
        printf("cublas error: u x s\n");
        return;
    
#endif


void calc_host(double* u, double* s, double* r) 
    cublasHandle_t cnpHandle;
    cublasStatus_t status = cublasCreate(&cnpHandle);
    cublasSetPointerMode(cnpHandle, CUBLAS_POINTER_MODE_DEVICE);

    if (status != CUBLAS_STATUS_SUCCESS)    
        printf("error while initializing cublas\n");
        return;
    

    status = cublasDdot(cnpHandle,5,u,1,s,1,r);
    cudaThreadSynchronize();
    if (status != CUBLAS_STATUS_SUCCESS)        
        printf("cublas error: u x s\n");
        return;
    



int main(int argc, char** argv) 
    const int n = 5;

    double u[n] = 0,1,2,4,8;
    double s[n] = 1, 0.64570312500000004,
              0.44203125000000004, 0.65804687500000003, 0.71976562500000008;
    double r = 0.0;
    double *dev_s,*dev_u,*dev_r;
    cudaMalloc( (void**)&dev_s, sizeof(double)*n);
    cudaMalloc( (void**)&dev_u, sizeof(double)*n);
    cudaMalloc( (void**)&dev_r, sizeof(double));
    cudaMemcpy( dev_s, s, n*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy( dev_u, u, n*sizeof(double), cudaMemcpyHostToDevice);

    #if __CUDA_ARCH__>=350
        calc_dev<<<1,1>>>(dev_s,dev_u,dev_r);
    #else
        calc_host(dev_s,dev_u,dev_r);
    #endif

    cudaMemcpy( &r, dev_r, sizeof(double), cudaMemcpyDeviceToHost);
    printf("%.3f\n",r);
    return 0;

Nsight：

make all 
Building file: ../main.cu
Invoking: NVCC Compiler
/usr/local/cuda-5.5/bin/nvcc -G -g -O0 -gencode arch=compute_20,code=sm_20 -gencode arch=compute_35,code=sm_35 -odir "" -M -o "main.d" "../main.cu"
/usr/local/cuda-5.5/bin/nvcc --device-c -G -O0 -g -gencode arch=compute_20,code=sm_20 -gencode arch=compute_35,code=sm_35  -x cu -o  "main.o" "../main.cu"
Finished building: ../main.cu

Building target: test_sm_compatibility
Invoking: NVCC Linker
/usr/local/cuda-5.5/bin/nvcc --cudart static --relocatable-device-code=true -gencode arch=compute_20,code=sm_20 -gencode arch=compute_35,code=sm_35 -link -o  "test_sm_compatibility"  ./main.o   -lcublas -lcublas_device
nvlink fatal   : could not find compatible device code in /usr/local/cuda-5.5/bin/../targets/x86_64-linux/lib/libcublas_device.a
make: *** [test_sm_compatibility] Error 255

【问题讨论】：

如果您想提供一个完整的示例（演示 cuda 代码的简单文件，以及产生错误的完整编译步骤），我会看一下。这个问题还不够，我宁愿不要猜测你的代码和宏控制是什么样的。

@RobertCrovella 谢谢，我添加了一个示例。我希望它显示了我正在尝试做的事情。

【参考方案1】：

现在在 CUDA 6.0 中可以做到这一点

您发布的示例使用您显示的编译命令按原样编译。

唯一的困难是必须忽略各种 nvlink 警告：

nvlink warning : SM Arch ('sm_20') not found in '/usr/local/cuda/bin/..//lib64/libcublas_device.a'

但是正确的可执行文件是正确构建的。