OpenCL 第5课:向量相加

OpenCL程序分为两个部份，一部份是内核代码，负责具体算法。另一部份是主程序负责初始化OpenCL和准备数据。主程序加载内核代码，并按照即定方法进行运算。

内核代码可以写在主程序里面，也可以写在另一个文本文件里，有点像DX中的HLSL和OPENGL里的GLSL。哈哈，明白意思就行了。我们用第一种方法，把代码跟源程序分开写。

调用OpenCL大至分7个步骤

1：初始化OpenCL

2：创建上下文设备

3：创建命令队列

4：创建数据缓冲区

5：将数据上传到缓冲区

6：加载编译代码，创建内核调用函数

7：设置参数，执行内核

8：读回计算结果。

下面我们通过一个向量相加的程序来了解OpenCL 。有A，B两个四维向量，相加后值存在C向量里。OpenCL会根据用户提供的维数，将向量分解成多个任务分发给多个CPU计算。

源码分两部份

（一）vecadd.cl核心代码。

?

1

2

3

4

5

6

7

8

__kernelvoidvecAdd(__global int* A,

        __globalint* B,

        __globalint* C)

{

    //获取当前工作项所在位置（线程索引号）

    intidx = get_global_id(0);

    C[idx] = A[idx] + B[idx];

}

__kernel 指明这是一个OpenCL内核，__global 说明指针指向的是全局的设备内存空间，其它的就是C语言的函数的语法。kernel必须返回空类型。

 

（二）main.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

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

117

118

119

120

121

122

123

124

125

126

127

128

129

130

131

132

133

134

135

136

137

138

139

140

141

142

143

144

145

146

147

148

149

150

151

152

153

154

155

156

157

158

159

160

161

162

163

164

165

166

167

168

#include <iostream>

#include <stdio.h>

#include <string.h>

#include <string>

#include <CL/cl.h>//包含CL的头文件

 

usingnamespace std;

 

//四维向量

#define elements 4

 

//从外部文件获取cl内核代码

boolGetFileData(constchar* fname,string& str)

{

    FILE* fp = fopen(fname,"r");

    if(fp==NULL)

    {

        printf("no found file\n");

        returnfalse;

    }

 

    intn=0;

    while(feof(fp)==0)

    {

        str += fgetc(fp);

    }

 

    returntrue;

}

 

intmain()

{

    //先读外部CL核心代码，如果失败则退出。

    //代码存buf_code里面

    string code_file;

 

    if(false== GetFileData("vecadd.cl",code_file))

        return0;

 

    char* buf_code = newchar[code_file.size()];

    strcpy(buf_code,code_file.c_str());

    buf_code[code_file.size()-1] = NULL;

 

    //声明CL所需变量。

    cl_device_id device;

    cl_platform_id platform_id = NULL;

    cl_context context;

    cl_command_queue cmdQueue;

    cl_mem bufferA,bufferB,bufferC;

    cl_program program;

    cl_kernel kernel = NULL;

 

    //我们使用的是一维向量

    //设定向量大小（维数）

    size_tglobalWorkSize[1];

    globalWorkSize[0] = elements;

 

    cl_int err;

 

    /*

        定义输入变量和输出变量，并设定初值

    */

    int* buf_A = newint[elements];

    int* buf_B = newint[elements];

    int* buf_C = newint[elements];

 

    size_tdatasize = sizeof(int) * elements;

 

    buf_A[0] = 1;

    buf_A[1] = 2;

    buf_A[2] = 3;

    buf_A[3] = 4;

 

    buf_B[0] = 5;

    buf_B[1] = 6;

    buf_B[2] = 7;

    buf_B[3] = 8;

 

    //step 1:初始化OpenCL

    err = clGetPlatformIDs(1,&platform_id,NULL);

 

    if(err!=CL_SUCCESS)

    {

        cout<<"clGetPlatformIDs error"<<endl;

        return0;

    }

 

    //这次我们只用CPU来进行并行运算，当然你也可以该成GPU

    clGetDeviceIDs(platform_id,CL_DEVICE_TYPE_CPU,1,&device,NULL);

 

    //step 2:创建上下文

    context = clCreateContext(NULL,1,&device,NULL,NULL,NULL);

 

    //step 3:创建命令队列

    cmdQueue = clCreateCommandQueue(context,device,0,NULL);

 

    //step 4:创建数据缓冲区

    bufferA = clCreateBuffer(context,

                             CL_MEM_READ_ONLY,

                             datasize,NULL,NULL);

 

    bufferB = clCreateBuffer(context,

                             CL_MEM_READ_ONLY,

                             datasize,NULL,NULL);

 

    bufferC = clCreateBuffer(context,

                             CL_MEM_WRITE_ONLY,

                             datasize,NULL,NULL);

 

    //step 5:将数据上传到缓冲区

    clEnqueueWriteBuffer(cmdQueue,

                         bufferA,CL_FALSE,

                         0,datasize,

                         buf_A,0,

                         NULL,NULL);

 

    clEnqueueWriteBuffer(cmdQueue,

                         bufferB,CL_FALSE,

                         0,datasize,

                         buf_B,0,

                         NULL,NULL);

 

    //step 6:加载编译代码,创建内核调用函数

 

    program = clCreateProgramWithSource(context,1,

                                        (constchar**)&buf_code,

                                        NULL,NULL);

 

    clBuildProgram(program,1,&device,NULL,NULL,NULL);

 

    kernel = clCreateKernel(program,"vecAdd",NULL);

 

    //step 7:设置参数，执行内核

    clSetKernelArg(kernel,0,sizeof(cl_mem),&bufferA);

    clSetKernelArg(kernel,1,sizeof(cl_mem),&bufferB);

    clSetKernelArg(kernel,2,sizeof(cl_mem),&bufferC);

 

    clEnqueueNDRangeKernel(cmdQueue,kernel,

                           1,NULL,

                           globalWorkSize,

                           NULL,0,NULL,NULL);

 

    //step 8:取回计算结果

    clEnqueueReadBuffer(cmdQueue,bufferC,CL_TRUE,0,

                        datasize,buf_C,0,NULL,NULL);

 

    //输出计算结果

    cout<<"["<<buf_A[0]<<","<<buf_A[1]<<","<<buf_A[2]<<","<<buf_A[3]<<"]+["

    <<buf_B[0]<<","<<buf_B[1]<<","<<buf_B[2]<<","<<buf_B[3]<<"]=["

    <<buf_C[0]<<","<<buf_C[1]<<","<<buf_C[2]<<","<<buf_C[3]<<"]"<<endl;

 

    //释放所有调用和内存

 

    clReleaseKernel(kernel);

    clReleaseProgram(program);

    clReleaseCommandQueue(cmdQueue);

    clReleaseMemObject(bufferA);

    clReleaseMemObject(bufferB);

    clReleaseMemObject(bufferC);

    clReleaseContext(context);

 

    deletebuf_A;

    deletebuf_B;

    deletebuf_C;

    deletebuf_code;

 

    return0;

}

运算结果：

 [1,2,3,4] + [5,6,7,8] = [6,8,10,12]