混合语言CUDA编程

Question

如何将CUDA代码与Fortran和C代码(混合语言编程)结合起来。Fortran代码调用C函数，然后调用CUDA内核。例如。

Fortran函数：

if(flag.eq.1) call c_func

C职能：

void c_func()
{
  /* copy data to device 
  ....
  cuda_kernel<<< kernel parameters>>>();

  /* copy data from device to Host
  ........
}

编译这种类型的代码的方法是什么？

Answer 1

我相信对此有很多可能的方法。但是，按照您给出的示例，它应该是相当简单的。

该任务可分为两部分：

1. 如何从Fortran调用C函数
2. 如何从C调用CUDA函数

我认为你的问题可能是围绕第一件，因此，它并不是真正的数据自动化系统-具体。当然，对于第二部分，在cuda标记上有很多示例，还有cuda样本码和程序编制指南。

一种可能有助于简化第一部分的方法是使用ISO_C_BINDING 内建模块，它内置于许多当前的fortran发行版中。这个模块定义了许多类型，这些类型对于在C和Fortran之间传递数据非常有用。

然后，您可以创建一个INTERFACE块来定义您希望从fortran调用的C函数的参数。下面是一个有用的示例：

$ cat cuda_test.f90
!=======================================================================================================================
!Interface to cuda C functions
!=======================================================================================================================
module cuda_test

  use iso_c_binding

  interface
     !
     integer(c_int) function cudatestfunc(idata, isize) bind(C, name="cudatestfunc")
       use iso_c_binding
       implicit none
       type(c_ptr),value :: idata
       integer(c_int),value :: isize
     end function cudatestfunc
     !
  end interface

end module cuda_test



!=======================================================================================================================
program main
!=======================================================================================================================

  use iso_c_binding

  use cuda_test

  type(c_ptr) :: mydata
  integer*4, target   :: mysize,myresult
  integer*4,dimension(:),allocatable,target :: darray
  mysize = 100
  allocate(darray(mysize))
  darray = (/ (1, I = 1, mysize) /)
  mydata = c_loc(darray)
  myresult = cudatestfunc(mydata, mysize)

  write (*, '(A, I10)') "  result: ", myresult
  write (*,*)

end program main
$ cat cuda_test.cu
#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void testkernel(int *data, int size){

  for (int i = 1; i < size; i++) data[0] += data[i];
}
extern "C" {
int cudatestfunc(int *data, int size){

  int *d_data;
  cudaMalloc(&d_data, size*sizeof(int));
  cudaMemcpy(d_data, data, size*sizeof(int), cudaMemcpyHostToDevice);
  testkernel<<<1,1>>>(d_data, size);
  int result;
  cudaMemcpy(&result, d_data, sizeof(int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cuda error");
  return result;
}
}
$ gfortran -c cuda_test.f90 -o cuda_testf.o                 
$ nvcc -c cuda_test.cu -o cuda_testc.o                      
$ gfortran cuda_testc.o cuda_testf.o -o cuda_test -L/usr/local/cuda/lib64 -lcudart -lstdc++
$ ./cuda_test
  result:        100

$

(在RHEL 6.2、GNU 4.4.7、CUDA 7.0上测试)

附注/其他备选方案：

1. 如果您只需要调用CUDA运行时API函数，则可以直接从fortran调用这些函数，而不需要任何C/C++文件(如果您创建自己的绑定)。例如这里。
2. 如果您只需要调用CUSPARSE或CUBLAS库函数，则会为您创建一些绑定，这些绑定都包含在CUDA发行版中。默认情况下，这些绑定在linux上安装在/usr/local/cuda/src上。cublas绑定的工作示例包含在cublas文档中。
3. 如果您需要直接从fortran调用其他CUDA库函数，那么创建您自己的绑定并不是非常困难。对于CUSOLVER中的一组简单操作来说，一个有效的示例是这里。
4. 您还可以直接编写库达·福特兰代码。(这里就是一个例子。)这需要来自PGI的CUDA编译器。
5. 您还可以编写OpenACC Fortran代码。这需要一个可用的OpenACC编译器，例如来自PGI的编译器。这里提供免费学术使用或试用版本。