让 Windows 的 R 用上 CUDA

让 Windows 的 R 用上 CUDA

  • 2019-08-12

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 语句，并设置了核函数启动配置。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
// 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 文件中。

1
2
3
4
// kernel.h
#pragma once
#include <cuda_runtime.h>
bool addWithCuda(int *c, const int *a, const int *b, unsigned int size);

然后需要编写 MSVC 的接口函数了。在 add.h 文件中做如下声明：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
// 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 中进行定义：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
// 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 的主函数中进行调用

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
// 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 参数的后面，方法如下：

1
g++ -I"./AddCUDA" -L"./x64/Release" mingw.o -o cudaMinGWC -lAddCUDA

这样就生成了 cudaMinGWC.exe 文件。运行一下可以得到结果：

1
2
result:  
  110  110  110  110  110  110  110  110  110  110

抽象类的写法

抽象类的写法主要用到了多态的特性。首先需要创建一个抽象基类 IAdd 和派生类 CAdd 。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
// 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;
}

在抽象类中完全不实现类的任何接口，都标记为纯虚函数。 同时，在抽象类的外面，定义一套工厂函数，用于创建和销毁这个抽象类派生类的对象。

派生类的定义如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
// 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 中调用如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
#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;
}

程序运行结果：

1
2
result:
  110  110  110  110  110  110  110  110  110  110

可见已经可以运行了。

总结

总体而言，这种方式调用方式的开销还是比较大的。 不仅在内存中复制了一份数据，在传递数据的过程中是一个一个传递的，比直接内存拷贝开销大很多。 另外如果采用抽象类的方式，还有虚函数调用的开销。 因此，如果不是必须在 Windows 上用 MinGW 调用 CUDA 程序，尽量还是使用 MSVC 编译器。

对于 Rcpp 而言，恰恰是必须在 Windows 使用 MinGW ，此使想调用 CUDA 程序，则需要通过这种方式。

本文所涉及代码已发布在 GitHub 上。

感谢您的阅读，本文由 HPDell 的个人博客 版权所有。如若转载，请注明出处：HPDell 的个人博客（http://hpdell.github.io/编程/windows-r-cuda/）
HTML5 播放 RTSP 视频
探索性因子分析