关于OpenCL 的 work_item 的数量

[hehykissool]

本来Max_work_item[0] = 256;
Max_work_item[1] = 256;
Max_work_item[2] = 256;

但是可以设置
global_work_item[3] = {5000,80,80}
仍能运行
但是继续增加就会无法运行，开始运行后很久都没有计算结果
求问这是什么原因

以及，若我需要运行的item很多很多，应该从哪里下手

[hqk2012]

没有遇到过这种问题

[卞亚涛]

1，Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256;指的是一个workgroup内不同维度的#item最大值 global_work_item的大小是 workgroup大小与workgroup个数的乘积，所以可以这样设置 2，虽然有规定每个维度的最大值，但是可用的最大#group以及#item最终受硬件限制。有时候虽然你可以设置很多item，但是这时候过多的item（若已经超过资源限制）于性能提升无益处，因为此时系统是先执行一部分，再执行另一部分。 3，运行很久没有结果，你可以认真分析一下数据相关性等等

引用 楼主 hehykissool 的回复:

本来Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256; 但是可以设置 global_work_item[3] = {5000,80,80} 仍能运行 但是继续增加就会无法运行，开始运行后很久都没有计算结果 求问这是什么原因 以及，若我需要运行的item很多很多，应该从哪里下手 ……

[outstander]

引用 楼主 hehykissool 的回复:

本来Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256; 但是可以设置 global_work_item[3] = {5000,80,80} 仍能运行 但是继续增加就会无法运行，开始运行后很久都没有计算结果 求问这是什么原因 以及，若我需要运行的item很多很多，应该从哪里下手 ……

Max数组是什么?怎么得到的？ 看数据是workgroup 大小的限制

[hehykissool]

引用 3 楼 outstander 的回复:

引用 楼主 hehykissool 的回复:本来Max_work_item[0] = 256;
Max_work_item[1] = 256;
Max_work_item[2] = 256;

但是可以设置
global_work_item[3] = {5000,80,80}
仍能运行
但是继续增加就会无法运行，开始运行后很久都没有计算结果
求问这是什么原……

在clinfo中写的
Microsoft Windows [版本 6.1.7600]
版权所有 (c) 2009 Microsoft Corporation。保留所有权利。

C:\Users\hehykissool>clinfo
Number of platforms: 1
Platform Profile: FULL_PROFILE
Platform Version: OpenCL 1.2 AMD-APP (923.1)
Platform Name: AMD Accelerated Parallel Proces
sing
Platform Vendor: Advanced Micro Devices, Inc.
Platform Extensions: cl_khr_icd cl_amd_event_callbac
k cl_amd_offline_devices cl_khr_d3d10_sharing


Platform Name: AMD Accelerated Parallel Proces
sing
Number of devices: 2
Device Type: CL_DEVICE_TYPE_GPU
Device ID: 4098
Max compute units: 5
Max work items dimensions: 3
Max work items[0]: 256
Max work items[1]: 256
Max work items[2]: 256
Max work group size: 256
Preferred vector width char: 16
Preferred vector width short: 8
Preferred vector width int: 4
Preferred vector width long: 2
Preferred vector width float: 4
Preferred vector width double: 0
Native vector width char: 16
Native vector width short: 8
Native vector width int: 4
Native vector width long: 2
Native vector width float: 4
Native vector width double: 0
Max clock frequency: 550Mhz
Address bits: 32
Max memory allocation: 536870912
Image support: Yes
Max number of images read arguments: 128
Max number of images write arguments: 8
Max image 2D width: 8192
Max image 2D height: 8192
Max image 3D width: 2048
Max image 3D height: 2048
Max image 3D depth: 2048
Max samplers within kernel: 16
Max size of kernel argument: 1024
Alignment (bits) of base address: 2048
Minimum alignment (bytes) for any datatype: 128
Single precision floating point capability
Denorms: No
Quiet NaNs: Yes
Round to nearest even: Yes
Round to zero: Yes
Round to +ve and infinity: Yes
IEEE754-2008 fused multiply-add: Yes
Cache type: None
Cache line size: 0
Cache size: 0
Global memory size: 1073741824
Constant buffer size: 65536
Max number of constant args: 8
Local memory type: Scratchpad
Local memory size: 32768
Kernel Preferred work group size multiple: 64
Error correction support: 0
Unified memory for Host and Device: 0
Profiling timer resolution: 1
Device endianess: Little
Available: Yes
Compiler available: Yes
Execution capabilities:
Execute OpenCL kernels: Yes
Execute native function: No
Queue properties:
Out-of-Order: No
Profiling : Yes
Platform ID: 000007FEE85D2A08
Name: Redwood
Vendor: Advanced Micro Devices, Inc.
Driver version: CAL 1.4.1385 (VM)
Profile: FULL_PROFILE
Version: OpenCL 1.2 AMD-APP (923.1)
Extensions: cl_khr_global_int32_base_atomic
s cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_lo
cal_int32_extended_atomics cl_khr_3d_image_writes cl_khr_byte_addressable_store
cl_khr_gl_sharing cl_ext_atomic_counters_32 cl_amd_device_attribute_query cl_amd
_vec3 cl_amd_printf cl_amd_media_ops cl_amd_popcnt cl_khr_d3d10_sharing

[hehykissool]

引用 2 楼 bianjunlan 的回复:

1，Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256;指的是一个workgroup内不同维度的#item最大值 global_work_item的大小是 workgroup大小与workgroup个数的乘积，所以可以这样设置 2，虽然有规定每个维度的最大值，但是可用……

怎样去寻找设备的运行上限呢，或者只能靠实验测试？

[卞亚涛]

引用 5 楼 hehykissool 的回复:

引用 2 楼 bianjunlan 的回复:1，Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256;指的是一个workgroup内不同维度的#item最大值 global_work_item的大小是 workgroup大小与workgroup个数的乘积，所以可以这样设置 ……

opencl里没有直接的相关函数，因为这个上限跟核心数目，global memoey，local memory，private memory都有关，所以一定跟具体的程序占用资源量有关。

[hehykissool]

引用 6 楼 bianjunlan 的回复:

引用 5 楼 hehykissool 的回复:引用 2 楼 bianjunlan 的回复:1，Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256;指的是一个workgroup内不同维度的#item最大值 global_work_item的大小是 workgroup大小与wo……

感谢您的解答，我还想请问为什么要采用work_group，有什么益处吗，因为如果不设置的话，系统会自动分配，自动分配与手动设置的区别在哪里

[outstander]

受限于work item数的话有一种方式可以进行处理: 在kernel中写一个循环，每个work item处理完本来分配的任务后不退出， 继续处理掉剩下的任务即可。

[outstander]

引用 6 楼 bianjunlan 的回复:

引用 5 楼 hehykissool 的回复:引用 2 楼 bianjunlan 的回复:1，Max_work_item[0] = 256;
Max_work_item[1] = 256;
Max_work_item[2] = 256;指的是一个workgroup内不同维度的#item最大值

global_work_item的大小是 workgroup大小与wo……

应该有函数进行查找的， clinfo程序其实就是用那些查找设备和平台特性的函数写的，统一输出了而已

[hehykissool]

引用 8 楼 outstander 的回复:

受限于work item数的话有一种方式可以进行处理: 在kernel中写一个循环，每个work item处理完本来分配的任务后不退出， 继续处理掉剩下的任务即可。

能稍微示例一下吗，怎样继续处理不退出，没看懂

[hehykissool]

引用 8 楼 outstander 的回复:

受限于work item数的话有一种方式可以进行处理: 在kernel中写一个循环，每个work item处理完本来分配的任务后不退出， 继续处理掉剩下的任务即可。

保存kernel的话应该也会绑定数据吧

[卞亚涛]

引用 7 楼 hehykissool 的回复:

引用 6 楼 bianjunlan 的回复:引用 5 楼 hehykissool 的回复:引用 2 楼 bianjunlan 的回复:1，Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_item[2] = 256;指的是一个workgroup内不同维度的#item最大值 global_work_i……

关于为什么使用work group，我觉得有以下原因： 1， work group可以让你对 work item进行同步，例如设置barrier等 2， workgroup可以利用local memory进行数据交互， local memory的速度极快 3， 对于某些特定的算法，使用workgroup方便算法的设计，当然也跟前两点有关系。比如并行的粒子滤波算法，可以用一个workgroup进行一个粒子的运算，这里粒子相关的计算，同步，资源等等均可以在此workgroup内进行。 ps: 全局的数据交互可以使用 global memory，但是这个很慢，最好先workgroup内交互，再全局交互

[卞亚涛]

引用 9 楼 outstander 的回复:

本帖最后由 outstander 于 2013-04-10 23:05:23 编辑 引用 6 楼 bianjunlan 的回复:引用 5 楼 hehykissool 的回复:引用 2 楼 bianjunlan 的回复:1，Max_work_item[0] = 256; Max_work_item[1] = 256; Max_work_i……

是有函数查找，不过查找的是资源的上限。他问的是 对于特定问题可以使用的#workgroup的上限，这个跟具体问题的资源使用量有关。

[outstander]

引用 10 楼 hehykissool 的回复:

引用 8 楼 outstander 的回复:受限于work item数的话有一种方式可以进行处理: 在kernel中写一个循环，每个work item处理完本来分配的任务后不退出， 继续处理掉剩下的任务即可。 能稍微示例一下吗，怎样继续处理不退出，没看懂

假设对一个数组 Array[N] 乘以一个常数alpha。 如果只能发起了512个线程，可以用以下代码 for（int i = get_global_id(0); i < N; i+=512） { Array[i] *= alpha; } 这样用判断边界的方式，用有限的线程通过循环处理任意大小的任务