混合編程[python+cpp+cuda]

---

不少時候，咱們是基於python進行模型的設計和運行，但是基於python自己的速度問題，使得原生態python代碼沒法知足生產需求，不過咱們能夠藉助其餘編程語言來緩解python開發的性能瓶頸。這裏簡單介紹個例子，以此完成如何先基於cuda編寫瓶頸函數，而後在將接口經過cpp進行封裝，最後以庫的形式被python調用。

1 cpp+python

首先，介紹下如何python調用cpp的代碼。這裏極力推薦pybind11。由於pybind11是一個輕量級，只包含頭文件的庫，他能夠在C++中調用python，或者python中調用C++代碼。其語法相似Boost.Python。但是不一樣的是Boost是一個重量級的庫，由於爲了兼容幾乎全部的C++編譯器，因此須要支持哪些最老的，bug最多的編譯器。該做者考慮到如今c11都很普及了，因此丟棄那些以前的東西，從而打造這麼一個輕量級的庫。咱們經過代碼統計：

首先是對pybind11的安裝：

git clone https://github.com/pybind/pybind11.git
cd pybind11 
mkdir build && cd build
cmake ../
make -j32

上述cmake須要3.2及以上版本。最後輸出結果以下圖所示：

這裏簡單呈現下一級目錄：

爲了實現python調用cpp，咱們先創建個文件名叫test.cpp

#include<pybind11/pybind11.h>

namespace py = pybind11;

int
add(int i, int j){
  return i+j;
}

// 該宏會在python的import語句觸發
PYBIND11_MODULE(example, m){
  m.doc() = "pybind11 example plugin";
  m.def("add", &add, "a function which adds two numbers",
        py::arg("i"), py::arg("j"));
}

而後執行：

g++  -Wall -shared -std=c++11 -fPIC \
         -I/home/zzc/software/pybind11/include \
        `cd /home/zzc/software/pybind11 && python3 -m pybind11 --includes`  \
        test.cpp \
       -o example`python3-config --extension-suffix`

結果以下圖

接下來，咱們將其改爲參數支持numpy,可參考 官網文檔； pybind11—python numpy與C++數據傳遞：

#include<pybind11/pybind11.h>
#include<pybind11/numpy.h>

namespace py = pybind11;

int
add(py::array_t<float> &array, int col){

  py::buffer_info buf1 = array.request();
  float *p = (float *)buf1.ptr;
  for (int i=0; i<col; i++){
    printf("cur value %lf\n", *p++);
  }
  return 0;
}



PYBIND11_MODULE(example, m){
  m.doc() = "pybind11 example plugin";
  m.def("add", &add, "a function which adds two numbers");
}

而後依然用上述命令編譯成so，調用結果以下圖：

更詳細的pybind11使用方法，可閱讀官方文檔

2 cuda+cpp+python

這裏只介紹如何編寫cuda的代碼，而後提供python接口。經過調查pybind11的issues：alias template error with Intel 2016.0.3 compilers,若是直接編寫cu代碼，而後一步到位，會觸發不少問題。而如這裏最後所述，較好的方式就是分開：

• 編寫cuda代碼，並生成動態連接庫；
• 編寫cpp代碼，經過函數引用方式用pybind11進行接口封裝；
• python導入對應模塊便可使用。

如上圖所示，首先，編寫cuda代碼,這裏爲了簡潔，咱們只寫一個printf

// cuda_test.cu
#include<cuda_runtime.h>
#include<stdio.h>

__global__ void
kernel(){
    printf("inside in kernel\n");
}

int
cuda(int a, int b){

   kernel<<<1,10>>>();
   cudaDeviceSynchronize();

   return 0;
}

對應頭文件：

//cuda_test.h
int cuda(int, int);

而後咱們將其用nvcc編譯成動態連接庫

nvcc --shared -Xcompiler -fPIC cuda_test.cu -o libcutest.so

結果如上圖
接着，咱們藉助pybind11，此時增長了幾行

#include<pybind11/pybind11.h>
#include"cuda_test.h" //新增的

namespace py = pybind11;

int
add(int i, int j){
  return i+j;
}

PYBIND11_MODULE(example, m){
  m.doc() = "pybind11 example plugin";
  m.def("add", &add, "a function which adds two numbers",
        py::arg("i"), py::arg("j"));
  m.def("cuda", &cuda,"testing",
        py::arg("a"), py::arg("b")); //新增的
}

而後輸入以下編譯方式：

g++  -Wall -shared -std=c++11 -fPIC \
        -L.  -lcutest \
        -I/home/zzc/software/pybind11/include \
       `cd /home/zzc/software/pybind11 && python3 -mpybind11 --includes`  \
       test.cpp \
       -o  example`python3-config --extension-suffix`

此時生成結果

而後使用