1#!amber
2# Copyright 2019 The Amber Authors.
3#
4# Licensed under the Apache License, Version 2.0 (the "License");
5# you may not use this file except in compliance with the License.
6# You may obtain a copy of the License at
7#
8#     https://www.apache.org/licenses/LICENSE-2.0
9#
10# Unless required by applicable law or agreed to in writing, software
11# distributed under the License is distributed on an "AS IS" BASIS,
12# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13# See the License for the specific language governing permissions and
14# limitations under the License.
15
16SHADER compute my_shader OPENCL-C
17kernel void foo(global int* in, global int* out) {
18  unsigned int local_size_x = get_local_size(0);
19  unsigned int local_size_y = get_local_size(1);
20  unsigned int local_size_z = get_local_size(2);
21  unsigned int local_id_x = get_local_id(0);
22  unsigned int local_id_y = get_local_id(1);
23  unsigned int local_id_z = get_local_id(2);
24  unsigned int group_id_x = get_group_id(0);
25  unsigned int group_id_y = get_group_id(1);
26  unsigned int group_id_z = get_group_id(2);
27  unsigned int global_id_x = get_global_id(0);
28  unsigned int global_id_y = get_global_id(1);
29  unsigned int global_id_z = get_global_id(2);
30  unsigned int wgs_x = get_num_groups(0);
31  unsigned int wgs_y = get_num_groups(1);
32  unsigned int wgs_z = get_num_groups(2);
33
34  unsigned int in_wg_id = (local_id_z * local_size_x * local_size_y) +
35    (local_id_y * local_size_x) +
36    local_id_x;
37  unsigned int prev_ids = (local_size_x * local_size_y * local_size_z) *
38    (group_id_z * wgs_y * wgs_x + group_id_y * wgs_x + group_id_x);
39  unsigned int linear_id = in_wg_id + prev_ids;
40  out[linear_id] = in[linear_id];
41}
42END
43
44BUFFER in_buf DATA_TYPE uint32 SIZE 64 SERIES_FROM 1 INC_BY 1
45BUFFER out_buf DATA_TYPE uint32 SIZE 64 FILL 0
46
47PIPELINE compute my_pipeline
48  ATTACH my_shader ENTRY_POINT foo \
49    SPECIALIZE 0 AS uint32 2 \
50    SPECIALIZE 1 AS uint32 2 \
51    SPECIALIZE 2 AS uint32 2
52
53  BIND BUFFER in_buf KERNEL ARG_NAME in
54  BIND BUFFER out_buf KERNEL ARG_NUMBER 1
55END
56
57RUN my_pipeline 2 2 2
58
59EXPECT out_buf EQ_BUFFER in_buf
60
61