1 // Copyright (c) 2016-2020 Intel Corporation
3 // Licensed under the Apache License, Version 2.0 (the "License");
4 // you may not use this file except in compliance with the License.
5 // You may obtain a copy of the License at
7 // http://www.apache.org/licenses/LICENSE-2.0
9 // Unless required by applicable law or agreed to in writing, software
10 // distributed under the License is distributed on an "AS IS" BASIS,
11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 // See the License for the specific language governing permissions and
13 // limitations under the License.
16 #include "include/common.cl"
18 #include "include/data_types.cl"
19 #include "include/fetch.cl"
20 #include "include/mmad.cl"
22 #define INPUT_PACKED_TYPE_8 CAT(INPUT_PACKED_TYPE, 8)
23 #define FILTER_PACKED_TYPE_8 CAT(FILTER_PACKED_TYPE, 8)
25 #define AS_TYPE(type, val) CAT(as_, type)(val)
27 __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE)))
28 KERNEL(fully_connected_gpu_MMAD)(
29 const __global INPUT0_TYPE* input,
30 __global OUTPUT_TYPE* output,
31 const __global FILTER_TYPE* weights
33 , const __global BIAS_TYPE* biases
35 #if HAS_FUSED_OPS_DECLS
40 #if OUTPUT_BATCH_NUM == 1
41 const uint f = (uint)get_global_id(0);
44 const uint f = (uint)get_global_id(0);
45 const uint b = (uint)get_global_id(1);
50 const uint filter_offset = FILTER_GET_OFFSET(f);
52 const uint input_offset = INPUT0_GET_INDEX(b, 0, 0, 0, 0);
54 const uint input_offset = INPUT0_GET_INDEX(b, 0, 0, 0);
58 for (uint k = 0; k < FEATURE_BLOCKS_COUNT; ++k) {
60 for (uint spatial = 0; spatial < FILTER_SPATIAL_SIZE; ++spatial) {
62 for (uint zi = 0; zi < FILTER_SIZE_Z; ++zi)
63 for (uint yi = 0; yi < FILTER_SIZE_Y; ++yi)
64 for (uint xi = 0; xi < FILTER_SIZE_X; ++xi) {
65 const uint spatial = xi + yi * FILTER_SIZE_X + zi * FILTER_SIZE_X * FILTER_SIZE_Y;
67 #else // SPATIAL_MAJOR
69 for (uint spatial = 0; spatial < FILTER_SPATIAL_SIZE; ++spatial) {
71 for (uint zi = 0; zi < FILTER_SIZE_Z; ++zi)
72 for (uint yi = 0; yi < FILTER_SIZE_Y; ++yi)
73 for (uint xi = 0; xi < FILTER_SIZE_X; ++xi) {
74 const uint spatial = xi + yi * FILTER_SIZE_X + zi * FILTER_SIZE_X * FILTER_SIZE_Y;
76 for (uint k = 0; k < FEATURE_BLOCKS_COUNT; ++k) {
79 uint input_idx = input_offset + spatial * MMAD_INPUT_SPATIAL_PITCH + k * MMAD_INPUT_FBLOCK_PITCH;
81 uint input_idx = input_offset + k * MMAD_INPUT_FBLOCK_PITCH + zi * MMAD_INPUT_Z_PITCH + yi * MMAD_INPUT_Y_PITCH + xi * MMAD_INPUT_X_PITCH;
83 uint filter_idx = filter_offset + spatial * MMAD_FILTER_SPATIAL_PITCH + k * MMAD_FILTER_FBLOCK_PITCH;
85 uint input_data_u = intel_sub_group_block_read((const __global uint*)(input + input_idx));
86 INPUT_PACKED_TYPE input_data = AS_TYPE(INPUT_PACKED_TYPE, input_data_u);
88 INPUT_PACKED_TYPE_8 activations; //activations of all lanes
89 activations.s0 = sub_group_broadcast(input_data, 0);
90 activations.s1 = sub_group_broadcast(input_data, 1);
91 activations.s2 = sub_group_broadcast(input_data, 2);
92 activations.s3 = sub_group_broadcast(input_data, 3);
93 activations.s4 = sub_group_broadcast(input_data, 4);
94 activations.s5 = sub_group_broadcast(input_data, 5);
95 activations.s6 = sub_group_broadcast(input_data, 6);
96 activations.s7 = sub_group_broadcast(input_data, 7);
98 uint8 weights_data_u = intel_sub_group_block_read8((const __global uint*)(weights + filter_idx));
99 FILTER_PACKED_TYPE_8 weights_data = AS_TYPE(FILTER_PACKED_TYPE_8, weights_data_u);
101 dotProd = MMAD_8(activations, weights_data, dotProd);
105 #if HAS_FEATURE_LEFTOVERS
106 const uint lid = get_sub_group_local_id();
109 for (uint spatial = 0; spatial < FILTER_SPATIAL_SIZE; ++spatial) {
111 for (uint zi = 0; zi < FILTER_SIZE_Z; ++zi)
112 for (uint yi = 0; yi < FILTER_SIZE_Y; ++yi)
113 for (uint xi = 0; xi < FILTER_SIZE_X; ++xi) {
114 const uint spatial = xi + yi * FILTER_SIZE_X + zi * FILTER_SIZE_X * FILTER_SIZE_Y;
115 #endif // !SPLIT_SPATIAL
117 #else // SPATIAL_MAJOR
119 for (uint spatial = 0; spatial < FILTER_SPATIAL_SIZE; ++spatial) {
120 #else // !SPLIT_SPATIAL
121 for (uint zi = 0; zi < FILTER_SIZE_Z; ++zi)
122 for (uint yi = 0; yi < FILTER_SIZE_Y; ++yi)
123 for (uint xi = 0; xi < FILTER_SIZE_X; ++xi) {
124 const uint spatial = xi + yi * FILTER_SIZE_X + zi * FILTER_SIZE_X * FILTER_SIZE_Y;
125 #endif // !SPLIT_SPATIAL
126 #endif // SPATIAL_MAJOR
129 uint input_idx = input_offset + spatial * MMAD_INPUT_SPATIAL_PITCH + FEATURE_BLOCKS_COUNT * INPUT0_FEATURE_PITCH;
130 #else // !SPLIT_SPATIAL
131 uint input_idx = input_offset + FEATURE_BLOCK_COUNT * INPUT0_FEATURE_PITCH + zi * MMAD_INPUT_Z_PITCH + yi * MMAD_INPUT_Y_PITCH + xi * MMAD_INPUT_X_PITCH;
132 #endif // !SPLIT_SPATIAL
133 uint filter_idx = filter_offset + spatial * MMAD_FILTER_SPATIAL_PITCH + FEATURE_BLOCKS_COUNT * MMAD_FILTER_FBLOCK_PITCH;
135 MAKE_VECTOR_TYPE(INPUT0_TYPE, 4) input_data_u = (0, 0, 0, 0);
136 for (uint i = 0; i < 4; i++) {
137 if (FEATURE_BLOCKS_COUNT*32 + lid*4 + i < INPUT0_FEATURE_NUM) {
138 input_data_u[i] = input[input_idx + (lid*4 + i)*INPUT0_FEATURE_PITCH];
141 INPUT_PACKED_TYPE input_data = AS_TYPE(INPUT_PACKED_TYPE, input_data_u);
143 INPUT_PACKED_TYPE_8 activations; //activations of all lanes
144 activations.s0 = sub_group_broadcast(input_data, 0);
145 activations.s1 = sub_group_broadcast(input_data, 1);
146 activations.s2 = sub_group_broadcast(input_data, 2);
147 activations.s3 = sub_group_broadcast(input_data, 3);
148 activations.s4 = sub_group_broadcast(input_data, 4);
149 activations.s5 = sub_group_broadcast(input_data, 5);
150 activations.s6 = sub_group_broadcast(input_data, 6);
151 activations.s7 = sub_group_broadcast(input_data, 7);
153 uint8 weights_data_u = intel_sub_group_block_read8((const __global uint*)(weights + filter_idx));
154 FILTER_PACKED_TYPE_8 weights_data = AS_TYPE(FILTER_PACKED_TYPE_8, weights_data_u);
156 dotProd = MMAD_8(activations, weights_data, dotProd);
158 #endif // HAS_FEATURE_LEFTOVERS
160 if (OUTPUT_FEATURE_NUM % SUB_GROUP_SIZE != 0 && f >= OUTPUT_FEATURE_NUM)
165 const uint bias_index = GET_DATA_INDEX(BIAS, b, f, 0, 0);
167 const uint bias_index = f;
170 float dequantized = (float)dotProd + biases[bias_index];
172 float dequantized = (float)dotProd;
175 const uint out_idx = OUTPUT_GET_INDEX(b, f, 0, 0);
179 OUTPUT_TYPE res = FUSED_OPS_RESULT;
181 output[out_idx] = res;
183 output[out_idx] = TO_OUTPUT_TYPE(dequantized);
187 #undef INPUT_PACKED_TYPE_8
188 #undef FILTER_PACKED_TYPE_8