{%- include "structs.wgsl" -%}

@group(0) @binding(0)
var<storage, read> input_0: Array;

@group(0) @binding(1)
var<storage, read> input_1: ArrayMatrix;

{%- if i_lens | length == 3 -%} // Bias
	@group(0) @binding(2)
	var<storage, read> input_2: ArrayVector;

	@group(0) @binding(3)
	var<storage, read_write> output_0: Array;

{%- else -%}
	@group(0) @binding(2)
	var<storage, read_write> output_0: Array;

{%- endif %}


@compute @workgroup_size(256, 1, 1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
	let gidx = global_id.x;
	if (gidx < {{ o_lens[0]/4 | int }}u) {
		let batch = gidx / {{ o_chunks[0][0] / 4 | int }}u; 
		let rest = gidx % {{ o_chunks[0][0] / 4 | int }}u; 

		let m = rest / {{ o_chunks[0][1] }}u;
		let xy = rest % {{ o_chunks[0][1] }}u;
			
		var result = Vec4(Scalar(), Scalar(), Scalar(), Scalar());
		
		let root_index = batch * {{ i_chunks[0][0] }}u + xy;
		let root_kernel_index = m * {{ channel / 16 * 4 | int }}u;

		for(var c: u32 = 0u; c < {{ channel / 16 | int }}u; c = c + 1u) {
			let base_index = root_index + c * {{ 16 * i_chunks[0][1] }}u;
			var base_kernel_index = root_kernel_index + c;
			
			var matrix_0 = input_1.data[base_kernel_index];
			var matrix_1 = input_1.data[base_kernel_index + {{ channel / 16 | int }}u];
			var matrix_2 = input_1.data[base_kernel_index + {{ 2 * channel / 16 | int }}u];
			var matrix_3 = input_1.data[base_kernel_index + {{ 3 * channel / 16 | int }}u];

			for(var index_c_vec: u32 = 0u; index_c_vec < 4u; index_c_vec = index_c_vec + 1u) {
				let base_index_2 = base_index + index_c_vec * {{ 4 * i_chunks[0][1] }}u;

				let tmp_vec = Vec4(
					input_0.data[base_index_2],
					input_0.data[base_index_2 + {{ i_chunks[0][1] }}u],
					input_0.data[base_index_2 + {{ 2 * i_chunks[0][1] }}u],
					input_0.data[base_index_2 + {{ 3 * i_chunks[0][1] }}u],
				);
			
				result = tmp_vec * Mat4x4(
					matrix_0[index_c_vec],
					matrix_1[index_c_vec],
					matrix_2[index_c_vec],
					matrix_3[index_c_vec],
				) + result; 
			}
		}

		{% if i_lens | length == 3 -%}
			result = result + input_2.data[m];
		{%- endif %}

		{% set activation_input = "result" %}
		{% set activation_output = "result" %}
		{% set activation_type = op_type | replace(from="Conv", to="") %}
		{%- include "snippets/activation_vec.wgsl" %}

		let base_index_3 = batch * {{ o_chunks[0][0] }}u + m * {{ o_chunks[0][1] * 4 }}u + xy;
		for(var index_vec: u32 = 0u; index_vec < 4u; index_vec = index_vec + 1u) {
			let index = base_index_3 + index_vec * {{ o_chunks[0][1] }}u;
			output_0.data[index] = result[index_vec];
		}
	}
}