如何从“Cuda Fortran 内核/设备例程”调用“Cuda C”设备例程……？答案

【问题标题】：How to call “Cuda C” device routine from “Cuda Fortran Kernel/device routine”…?如何从“Cuda Fortran 内核/设备例程”调用“Cuda C”设备例程……？
【发布时间】：2022-01-21 15:11:45
【问题描述】：

我正在尝试从 Cuda Fortran 中的内核调用 Cuda C 中的设备例程，但出现链接错误。有人可以帮我解决这个问题吗...？

文件 1:SortbyKey.f95

MODULE TEST
   USE CUDAFOR
   CONTAINS 
   ATTRIBUTES(GLOBAL) SUBROUTINE SORTING()
   IMPLICIT NONE

   INTERFACE
      ATTRIBUTES(DEVICE) SUBROUTINE SORTBYKEY(D,M) BIND(c, name='SORTBYKEY_')
         USE, INTRINSIC :: iso_c_binding
         INTEGER(KIND=4),VALUE::M
         REAL(KIND=8)::D(M)
      END SUBROUTINE SORTBYKEY
   END INTERFACE

   INTEGER(KIND=4),PARAMETER::N=10
   REAL(KIND=8)::B(N)

   B=(/100,98,65,50,54,43,40,30,20,10/)

   CALL SORTBYKEY(B,N)

   END SUBROUTINE SORTING

END MODULE TEST

PROGRAM MAIN
   USE TEST
   IMPLICIT NONE 
   INTEGER(KIND=4)::I

   CALL SORTING <<<1,1>>> ()
   I=CUDADEVICESYNCHRONIZE()
END PROGRAM MAIN

文件2:SbyK.cu

#include <iostream>

#include <thrust/sort.h>
#include <thrust/execution_policy.h>

extern "C"
{
   __device__ void SORTBYKEY_(double *B, int N)
   {
      printf("\nB values are\n");
      for(int i=0;i<N;i++)
      {
         printf("%f\t",B[i]);
      }
     //thrust::sort(thrust::device, B, B + N);
   }
}

编译：

nvfortran -Mcuda -c SortbyKey.f95
nvcc -c SbyK.cu

链接和错误：

nvfortran -Mcuda -c++libs SbyK.o SortbyKey.o -o Sort
nvlink error   : Undefined reference to 'SORTBYKEY_' in 'SortbyKey.o'
pgacclnk: child process exit status 2: /opt/nvidia/hpc_sdk/Linux_x86_64/21.7/compilers/bin/tools/nvdd

我的目标是使用 SORTBYKEY_ 函数中的推力库对数组 B 进行操作。例如，如果我启动 1000 个线程，我想使用推力库对每个线程中的数组 B 进行排序。

【问题讨论】：

您必须启用单独的编译和设备链接才能使其工作
this 可能感兴趣

标签： cuda fortran nvcc

【解决方案1】：

通过使用带有-dc 标志的*cu 文件进行编译，我们可以启用分离编译和设备链接。

编译和链接

nvfortran -Mcuda -c SortbyKey.f95
nvcc -dc SbyK.cu
nvfortran -Mcuda -c++libs SbyK.o SortbyKey.o -o Sort

【讨论】：