OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

broadcom/compiler: handle load_workgroup_size

Browse Source 
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29554>
This commit is contained in:
Karol Herbst
2024-06-05 18:40:47 +02:00
committed by Marge Bot
parent 3aafe75471
commit 83883a6cc2
2 changed files with 15 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
src/broadcom/compiler/nir_to_vir.c
View File

@@ -3668,6 +3668,18 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
break;
}
case nir_intrinsic_load_workgroup_size: {
struct qreg x = vir_uniform(c, QUNIFORM_WORK_GROUP_SIZE, 0);
ntq_store_def(c, &instr->def, 0, x);
struct qreg y = vir_uniform(c, QUNIFORM_WORK_GROUP_SIZE, 1);
ntq_store_def(c, &instr->def, 1, y);
struct qreg z = vir_uniform(c, QUNIFORM_WORK_GROUP_SIZE, 2);
ntq_store_def(c, &instr->def, 2, z);
break;
}
case nir_intrinsic_load_local_invocation_index:
ntq_store_def(c, &instr->def, 0,
emit_load_local_invocation_index(c));

3
src/broadcom/compiler/v3d_compiler.h
View File

@@ -311,6 +311,9 @@ enum quniform_contents {
*/
QUNIFORM_WORK_GROUP_BASE,
/* Workgroup size for variable workgroup support */
QUNIFORM_WORK_GROUP_SIZE,
/**
* Returns the the offset of the scratch buffer for register spilling.
*/