1 // 2 // Copyright (c) 2021 The Khronos Group Inc. 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 // http://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 // 16 #include "subgroup_common_kernels.h" 17 18 const char* bcast_source = 19 "__kernel void test_bcast(const __global Type *in, " 20 "__global int4 *xy, __global Type *out)\n" 21 "{\n" 22 " int gid = get_global_id(0);\n" 23 " XY(xy,gid);\n" 24 " Type x = in[gid];\n" 25 " uint which_sub_group_local_id = xy[gid].z;\n" 26 " out[gid] = sub_group_broadcast(x, which_sub_group_local_id);\n" 27 28 "}\n"; 29 30 const char* redadd_source = "__kernel void test_redadd(const __global Type " 31 "*in, __global int4 *xy, __global Type *out)\n" 32 "{\n" 33 " int gid = get_global_id(0);\n" 34 " XY(xy,gid);\n" 35 " out[gid] = sub_group_reduce_add(in[gid]);\n" 36 "}\n"; 37 38 const char* redmax_source = "__kernel void test_redmax(const __global Type " 39 "*in, __global int4 *xy, __global Type *out)\n" 40 "{\n" 41 " int gid = get_global_id(0);\n" 42 " XY(xy,gid);\n" 43 " out[gid] = sub_group_reduce_max(in[gid]);\n" 44 "}\n"; 45 46 const char* redmin_source = "__kernel void test_redmin(const __global Type " 47 "*in, __global int4 *xy, __global Type *out)\n" 48 "{\n" 49 " int gid = get_global_id(0);\n" 50 " XY(xy,gid);\n" 51 " out[gid] = sub_group_reduce_min(in[gid]);\n" 52 "}\n"; 53 54 const char* scinadd_source = 55 "__kernel void test_scinadd(const __global Type *in, __global int4 *xy, " 56 "__global Type *out)\n" 57 "{\n" 58 " int gid = get_global_id(0);\n" 59 " XY(xy,gid);\n" 60 " out[gid] = sub_group_scan_inclusive_add(in[gid]);\n" 61 "}\n"; 62 63 const char* scinmax_source = 64 "__kernel void test_scinmax(const __global Type *in, __global int4 *xy, " 65 "__global Type *out)\n" 66 "{\n" 67 " int gid = get_global_id(0);\n" 68 " XY(xy,gid);\n" 69 " out[gid] = sub_group_scan_inclusive_max(in[gid]);\n" 70 "}\n"; 71 72 const char* scinmin_source = 73 "__kernel void test_scinmin(const __global Type *in, __global int4 *xy, " 74 "__global Type *out)\n" 75 "{\n" 76 " int gid = get_global_id(0);\n" 77 " XY(xy,gid);\n" 78 " out[gid] = sub_group_scan_inclusive_min(in[gid]);\n" 79 "}\n"; 80 81 const char* scexadd_source = 82 "__kernel void test_scexadd(const __global Type *in, __global int4 *xy, " 83 "__global Type *out)\n" 84 "{\n" 85 " int gid = get_global_id(0);\n" 86 " XY(xy,gid);\n" 87 " out[gid] = sub_group_scan_exclusive_add(in[gid]);\n" 88 "}\n"; 89 90 const char* scexmax_source = 91 "__kernel void test_scexmax(const __global Type *in, __global int4 *xy, " 92 "__global Type *out)\n" 93 "{\n" 94 " int gid = get_global_id(0);\n" 95 " XY(xy,gid);\n" 96 " out[gid] = sub_group_scan_exclusive_max(in[gid]);\n" 97 "}\n"; 98 99 const char* scexmin_source = 100 "__kernel void test_scexmin(const __global Type *in, __global int4 *xy, " 101 "__global Type *out)\n" 102 "{\n" 103 " int gid = get_global_id(0);\n" 104 " XY(xy,gid);\n" 105 " out[gid] = sub_group_scan_exclusive_min(in[gid]);\n" 106 "}\n"; 107