Home
last modified time | relevance | path

Searched refs:local_size (Results 1 – 25 of 136) sorted by relevance

123456

/third_party/mesa3d/src/intel/compiler/
Dtest_simd_selection.cpp71 prog_data->local_size[0] = 32; in SIMDSelectionCS()
72 prog_data->local_size[1] = 1; in SIMDSelectionCS()
73 prog_data->local_size[2] = 1; in SIMDSelectionCS()
92 prog_data->local_size[0] = devinfo->max_cs_workgroup_threads; in TEST_F()
93 prog_data->local_size[1] = 32; in TEST_F()
94 prog_data->local_size[2] = 1; in TEST_F()
106 prog_data->local_size[0] = 1; in TEST_F()
107 prog_data->local_size[1] = 1; in TEST_F()
108 prog_data->local_size[2] = 1; in TEST_F()
120 prog_data->local_size[0] = 8; in TEST_F()
[all …]
Dbrw_simd_selection.c66 const bool workgroup_size_variable = prog_data->local_size[0] == 0; in brw_simd_should_compile()
75 const unsigned workgroup_size = prog_data->local_size[0] * in brw_simd_should_compile()
76 prog_data->local_size[1] * in brw_simd_should_compile()
77 prog_data->local_size[2]; in brw_simd_should_compile()
170 if (prog_data->local_size[0] == sizes[0] && in brw_simd_select_for_workgroup_size()
171 prog_data->local_size[1] == sizes[1] && in brw_simd_select_for_workgroup_size()
172 prog_data->local_size[2] == sizes[2]) in brw_simd_select_for_workgroup_size()
179 cloned.local_size[i] = sizes[i]; in brw_simd_select_for_workgroup_size()
Dintel_clc.c159 cs_prog_data->local_size[0], in print_cs_prog_data_fields()
160 cs_prog_data->local_size[1], in print_cs_prog_data_fields()
161 cs_prog_data->local_size[2]); in print_cs_prog_data_fields()
543 kernel.prog_data.local_size[0], in main()
544 kernel.prog_data.local_size[1], in main()
545 kernel.prog_data.local_size[2]); in main()
/third_party/mesa3d/src/gallium/drivers/freedreno/a5xx/
Dfd5_compute.c145 const unsigned *local_size = in fd5_launch_grid() local
152 A5XX_HLSQ_CS_NDRANGE_0_LOCALSIZEX(local_size[0] - 1) | in fd5_launch_grid()
153 A5XX_HLSQ_CS_NDRANGE_0_LOCALSIZEY(local_size[1] - 1) | in fd5_launch_grid()
154 A5XX_HLSQ_CS_NDRANGE_0_LOCALSIZEZ(local_size[2] - 1)); in fd5_launch_grid()
156 A5XX_HLSQ_CS_NDRANGE_1_GLOBALSIZE_X(local_size[0] * num_groups[0])); in fd5_launch_grid()
159 A5XX_HLSQ_CS_NDRANGE_3_GLOBALSIZE_Y(local_size[1] * num_groups[1])); in fd5_launch_grid()
162 A5XX_HLSQ_CS_NDRANGE_5_GLOBALSIZE_Z(local_size[2] * num_groups[2])); in fd5_launch_grid()
179 A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEX(local_size[0] - 1) | in fd5_launch_grid()
180 A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEY(local_size[1] - 1) | in fd5_launch_grid()
181 A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEZ(local_size[2] - 1)); in fd5_launch_grid()
/third_party/pcre2/pcre2/src/sljit/
DsljitNativeX86_64.c375 sljit_s32 fscratches, sljit_s32 fsaveds, sljit_s32 local_size) in sljit_emit_enter() argument
389 …ljit_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size)); in sljit_emit_enter()
390 set_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size); in sljit_emit_enter()
425 local_size += SLJIT_LOCALS_OFFSET; in sljit_emit_enter()
429 saved_float_regs_offset = ((local_size + 0xf) & ~0xf); in sljit_emit_enter()
430 local_size = saved_float_regs_offset + saved_float_regs_size; in sljit_emit_enter()
492 local_size = ((local_size + saved_regs_size + 0xf) & ~0xf) - saved_regs_size; in sljit_emit_enter()
493 compiler->local_size = local_size; in sljit_emit_enter()
496 if (local_size > 0) { in sljit_emit_enter()
497 if (local_size <= 4 * 4096) { in sljit_emit_enter()
[all …]
DsljitNativeX86_32.c264 sljit_s32 fscratches, sljit_s32 fsaveds, sljit_s32 local_size) in sljit_emit_enter() argument
275 …ljit_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size)); in sljit_emit_enter()
276 set_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size); in sljit_emit_enter()
372 local_size = ((SLJIT_LOCALS_OFFSET_BASE + local_size + size + 0xf) & ~0xf) - size; in sljit_emit_enter()
373 compiler->local_size = local_size; in sljit_emit_enter()
392 if (word_arg_count == 3 && local_size > 4 * 4096) in sljit_emit_enter()
393 r2_offset = local_size + args_size; in sljit_emit_enter()
412 SLJIT_ASSERT(r2_offset == -1 || local_size > 4 * 4096); in sljit_emit_enter()
414 if (local_size > 4096) { in sljit_emit_enter()
415 if (local_size <= 4 * 4096) { in sljit_emit_enter()
[all …]
DsljitNativeARM_64.c972 sljit_s32 fscratches, sljit_s32 fsaveds, sljit_s32 local_size) in sljit_emit_enter() argument
979 …ljit_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size)); in sljit_emit_enter()
980 set_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size); in sljit_emit_enter()
985 local_size = (local_size + saved_regs_size + 0xf) & ~0xf; in sljit_emit_enter()
986 compiler->local_size = local_size; in sljit_emit_enter()
988 if (local_size <= 512) { in sljit_emit_enter()
990 | RN(SLJIT_SP) | (sljit_ins)((-(local_size >> 3) & 0x7f) << 15))); in sljit_emit_enter()
991 offs = (sljit_ins)(local_size - 2 * SSIZE_OF(sw)) << (15 - 3); in sljit_emit_enter()
992 local_size = 0; in sljit_emit_enter()
998 local_size -= saved_regs_size; in sljit_emit_enter()
[all …]
DsljitNativeARM_T2_32.c1139 sljit_s32 fscratches, sljit_s32 fsaveds, sljit_s32 local_size) in sljit_emit_enter() argument
1154 …ljit_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size)); in sljit_emit_enter()
1155 set_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size); in sljit_emit_enter()
1188 local_size = ((size + local_size + 0x7) & ~0x7) - size; in sljit_emit_enter()
1189 compiler->local_size = local_size; in sljit_emit_enter()
1296 if (local_size >= 4096) { in sljit_emit_enter()
1302 if (local_size < 4 * 4096) { in sljit_emit_enter()
1303 if (local_size > 2 * 4096) { in sljit_emit_enter()
1304 if (local_size > 3 * 4096) { in sljit_emit_enter()
1313 FAIL_IF(load_immediate(compiler, TMP_REG2, ((sljit_uw)local_size >> 12) - 1)); in sljit_emit_enter()
[all …]
DsljitNativeRISCV_common.c599 sljit_s32 fscratches, sljit_s32 fsaveds, sljit_s32 local_size) in sljit_emit_enter() argument
605 …ljit_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size)); in sljit_emit_enter()
606 set_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size); in sljit_emit_enter()
608 local_size += GET_SAVED_REGISTERS_SIZE(scratches, saveds - saved_arg_count, 1); in sljit_emit_enter()
611 if ((local_size & SSIZE_OF(sw)) != 0) in sljit_emit_enter()
612 local_size += SSIZE_OF(sw); in sljit_emit_enter()
613 local_size += GET_SAVED_FLOAT_REGISTERS_SIZE(fscratches, fsaveds, sizeof(sljit_f64)); in sljit_emit_enter()
616 local_size += GET_SAVED_FLOAT_REGISTERS_SIZE(fscratches, fsaveds, sizeof(sljit_f64)); in sljit_emit_enter()
618 local_size = (local_size + SLJIT_LOCALS_OFFSET + 15) & ~0xf; in sljit_emit_enter()
619 compiler->local_size = local_size; in sljit_emit_enter()
[all …]
DsljitNativePPC_common.c752 sljit_s32 fscratches, sljit_s32 fsaveds, sljit_s32 local_size) in sljit_emit_enter() argument
762 …ljit_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size)); in sljit_emit_enter()
763 set_emit_enter(compiler, options, arg_types, scratches, saveds, fscratches, fsaveds, local_size); in sljit_emit_enter()
765 local_size += GET_SAVED_REGISTERS_SIZE(scratches, saveds - saved_arg_count, 0) in sljit_emit_enter()
769 local_size += SSIZE_OF(sw); in sljit_emit_enter()
771 local_size = (local_size + SLJIT_LOCALS_OFFSET + 15) & ~0xf; in sljit_emit_enter()
772 compiler->local_size = local_size; in sljit_emit_enter()
777 offset = local_size; in sljit_emit_enter()
779 if (local_size <= STACK_MAX_DISTANCE) { in sljit_emit_enter()
781 FAIL_IF(push_inst(compiler, STWU | S(SLJIT_SP) | A(SLJIT_SP) | IMM(-local_size))); in sljit_emit_enter()
[all …]
/third_party/mesa3d/src/gallium/drivers/freedreno/a4xx/
Dfd4_compute.c164 const unsigned *local_size = in fd4_launch_grid() local
171 A4XX_HLSQ_CL_NDRANGE_0_LOCALSIZEX(local_size[0] - 1) | in fd4_launch_grid()
172 A4XX_HLSQ_CL_NDRANGE_0_LOCALSIZEY(local_size[1] - 1) | in fd4_launch_grid()
173 A4XX_HLSQ_CL_NDRANGE_0_LOCALSIZEZ(local_size[2] - 1)); in fd4_launch_grid()
175 A4XX_HLSQ_CL_NDRANGE_1_SIZE_X(local_size[0] * num_groups[0])); in fd4_launch_grid()
178 A4XX_HLSQ_CL_NDRANGE_3_SIZE_Y(local_size[1] * num_groups[1])); in fd4_launch_grid()
181 A4XX_HLSQ_CL_NDRANGE_5_SIZE_Z(local_size[2] * num_groups[2])); in fd4_launch_grid()
194 A4XX_CP_EXEC_CS_INDIRECT_2_LOCALSIZEX(local_size[0] - 1) | in fd4_launch_grid()
195 A4XX_CP_EXEC_CS_INDIRECT_2_LOCALSIZEY(local_size[1] - 1) | in fd4_launch_grid()
196 A4XX_CP_EXEC_CS_INDIRECT_2_LOCALSIZEZ(local_size[2] - 1)); in fd4_launch_grid()
/third_party/mesa3d/src/freedreno/computerator/
Dir3_asm.c41 kernel->base.local_size[0] = v->local_size[0]; in ir3_asm_assemble()
42 kernel->base.local_size[1] = v->local_size[1]; in ir3_asm_assemble()
43 kernel->base.local_size[2] = v->local_size[2]; in ir3_asm_assemble()
Da4xx.c287 const unsigned *local_size = kernel->local_size; in a4xx_emit_grid() local
299 A4XX_HLSQ_CL_NDRANGE_0_LOCALSIZEX(local_size[0] - 1) | in a4xx_emit_grid()
300 A4XX_HLSQ_CL_NDRANGE_0_LOCALSIZEY(local_size[1] - 1) | in a4xx_emit_grid()
301 A4XX_HLSQ_CL_NDRANGE_0_LOCALSIZEZ(local_size[2] - 1)); in a4xx_emit_grid()
303 A4XX_HLSQ_CL_NDRANGE_1_SIZE_X(local_size[0] * num_groups[0])); in a4xx_emit_grid()
306 A4XX_HLSQ_CL_NDRANGE_3_SIZE_Y(local_size[1] * num_groups[1])); in a4xx_emit_grid()
309 A4XX_HLSQ_CL_NDRANGE_5_SIZE_Z(local_size[2] * num_groups[2])); in a4xx_emit_grid()
Da6xx.c413 const unsigned *local_size = kernel->local_size; in a6xx_emit_grid() local
425 A6XX_HLSQ_CS_NDRANGE_0_LOCALSIZEX(local_size[0] - 1) | in a6xx_emit_grid()
426 A6XX_HLSQ_CS_NDRANGE_0_LOCALSIZEY(local_size[1] - 1) | in a6xx_emit_grid()
427 A6XX_HLSQ_CS_NDRANGE_0_LOCALSIZEZ(local_size[2] - 1)); in a6xx_emit_grid()
429 A6XX_HLSQ_CS_NDRANGE_1_GLOBALSIZE_X(local_size[0] * num_groups[0])); in a6xx_emit_grid()
432 A6XX_HLSQ_CS_NDRANGE_3_GLOBALSIZE_Y(local_size[1] * num_groups[1])); in a6xx_emit_grid()
435 A6XX_HLSQ_CS_NDRANGE_5_GLOBALSIZE_Z(local_size[2] * num_groups[2])); in a6xx_emit_grid()
/third_party/mesa3d/src/gallium/drivers/freedreno/a6xx/
Dfd6_compute.c151 const unsigned *local_size = in fd6_launch_grid() local
158 A6XX_HLSQ_CS_NDRANGE_0_LOCALSIZEX(local_size[0] - 1) | in fd6_launch_grid()
159 A6XX_HLSQ_CS_NDRANGE_0_LOCALSIZEY(local_size[1] - 1) | in fd6_launch_grid()
160 A6XX_HLSQ_CS_NDRANGE_0_LOCALSIZEZ(local_size[2] - 1)); in fd6_launch_grid()
162 A6XX_HLSQ_CS_NDRANGE_1_GLOBALSIZE_X(local_size[0] * num_groups[0])); in fd6_launch_grid()
165 A6XX_HLSQ_CS_NDRANGE_3_GLOBALSIZE_Y(local_size[1] * num_groups[1])); in fd6_launch_grid()
168 A6XX_HLSQ_CS_NDRANGE_5_GLOBALSIZE_Z(local_size[2] * num_groups[2])); in fd6_launch_grid()
186 A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEX(local_size[0] - 1) | in fd6_launch_grid()
187 A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEY(local_size[1] - 1) | in fd6_launch_grid()
188 A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEZ(local_size[2] - 1)); in fd6_launch_grid()
/third_party/vk-gl-cts/external/openglcts/modules/gles31/
Des31cComputeShaderTests.cpp746 std::string GenSource(const uvec3& local_size, const uvec3& num_groups) in GenSource() argument
748 const uvec3 global_size = local_size * num_groups; in GenSource()
750 ss << NL "layout(local_size_x = " << local_size.x() << ", local_size_y = " << local_size.y() in GenSource()
751 …<< ", local_size_z = " << local_size.z() << ") in;" NL "const uvec3 kGlobalSize = uvec3(" << globa… in GenSource()
786 bool RunIteration(const uvec3& local_size, const uvec3& num_groups, bool dispatch_indirect) in RunIteration() argument
790 m_program = CreateComputeProgram(GenSource(local_size, num_groups)); in RunIteration()
804 …local_size.x() * num_groups.x() * local_size.y() * num_groups.y() * local_size.z() * num_groups.z(… in RunIteration()
847 for (GLuint z = 0; z < local_size.z() * num_groups.z(); ++z) in RunIteration()
849 for (GLuint y = 0; y < local_size.y() * num_groups.y(); ++y) in RunIteration()
851 for (GLuint x = 0; x < local_size.x() * num_groups.x(); ++x) in RunIteration()
[all …]
/third_party/mesa3d/src/microsoft/clc/
Dcompute_test.cpp487 if (!shader.dxil->metadata.local_size[0]) in run_shader_with_raw_args()
488 conf.local_size[0] = compile_args.x; in run_shader_with_raw_args()
490 conf.local_size[0] = shader.dxil->metadata.local_size[0]; in run_shader_with_raw_args()
492 if (!shader.dxil->metadata.local_size[1]) in run_shader_with_raw_args()
493 conf.local_size[1] = compile_args.y; in run_shader_with_raw_args()
495 conf.local_size[1] = shader.dxil->metadata.local_size[1]; in run_shader_with_raw_args()
497 if (!shader.dxil->metadata.local_size[2]) in run_shader_with_raw_args()
498 conf.local_size[2] = compile_args.z; in run_shader_with_raw_args()
500 conf.local_size[2] = shader.dxil->metadata.local_size[2]; in run_shader_with_raw_args()
502 if (compile_args.x % conf.local_size[0] || in run_shader_with_raw_args()
[all …]
/third_party/vk-gl-cts/external/openglcts/modules/gl/
Dgl4cComputeShaderTests.cpp744 std::string GenSource(const uvec3& local_size, const uvec3& num_groups) in GenSource() argument
746 const uvec3 global_size = local_size * num_groups; in GenSource()
748 ss << NL "layout(local_size_x = " << local_size.x() << ", local_size_y = " << local_size.y() in GenSource()
749 …<< ", local_size_z = " << local_size.z() << ") in;" NL "const uvec3 kGlobalSize = uvec3(" << globa… in GenSource()
784 bool RunIteration(const uvec3& local_size, const uvec3& num_groups, bool dispatch_indirect) in RunIteration() argument
788 m_program = CreateComputeProgram(GenSource(local_size, num_groups)); in RunIteration()
811 …local_size.x() * num_groups.x() * local_size.y() * num_groups.y() * local_size.z() * num_groups.z(… in RunIteration()
852 for (GLuint z = 0; z < local_size.z() * num_groups.z(); ++z) in RunIteration()
854 for (GLuint y = 0; y < local_size.y() * num_groups.y(); ++y) in RunIteration()
856 for (GLuint x = 0; x < local_size.x() * num_groups.x(); ++x) in RunIteration()
[all …]
/third_party/mesa3d/src/freedreno/ir3/
Dir3_assembler.c60 v->local_size[0] = v->local_size[1] = v->local_size[2] = 1; in ir3_parse_asm()
/third_party/mesa3d/src/mapi/glapi/gen/
Dgl_x86-64_asm.py42 def local_size(registers): function
56 adjust_stack = local_size(registers)
67 adjust_stack = local_size(registers)
/third_party/glslang/Test/baseResults/
DnegativeArraySize.comp.out7 local_size = (1, 1, 1)
18 local_size = (1, 1, 1)
D430.comp.out2 ERROR: 0:4: 'local_size' : cannot change previously set size
3 ERROR: 0:5: 'local_size' : too large; see gl_MaxComputeWorkGroupSize
11 ERROR: 0:52: 'local_size' : cannot change previously set size
12 ERROR: 0:54: 'local_size' : can only apply to 'in'
13 ERROR: 0:54: 'local_size' : can only apply to 'in'
14 ERROR: 0:54: 'local_size' : can only apply to 'in'
23 local_size = (2, 1, 4096)
155 local_size = (2, 1, 4096)
D450.comp.out9 local_size = (1, 1, 1)
20 local_size = (1, 1, 1)
Dhlsl.localStructuredBuffer.comp.out7 local_size = (1, 1, 1)
23 local_size = (1, 1, 1)
/third_party/mesa3d/src/compiler/glsl/
Dast_type.cpp423 if (q.flags.q.local_size & (1 << i)) { in merge_qualifier()
424 if (this->local_size[i] in merge_qualifier()
426 this->local_size[i]->merge_qualifier(q.local_size[i]); in merge_qualifier()
428 this->local_size[i] = q.local_size[i]; in merge_qualifier()
659 valid_in_mask.flags.q.local_size = 7; in validate_in_qualifier()
775 if (state->in_qualifier->flags.q.local_size) { in merge_into_in_qualifier()
777 state->in_qualifier->local_size); in merge_into_in_qualifier()
778 state->in_qualifier->flags.q.local_size = 0; in merge_into_in_qualifier()
780 state->in_qualifier->local_size[i] = NULL; in merge_into_in_qualifier()
876 Q(local_size); in validate_flags()

123456