AMD OpenCL大学课程5OpenCL概述 续篇.docx

AMD OpenCL大学课程5OpenCL概述 续篇.docx

《AMD OpenCL大学课程5OpenCL概述 续篇.docx》由会员分享，可在线阅读，更多相关《AMD OpenCL大学课程5OpenCL概述 续篇.docx（11页珍藏版）》请在冰豆网上搜索。

AMDOpenCL大学课程5OpenCL概述续篇

AMDOpenCL大学课程（5）:

OpenCL概述续篇

OpenCL内存模型

OpenCL的内存模型定义了各种各样内存类型，各种内存模型之间有层级关系。

各种内存之间的数据传输必须是显式进行的，比如从hostmemory到devicememory，从globalmemory到localmemory等等。

WorkGroup被映射到硬件的CU上执行（在AMD5xxx系列显卡上，CU就是simd，一个simd中有16个pe，或者说是streamcore），OpenCL并不提供各个workgroup之间的一致性，如果我们需要在各个workgroup之间共享数据或者通信之类的，要自己通过软件实现。

Kernel函数的写法

每个线程（workitem）都有一个kenerl函数的实例。

下面我们看下kernel的写法：

1:

__kernelvoidvecadd（__globalconstfloat*A,__globalconstfloat*B,__globalfloat*C）

2:

{

3:

intid=get_global_id（0）;

4:

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

5:

}

每个Kernel函数都必须以__kernel开始，而且必须返回void。

每个输入参数都必须声明使用的内存类型。

通过一些API，比如get_global_id之类的得到线程id。

内存对象地址空间标识符有以下几种：

__global–memoryallocatedfromglobaladdressspace

__constant–aspecialtypeofread-onlymemory

__local–memorysharedbyawork-group

__private–privateperwork-itemmemory

__read_only/__write_only–usedforimages

Kernel函数参数如果是内存对象，那么一定是__global,__local或者constant。

运行Kernel

首先要设置线程索引空间的维数以及workgroup大小等。

我们通过函数clEnqueueNDRangeKerne把Kernel放在一个队列里，但不保证它马上执行，OpenCLdriver会管理队列，调度Kernel的执行。

注意：

每个线程执行的代码都是相同的，但是它们执行数据却是不同的。

该函数把要执行的Kernel函数放在指定的命令队列中，globald大小（线程索引空间）必须指定，local大小（workgroup）可以指定，也可以为空。

如果为空，则系统会自动根据硬件选择合适的大小。

event_wait_list用来选定一些events，只有这些events执行完后，该kernel才可能被执行，也就是通过事件机制来实现不同kernel函数之间的同步。

当Kernel函数执行完毕后，我们要把数据从devicememory中拷贝到hostmemory中去。

释放资源：

大多数的OpenCL资源都是指针，不使用的时候需要释放掉。

当然，程序关闭的时候这些对象也会被自动释放掉。

释放资源的函数是：

clRelase{Resource}，比如:

clReleaseProgram（）,clReleaseMemObject（）等。

错误捕捉：

如果OpenCL函数执行失败，会返回一个错误码，一般是个负值，返回0则表示执行成功。

我们可以根据该错误码知道什么地方出错了，需要修改。

错误码在cl.h中定义，下面是几个错误码的例子.

CL_DEVICE_NOT_FOUND-1

CL_DEVICE_NOT_AVAILABLE-2

CL_COMPILER_NOT_AVAILABLE-3

CL_MEM_OBJECT_ALLOCATION_FAILURE-4

…

下面是一个OpenCL机制的示意图

程序模型

数据并行：

workitem和内存对象元素之间是一一映射关系；workgroup可以显示指定，也可以隐式指定。

任务并行：

kernel的执行独立于线程索引空间；用其他方法表示并行，比如把不同的任务放入队列，用设备指定的特殊的向量类型等等。

同步：

workgroup内workitem之间的同步；命令队列中不同命令之间的同步。

完整代码如下：

1:

#include"stdafx.h"

2:

#include

3:

#include

4:

#include

5:

#include

6:

#include

7:

#include

8:

9:

usingnamespacestd;

10:

#defineNWITEMS262144

11:

12:

#pragmacomment（lib,"OpenCL.lib"）

13:

14:

//把文本文件读入一个string中

15:

intconvertToString（constchar*filename,std:

:

string&s）

16:

{

17:

size_tsize;

18:

char*str;

19:

20:

std:

:

fstreamf（filename,（std:

:

fstream:

:

in|std:

:

fstream:

:

binary））;

21:

22:

if（f.is_open（））

23:

{

24:

size_tfileSize;

25:

f.seekg（0,std:

:

fstream:

:

end）;

26:

size=fileSize=（size_t）f.tellg（）;

27:

f.seekg（0,std:

:

fstream:

:

beg）;

28:

29:

str=newchar[size+1];

30:

if（!

str）

31:

{

32:

f.close（）;

33:

returnNULL;

34:

}

35:

36:

f.read（str,fileSize）;

37:

f.close（）;

38:

str[size]='\0';

39:

40:

s=str;

41:

delete[]str;

42:

return0;

43:

}

44:

printf（"Error:

Failedtoopenfile%s\n",filename）;

45:

return1;

46:

}

47:

48:

intmain（intargc,char*argv[]）

49:

{

50:

//在host内存中创建三个缓冲区

51:

float*buf1=0;

52:

float*buf2=0;

53:

float*buf=0;

54:

55:

buf1=（float*）malloc（NWITEMS*sizeof（float））;

56:

buf2=（float*）malloc（NWITEMS*sizeof（float））;

57:

buf=（float*）malloc（NWITEMS*sizeof（float））;

58:

59:

//初始化buf1和buf2的内容

60:

inti;

61:

srand（（unsigned）time（NULL））;

62:

for（i=0;i

63:

buf1[i]=rand（）%65535;

64:

65:

srand（（unsigned）time（NULL）+1000）;

66:

for（i=0;i

67:

buf2[i]=rand（）%65535;

68:

69:

for（i=0;i

70:

buf[i]=buf1[i]+buf2[i];

71:

72:

cl_uintstatus;

73:

cl_platform_idplatform;

74:

75:

//创建平台对象

76:

status=clGetPlatformIDs（1,&platform,NULL）;

77:

78:

cl_device_iddevice;

79:

80:

//创建GPU设备

81:

clGetDeviceIDs（platform,CL_DEVICE_TYPE_GPU,

82:

1,

83:

&device,

84:

NULL）;

85:

//创建context

86:

cl_contextcontext=clCreateContext（NULL,

87:

1,

88:

&device,

89:

NULL,NULL,NULL）;

90:

//创建命令队列

91:

cl_command_queuequeue=clCreateCommandQueue（context,

92:

device,

93:

CL_QUEUE_PROFILING_ENABLE,NULL）;

94:

//创建三个OpenCL内存对象，并把buf1的内容通过隐式拷贝的方式

95:

//拷贝到clbuf1,buf2的内容通过显示拷贝的方式拷贝到clbuf2

96:

cl_memclbuf1=clCreateBuffer（context,

97:

CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,

98:

NWITEMS*sizeof（cl_float）,buf1,

99:

NULL）;

100:

101:

cl_memclbuf2=clCreateBuffer（context,

102:

CL_MEM_READ_ONLY,

103:

NWITEMS*sizeof（cl_float）,NULL,

104:

NULL）;

105:

106:

status=clEnqueueWriteBuffer（queue,clbuf2,1,

107:

0,NWITEMS*sizeof（cl_float）,buf2,0,0,0）;

108:

109:

cl_membuffer=clCreateBuffer（context,

110:

CL_MEM_WRITE_ONLY,

111:

NWITEMS*sizeof（cl_float）,

112:

NULL,NULL）;

113:

114:

constchar*filename="add.cl";

115:

std:

:

stringsourceStr;

116:

status=convertToString（filename,sourceStr）;

117:

constchar*source=sourceStr.c_str（）;

118:

size_tsourceSize[]={strlen（source）};

119:

120:

//创建程序对象

121:

cl_programprogram=clCreateProgramWithSource（

122:

context,

123:

1,

124:

&source,

125:

sourceSize,

126:

NULL）;

127:

//编译程序对象

128:

status=clBuildProgram（program,1,&device,NULL,NULL,NULL）;

129:

if（status!

=0）

130:

{

131:

printf（"clBuildfailed:

%d\n",status）;

132:

chartbuf[0x10000];

133:

clGetProgramBuildInfo（program,device,CL_PROGRAM_BUILD_LOG,0x10000,tbuf,NULL）;

134:

printf（"\n%s\n",tbuf）;

135:

return-1;

136:

}

137:

138:

//创建Kernel对象

139:

cl_kernelkernel=clCreateKernel（program,"vecadd",NULL）;

140:

//设置Kernel参数

141:

cl_intclnum=NWITEMS;

142:

clSetKernelArg（kernel,0,sizeof（cl_mem）,（void*）&clbuf1）;

143:

clSetKernelArg（kernel,1,sizeof（cl_mem）,（void*）&clbuf2）;

144:

clSetKernelArg（kernel,2,sizeof（cl_mem）,（void*）&buffer）;

145:

146:

//执行kernel

147:

cl_eventev;

148:

size_tglobal_work_size=NWITEMS;

149:

clEnqueueNDRangeKernel（queue,

150:

kernel,

151:

1,

152:

NULL,

153:

&global_work_size,

154:

NULL,0,NULL,&ev）;

155:

clFinish（queue）;

156:

157:

//数据拷回host内存

158:

cl_float*ptr;

159:

ptr=（cl_float*）clEnqueueMapBuffer（queue,

160:

buffer,

161:

CL_TRUE,

162:

CL_MAP_READ,

163:

0,

164:

NWITEMS*sizeof（cl_float）,

165:

0,NULL,NULL,NULL）;

166:

//结果验证，和cpu计算的结果比较

167:

if（!

memcmp（buf,ptr,NWITEMS））

168:

printf（"Verifypassed\n"）;

169:

elseprintf（"verifyfailed"）;

170:

171:

if（buf）

172:

free（buf）;

173:

if（buf1）

174:

free（buf1）;

175:

if（buf2）

176:

free（buf2）;

177:

178:

//删除OpenCL资源对象

179:

clReleaseMemObject（clbuf1）;

180:

clReleaseMemObject（clbuf2）;

181:

clReleaseMemObject（buffer）;

182:

clReleaseProgram（program）;

183:

clReleaseCommandQueue（queue）;

184:

clReleaseContext（context）;

185:

return0;

186:

}

187:

也可以在