R 是一個統計學經常用到的軟體，提供了非常多的統計學函數。 但是它是一個單線程解釋語言，面對大數據量的時候，往往性能跟不上，可以利用 Rcpp 編寫 C++ 包提供給 R 使用，可以大大提高性能。 而對於大規模數據的處理，使用 CUDA 則是一個非常好的解決方案。 在 Linux 和 macOS 下， CUDA 程序和 C++ 程序都使用 gcc 編譯器， 但是在 Windows 下，Rcpp 的包必須用 MinGW 編譯器， CUDA 的包必須用 MSVC 編譯器，需要一定的技巧才能讓 R 用上 CUDA。
本文介紹 MSVC 包和 MinGW 包的混合編譯，不僅適用於 R 語言，也不僅適用於 CUDA 程序， 也適用於其他需要通過 MinGW 的程序調用 MSVC 程序的情況。
MinGW 調用 MSVC 庫函數的條件MinGW 其實一直都是可以直接調用 MSVC 庫函數的，只是這樣的函數需要滿足幾個條件：
MinGW 可以調用 MSVC 編譯的動態庫，但是不能調用 MSVC 編譯的靜態庫， 因為 MinGW 和 MSVC 中會引用同樣的符號，但 MSVC 有的符號在 MinGW 中沒有，調用 MSVC 靜態庫時需要加載 MSVC 的符號，導致衝突。
MSVC 編譯的 DLL 必須導出 C 接口，否則 MinGW 中找不到符號，這是因為 C++ 會給函數名做修改以支持函數重載，但 MSVC 和 MinGW 對函數修改的方式不一樣。 也有人說時 __cdecl__ 和 __stdcall__ 的問題，但我試了一下不太行，還是找不到符號。 既然時 C 接口，那麼參數和返回值不能是 class ，只能用指向 class 類型的指針類型。
不能將 MinGW 中創建的指針傳遞到 MSVC 中進行操作，否則在 MSVC 中就相當於野指針。反之也不可以。 因為 MinGW 和 MSVC 編譯的庫，內存地址是兩套，指針不互通。
在 R 中，一定會大量用到矩陣，一有矩陣那必然涉及到指針，而且一定是在 MinGW 函數中創建的指針。那麼該怎麼將矩陣傳到 MSVC 的函數中呢？ 這其實非常類似於 CUDA 編程中的內存問題，我們只需要在兩邊分別開闢內存，然後將內存中的數據複製一下，相當於寫一個 cudaMalloc 和 cudaMemcpy 。 這樣相當於在全局創建了很多的變量，如果使用一些方法將這些全局變量統一管理會更好。 可以將這些全局變量保存在一個結構體中，工廠函數返回一個指向這個結構體的指針，並為這個結構體成員創建初始值。
也可以利用 C++ 類的封裝我們要在 R 中調用的函數，將這個類繼承自一個抽象類（所有函數都是純虛函數），將接口類導出， 同時導出一個 C 的工廠函數，返回指向這個接口類的指針， 利用 MSVC 和 MinGW 虛表結構一致的特點，就可以在 MinGW 的 C++ 代碼中使用這個接口類中的函數了（類沒有虛析構函數）。 將矩陣數據作為類的成員變量，利用 C++ 成員函數進行創建、賦值、銷毀。
非類的寫法首先演示一種純 C 非類的寫法。首先需要建立一個 VS 的 CUDA 工程，並設置該項目配置類型為「動態連結庫」。 項目目錄如下：
AddCUDA
add.cpp
add.h
kernel.cu
kernel.h
AddCUDA.vcxproj
文件 kernel.cu 使用的 CUDA 函數是 VS 中 CUDA 工程自帶的模板，做了一些修改， 主要去掉了 goto 語句，並設置了核函數啟動配置。
// kernel.cu#include "cuda_runtime.h"#include "device_launch_parameters.h"
#include <stdio.h>
#include "kernel.h"
__global__ void addKernel(int *c, const int *a, const int *b){int i = threadIdx.x; c[i] = a[i] + b[i];}
// Helper function for using CUDA to add vectors in parallel.bool addWithCuda(int *c, const int *a, const int *b, unsigned int size){int *dev_a = 0;int *dev_b = 0;int *dev_c = 0; cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }
// Allocate GPU buffers for three vectors (two input, one output) . cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }// Copy input vectors from host memory to GPU buffers. cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }
// Launch a kernel on the GPU with one thread for each element.dim3 blockSize(256), gridSize((size + blockSize.x - 1) / blockSize.x); addKernel<<<gridSize, blockSize >>>(dev_c, dev_a, dev_b);// Check for any errors launching the kernel cudaStatus = cudaGetLastError();if (cudaStatus != cudaSuccess) {fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus)); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }
// cudaDeviceSynchronize waits for the kernel to finish, and returns// any errors encountered during the launch. cudaStatus = cudaDeviceSynchronize();if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!"); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b);return false; }
return true;}將 addWithCuda 函數的聲明移動到 kernel.h 文件中。
// kernel.h#pragma once#include <cuda_runtime.h>bool addWithCuda(int *c, const int *a, const int *b, unsigned int size);然後需要編寫 MSVC 的接口函數了。在 add.h 文件中做如下聲明：
// add.h#pragma once
#ifdef DLL_EXPORT#define ADDCUDA_API extern "C" __declspec(dllexport)#else#define ADDCUDA_API extern "C" __declspec(dllimport)#endif // DLL_EXPORT
ADDCUDA_API int* createVector(int n);ADDCUDA_API void setVector(int* ptr, int i, int value);ADDCUDA_API int getVector(int* ptr, int i);ADDCUDA_API void deleteVector(int* ptr);ADDCUDA_API bool addVector(int* a, int* b, int n, int *c);當然， __declspec(dllimport) 可要可不要，要也可以，這樣這個 Dll 也可以給 Windows 程序使用。
在 add.cpp 中進行定義：
// add.cpp#include "add.h"#include "kernel.h"
int* createVector(int n){return new int[n];}
void deleteVector(int* ptr){delete[] ptr;}
void setVector(int* ptr, int i, int value){ ptr[i] = value;}
int getVector(int* ptr, int i){return ptr[i];}
bool addVector(int* a, int* b, int n, int *c){bool cudaStatus = addWithCuda(c, a, b, n);return cudaStatus;}雖然這幾個函數很短小，但是也不能寫在頭文件裡。否則 MinGW 中會報符號二義性錯誤。然後在 MinGW 的主函數中進行調用
// mingw.cpp#include "add.h"#include <stdio.h>
int main(int argc, char const *argv[]){int n = 100000;int *a = createVector(n);int *b = createVector(n);int *c = createVector(n);for (size_t i = 0; i < n; i++) { setVector(a, i, 10); setVector(b, i, 100); setVector(c, i, 0); } addVector(a, b, n, c);int *result = new int[n];for (size_t i = 0; i < n; i++) { result[i] = getVector(c, i); }printf("result:\n");for (size_t i = 0; i < 10; i++) {printf("%5d", result[i]); }printf("\n");
deleteVector(a); deleteVector(b); deleteVector(c);
return 0;}MinGW 在連結時，需要手動指定 MSVC 生成的 lib 文件，而且要放到 -o 參數的後面，方法如下：
g++ -I"./AddCUDA" -L"./x64/Release" mingw.o -o cudaMinGWC -lAddCUDA這樣就生成了 cudaMinGWC.exe 文件。運行一下可以得到結果：
result: 110 110 110 110 110 110 110 110 110 110抽象類的寫法抽象類的寫法主要用到了多態的特性。首先需要創建一個抽象基類 IAdd 和派生類 CAdd 。
// IAdd.h#pragma once
#ifdef DLL_EXPORT#define ADDCUDA_API __declspec(dllexport)#else#define ADDCUDA_API __declspec(dllimport)#endif // DLL_EXPORT
class ADDCUDA_API IAdd{public:virtual void SetA(int i, int value) = 0;virtual void SetB(int i, int value) = 0;virtual int GetC(int i) = 0;virtual bool Add() = 0;};
extern "C" ADDCUDA_API IAdd* Add_new(int n);extern "C" ADDCUDA_API void Add_del(IAdd* ptr);
// IAdd.cpp#include "IAdd.h"#include "CAdd.h"
IAdd* Add_new(int n){return new CAdd(n);}
void Add_del(IAdd* ptr){delete ptr;}在抽象類中完全不實現類的任何接口，都標記為純虛函數。 同時，在抽象類的外面，定義一套工廠函數，用於創建和銷毀這個抽象類派生類的對象。
派生類的定義如下：
// CAdd.h#pragma once
#include "IAdd.h"
class CAdd : public IAdd{private:int n;int* a;int* b;int* c;public: CAdd(int n); ~CAdd();
virtual void SetA(int i, int value);virtual void SetB(int i, int value);virtual int GetC(int i);virtual bool Add();};// CAdd.cpp#include "CAdd.h"#include "kernel.h"#include <memory.h>
CAdd::CAdd(int n){this->n = n; a = new int[n]; b = new int[n]; c = new int[n];memset(a, 0, sizeof(int) * n);memset(b, 0, sizeof(int) * n);memset(c, 0, sizeof(int) * n);}CAdd::~CAdd(){delete[] a;delete[] b;delete[] c;}
void CAdd::SetA(int i, int value){if (i < n) a[i] = value;}
void CAdd::SetB(int i, int value){if (i < n) b[i] = value;}
int CAdd::GetC(int i){return (i < n) ? c[i] : 0;}
bool CAdd::Add(){return addWithCuda(c, a, b, n);}在 MinGW 中調用如下：
#include <stdio.h>#include "IAdd.h"int main(int argc, char const *argv[]){int n = 100000; IAdd* ptr = Add_new(n);for (size_t i = 0; i < n; i++) { ptr->SetA(i, 10); ptr->SetB(i, 100); } ptr->Add();printf("result:\n");for (size_t i = 0; i < 10; i++) {printf("%5d", ptr->GetC(i)); }printf("\n");
Add_del(ptr);
return 0;}程序運行結果：
result:110 110 110 110 110 110 110 110 110 110可見已經可以運行了。
總結總體而言，這種方式調用方式的開銷還是比較大的。 不僅在內存中複製了一份數據，在傳遞數據的過程中是一個一個傳遞的，比直接內存拷貝開銷大很多。 另外如果採用抽象類的方式，還有虛函數調用的開銷。 因此，如果不是必須在 Windows 上用 MinGW 調用 CUDA 程序，儘量還是使用 MSVC 編譯器。
對於 Rcpp 而言，恰恰是必須在 Windows 使用 MinGW ，此使想調用 CUDA 程序，則需要通過這種方式。