opencl内存模型的几个问题

[losywar]

1. global、group、loacl、item、之间的关系，以及它们的使用方法
2. A10（APU）处理器传说统一了CPU和GPu的地址空间，opencl1.2里有没有这方面的sample，知道怎么使用么
3. 关于内存模型方面有没有更加深入理解的材料（网上的那些教程就算了吧）。
谢谢~~

[DigBug]

引用 楼主 losywar 的回复:

1. global、group、loacl、item、之间的关系，以及它们的使用方法 2. A10（APU）处理器传说统一了CPU和GPu的地址空间，opencl1.2里有没有这方面的sample，知道怎么使用么 3. 关于内存模型方面有没有更加深入理解的材料（网上的那些教程就算了吧）。 谢谢~~

1, global, local是内存模型的概念，对应于类比于CPU中，L2全局Cache与L1每个核的Cache；主要区别在访问权限例如是否可被全部的GPU核心所访问，及访问速度；group与item是线程模型的概念，一般而言一个item运行在一个GPU核心上，多个item可以组织成一个group；如果说他们之间的关系，可以这么说：global能被所有的item访问；local只能被一个group中的item访问（如果没有记错的话） 2，目前所有的APU，包括代号Llano和Trinity的A4、A6、A8、A10，都是将CPU与GPU进行物理整合，内存虚址空间仍然是分离的；代号Kaveri的APU会实现统一的CPU和GPU虚址空间；GPU和CPU指针可以互访，其实说白了就是没有内存和显存的区别，大家随便访问；

[outstander]

1.global、group、loacl是线程编号的空间，是逻辑概念。work-item是一个物理硬件概念，一个work-item在同一时刻只能运行一个线程，但可以切换到其他线程执行。 2.我没用过A10,暂不能给出统一地址空间的用法，等待其他大牛回答。 3.请参看AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide,该文档由AMD自己编写，主要深入介绍了AMD自己的硬件平台是如何支持opencl标准的。对内存模型有很深入的介绍。

[o_o_lll]

很深奥，一点都不懂

[masol]

1.1 global,local是相对于你的kernel而言的。这么说吧，迄今为止的opencl规范，你没有办法从宿主代码中初始化local内存，换言之，global内存可以在host中操作(clCreateBuffer,clEnqueueReadBuffer这些函数都是针对global的——kernel中使用global修饰符的那批变量)。clSetKernelArg如果设置local，只能设置尺寸，而不能设置初始值。一旦进入kernel视角，global,local没啥区别——前提是各个线程之间无关。如果有关系，建议你还是全部使用global,否则你需要使用async_work_group_copy来同步内存，不这样做，多核gpu环境下会有问题。使用global或local对性能不会有什么大的影响。对性能影响最大的是work-group size. 1.2 item,group的概念清晰你先需要了解GPU的计算模式。各家GPU的实现是完全不同的，opencl为了统一编程模型，给出了work-group size的概念。就是GPU能同时执行的最大单元——这些单元GPU会调度其在一个物理核心上同步执行，你可以通过查询MAX_WORK_ITEM_SIZES来获取当前硬件支持的最大数值，你的workgroup(总维度)不能超越这个值。这个值是由硬件实现的，合理的选择通常让你的性能差出一个数量级来，通常这个值不会小于32,而且你通常应该最少选择到32而不是让opencl自由选择.（ps: 这个值不是越大越好的，取决于硬件的实现，去看看各家指令发射如何实现的就会清晰为什么不是越大越好)。另外一方面，从程序角度，你通常考虑的是需要的任务数，例如你渲染一张1000X1000的图片，每个点单独计算，任务数1百万最好，这个任务数通常就是你的global work-size了，这个倒是没什么限制，只受到GPU内存分配的限制，但是你应该确保global work-size是local work-size的倍数，否则没法调度了——差异处理至少现阶段你必须在kernel中自行处理。不过global也不是越高越好的，现阶段的opencl只能发送任务，不能终止，去看下clEnqueueNDRangeKernel就清楚了，任务多了，中途无法执行宿主干预，因此，你需要一个合理的值，以便可以中途修正opencl代码——例如重新生成新的kernel以执行。 2. 内存模型和你有关系么？opencl的接口是固定的，你和kernel通信的办法就是通过opencl接口维护buffer,至于这个buffer是显存还是内存和你有啥关系？除非你要改变opencl规范，这个话题完全不需要考虑。其实现在不仅仅A10,CPU集成的GPU都是共享的宿主内存，你当然可以直接访问这部分内存——获取gdt，查到相应项去访问就可以了，还是上面的问题，有必要么？

[masol]

好几年没来，今天来转转，看到分高进来了，顺手写点，结果发现是12年的帖子，悲剧......