diff options
| author | Pan Xiuli <xiuli.pan@intel.com> | 2016-05-12 09:13:26 +0800 |
|---|---|---|
| committer | Yang Rong <rong.r.yang@intel.com> | 2016-06-13 17:02:16 +0800 |
| commit | b37871995fab6601936460d17180cccb7a5b0dee (patch) | |
| tree | f7ba10fc6752755b2f615038119124935ce5d8d1 /kernels | |
| parent | a2dd4f9bdc55026c9667db8b2d44a9dd529785e7 (diff) | |
Utest: Add subgroup work item test cases
Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Diffstat (limited to 'kernels')
| -rw-r--r-- | kernels/builtin_max_sub_group_size.cl | 7 | ||||
| -rw-r--r-- | kernels/builtin_num_sub_groups.cl | 7 | ||||
| -rw-r--r-- | kernels/builtin_sub_group_id.cl | 7 | ||||
| -rw-r--r-- | kernels/builtin_sub_group_size.cl | 7 |
4 files changed, 28 insertions, 0 deletions
diff --git a/kernels/builtin_max_sub_group_size.cl b/kernels/builtin_max_sub_group_size.cl new file mode 100644 index 00000000..c2f3b5ef --- /dev/null +++ b/kernels/builtin_max_sub_group_size.cl @@ -0,0 +1,7 @@ +__kernel void builtin_max_sub_group_size(global int *dst) +{ + int lid = get_local_linear_id(); + int lsz = get_local_size(0) * get_local_size(1) * get_local_size(2); + int gid = lid + lsz*(get_num_groups(1) * get_num_groups(0) * get_group_id(2) + get_num_groups(0) * get_group_id(1) + get_group_id(0)); + dst[gid] = get_max_sub_group_size(); +} diff --git a/kernels/builtin_num_sub_groups.cl b/kernels/builtin_num_sub_groups.cl new file mode 100644 index 00000000..08b56733 --- /dev/null +++ b/kernels/builtin_num_sub_groups.cl @@ -0,0 +1,7 @@ +__kernel void builtin_num_sub_groups(global int *dst) +{ + int lid = get_local_linear_id(); + int lsz = get_local_size(0) * get_local_size(1) * get_local_size(2); + int gid = lid + lsz*(get_num_groups(1) * get_num_groups(0) * get_group_id(2) + get_num_groups(0) * get_group_id(1) + get_group_id(0)); + dst[gid] = get_num_sub_groups(); +} diff --git a/kernels/builtin_sub_group_id.cl b/kernels/builtin_sub_group_id.cl new file mode 100644 index 00000000..accf3ad4 --- /dev/null +++ b/kernels/builtin_sub_group_id.cl @@ -0,0 +1,7 @@ +__kernel void builtin_sub_group_id(global int *dst) +{ + int lid = get_local_linear_id(); + int lsz = get_local_size(0) * get_local_size(1) * get_local_size(2); + int gid = lid + lsz*(get_num_groups(1) * get_num_groups(0) * get_group_id(2) + get_num_groups(0) * get_group_id(1) + get_group_id(0)); + dst[gid] = get_sub_group_id(); +} diff --git a/kernels/builtin_sub_group_size.cl b/kernels/builtin_sub_group_size.cl new file mode 100644 index 00000000..1e034bb5 --- /dev/null +++ b/kernels/builtin_sub_group_size.cl @@ -0,0 +1,7 @@ +__kernel void builtin_sub_group_size(global int *dst) +{ + int lid = get_local_linear_id(); + int lsz = get_local_size(0) * get_local_size(1) * get_local_size(2); + int gid = lid + lsz*(get_num_groups(1) * get_num_groups(0) * get_group_id(2) + get_num_groups(0) * get_group_id(1) + get_group_id(0)); + dst[gid] = get_sub_group_size(); +} |