Commit 974daec4 authored by Jason Ekstrand's avatar Jason Ekstrand

i965/fs: Implement basic SPIR-V subgroup intrinsics

Reviewed-by: Samuel Iglesias Gonsálvez's avatarSamuel Iglesias Gonsálvez <siglesias@igalia.com>
Reviewed-by: Iago Toral's avatarIago Toral Quiroga <itoral@igalia.com>
parent adc07779
......@@ -4501,6 +4501,14 @@ fs_visitor::nir_emit_intrinsic(const fs_builder &bld, nir_intrinsic_instr *instr
break;
}
case nir_intrinsic_first_invocation: {
fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_UD);
bld.exec_all().emit(SHADER_OPCODE_FIND_LIVE_CHANNEL, tmp);
bld.MOV(retype(dest, BRW_REGISTER_TYPE_UD),
fs_reg(component(tmp, 0)));
break;
}
default:
unreachable("unknown intrinsic");
}
......
......@@ -103,6 +103,24 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
break;
}
case nir_intrinsic_load_subgroup_id:
if (state->local_workgroup_size > 8)
continue;
/* For small workgroup sizes, we know subgroup_id will be zero */
sysval = nir_imm_int(b, 0);
break;
case nir_intrinsic_load_num_subgroups: {
unsigned local_workgroup_size =
nir->info.cs.local_size[0] * nir->info.cs.local_size[1] *
nir->info.cs.local_size[2];
unsigned num_subgroups =
DIV_ROUND_UP(local_workgroup_size, state->dispatch_width);
sysval = nir_imm_int(b, num_subgroups);
break;
}
default:
continue;
}
......
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment