1 /*
2 * Copyright © 2010 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 */
23
24 #include "compiler/glsl/ir.h"
25 #include "brw_fs.h"
26 #include "brw_nir.h"
27 #include "brw_eu.h"
28 #include "nir_search_helpers.h"
29 #include "util/u_math.h"
30 #include "util/bitscan.h"
31
32 using namespace brw;
33
34 void
emit_nir_code()35 fs_visitor::emit_nir_code()
36 {
37 emit_shader_float_controls_execution_mode();
38
39 /* emit the arrays used for inputs and outputs - load/store intrinsics will
40 * be converted to reads/writes of these arrays
41 */
42 nir_setup_outputs();
43 nir_setup_uniforms();
44 nir_emit_system_values();
45 last_scratch = ALIGN(nir->scratch_size, 4) * dispatch_width;
46
47 nir_emit_impl(nir_shader_get_entrypoint((nir_shader *)nir));
48 }
49
50 void
nir_setup_outputs()51 fs_visitor::nir_setup_outputs()
52 {
53 if (stage == MESA_SHADER_TESS_CTRL || stage == MESA_SHADER_FRAGMENT)
54 return;
55
56 unsigned vec4s[VARYING_SLOT_TESS_MAX] = { 0, };
57
58 /* Calculate the size of output registers in a separate pass, before
59 * allocating them. With ARB_enhanced_layouts, multiple output variables
60 * may occupy the same slot, but have different type sizes.
61 */
62 nir_foreach_shader_out_variable(var, nir) {
63 const int loc = var->data.driver_location;
64 const unsigned var_vec4s =
65 var->data.compact ? DIV_ROUND_UP(glsl_get_length(var->type), 4)
66 : type_size_vec4(var->type, true);
67 vec4s[loc] = MAX2(vec4s[loc], var_vec4s);
68 }
69
70 for (unsigned loc = 0; loc < ARRAY_SIZE(vec4s);) {
71 if (vec4s[loc] == 0) {
72 loc++;
73 continue;
74 }
75
76 unsigned reg_size = vec4s[loc];
77
78 /* Check if there are any ranges that start within this range and extend
79 * past it. If so, include them in this allocation.
80 */
81 for (unsigned i = 1; i < reg_size; i++) {
82 assert(i + loc < ARRAY_SIZE(vec4s));
83 reg_size = MAX2(vec4s[i + loc] + i, reg_size);
84 }
85
86 fs_reg reg = bld.vgrf(BRW_REGISTER_TYPE_F, 4 * reg_size);
87 for (unsigned i = 0; i < reg_size; i++) {
88 assert(loc + i < ARRAY_SIZE(outputs));
89 outputs[loc + i] = offset(reg, bld, 4 * i);
90 }
91
92 loc += reg_size;
93 }
94 }
95
96 void
nir_setup_uniforms()97 fs_visitor::nir_setup_uniforms()
98 {
99 /* Only the first compile gets to set up uniforms. */
100 if (push_constant_loc) {
101 assert(pull_constant_loc);
102 return;
103 }
104
105 uniforms = nir->num_uniforms / 4;
106
107 if (stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL) {
108 /* Add uniforms for builtins after regular NIR uniforms. */
109 assert(uniforms == prog_data->nr_params);
110
111 uint32_t *param;
112 if (nir->info.cs.local_size_variable &&
113 compiler->lower_variable_group_size) {
114 param = brw_stage_prog_data_add_params(prog_data, 3);
115 for (unsigned i = 0; i < 3; i++) {
116 param[i] = (BRW_PARAM_BUILTIN_WORK_GROUP_SIZE_X + i);
117 group_size[i] = fs_reg(UNIFORM, uniforms++, BRW_REGISTER_TYPE_UD);
118 }
119 }
120
121 /* Subgroup ID must be the last uniform on the list. This will make
122 * easier later to split between cross thread and per thread
123 * uniforms.
124 */
125 param = brw_stage_prog_data_add_params(prog_data, 1);
126 *param = BRW_PARAM_BUILTIN_SUBGROUP_ID;
127 subgroup_id = fs_reg(UNIFORM, uniforms++, BRW_REGISTER_TYPE_UD);
128 }
129 }
130
131 static bool
emit_system_values_block(nir_block * block,fs_visitor * v)132 emit_system_values_block(nir_block *block, fs_visitor *v)
133 {
134 fs_reg *reg;
135
136 nir_foreach_instr(instr, block) {
137 if (instr->type != nir_instr_type_intrinsic)
138 continue;
139
140 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
141 switch (intrin->intrinsic) {
142 case nir_intrinsic_load_vertex_id:
143 case nir_intrinsic_load_base_vertex:
144 unreachable("should be lowered by nir_lower_system_values().");
145
146 case nir_intrinsic_load_vertex_id_zero_base:
147 case nir_intrinsic_load_is_indexed_draw:
148 case nir_intrinsic_load_first_vertex:
149 case nir_intrinsic_load_instance_id:
150 case nir_intrinsic_load_base_instance:
151 case nir_intrinsic_load_draw_id:
152 unreachable("should be lowered by brw_nir_lower_vs_inputs().");
153
154 case nir_intrinsic_load_invocation_id:
155 if (v->stage == MESA_SHADER_TESS_CTRL)
156 break;
157 assert(v->stage == MESA_SHADER_GEOMETRY);
158 reg = &v->nir_system_values[SYSTEM_VALUE_INVOCATION_ID];
159 if (reg->file == BAD_FILE) {
160 const fs_builder abld = v->bld.annotate("gl_InvocationID", NULL);
161 fs_reg g1(retype(brw_vec8_grf(1, 0), BRW_REGISTER_TYPE_UD));
162 fs_reg iid = abld.vgrf(BRW_REGISTER_TYPE_UD, 1);
163 abld.SHR(iid, g1, brw_imm_ud(27u));
164 *reg = iid;
165 }
166 break;
167
168 case nir_intrinsic_load_sample_pos:
169 assert(v->stage == MESA_SHADER_FRAGMENT);
170 reg = &v->nir_system_values[SYSTEM_VALUE_SAMPLE_POS];
171 if (reg->file == BAD_FILE)
172 *reg = *v->emit_samplepos_setup();
173 break;
174
175 case nir_intrinsic_load_sample_id:
176 assert(v->stage == MESA_SHADER_FRAGMENT);
177 reg = &v->nir_system_values[SYSTEM_VALUE_SAMPLE_ID];
178 if (reg->file == BAD_FILE)
179 *reg = *v->emit_sampleid_setup();
180 break;
181
182 case nir_intrinsic_load_sample_mask_in:
183 assert(v->stage == MESA_SHADER_FRAGMENT);
184 assert(v->devinfo->gen >= 7);
185 reg = &v->nir_system_values[SYSTEM_VALUE_SAMPLE_MASK_IN];
186 if (reg->file == BAD_FILE)
187 *reg = *v->emit_samplemaskin_setup();
188 break;
189
190 case nir_intrinsic_load_work_group_id:
191 assert(v->stage == MESA_SHADER_COMPUTE ||
192 v->stage == MESA_SHADER_KERNEL);
193 reg = &v->nir_system_values[SYSTEM_VALUE_WORK_GROUP_ID];
194 if (reg->file == BAD_FILE)
195 *reg = *v->emit_cs_work_group_id_setup();
196 break;
197
198 case nir_intrinsic_load_helper_invocation:
199 assert(v->stage == MESA_SHADER_FRAGMENT);
200 reg = &v->nir_system_values[SYSTEM_VALUE_HELPER_INVOCATION];
201 if (reg->file == BAD_FILE) {
202 const fs_builder abld =
203 v->bld.annotate("gl_HelperInvocation", NULL);
204
205 /* On Gen6+ (gl_HelperInvocation is only exposed on Gen7+) the
206 * pixel mask is in g1.7 of the thread payload.
207 *
208 * We move the per-channel pixel enable bit to the low bit of each
209 * channel by shifting the byte containing the pixel mask by the
210 * vector immediate 0x76543210UV.
211 *
212 * The region of <1,8,0> reads only 1 byte (the pixel masks for
213 * subspans 0 and 1) in SIMD8 and an additional byte (the pixel
214 * masks for 2 and 3) in SIMD16.
215 */
216 fs_reg shifted = abld.vgrf(BRW_REGISTER_TYPE_UW, 1);
217
218 for (unsigned i = 0; i < DIV_ROUND_UP(v->dispatch_width, 16); i++) {
219 const fs_builder hbld = abld.group(MIN2(16, v->dispatch_width), i);
220 hbld.SHR(offset(shifted, hbld, i),
221 stride(retype(brw_vec1_grf(1 + i, 7),
222 BRW_REGISTER_TYPE_UB),
223 1, 8, 0),
224 brw_imm_v(0x76543210));
225 }
226
227 /* A set bit in the pixel mask means the channel is enabled, but
228 * that is the opposite of gl_HelperInvocation so we need to invert
229 * the mask.
230 *
231 * The negate source-modifier bit of logical instructions on Gen8+
232 * performs 1's complement negation, so we can use that instead of
233 * a NOT instruction.
234 */
235 fs_reg inverted = negate(shifted);
236 if (v->devinfo->gen < 8) {
237 inverted = abld.vgrf(BRW_REGISTER_TYPE_UW);
238 abld.NOT(inverted, shifted);
239 }
240
241 /* We then resolve the 0/1 result to 0/~0 boolean values by ANDing
242 * with 1 and negating.
243 */
244 fs_reg anded = abld.vgrf(BRW_REGISTER_TYPE_UD, 1);
245 abld.AND(anded, inverted, brw_imm_uw(1));
246
247 fs_reg dst = abld.vgrf(BRW_REGISTER_TYPE_D, 1);
248 abld.MOV(dst, negate(retype(anded, BRW_REGISTER_TYPE_D)));
249 *reg = dst;
250 }
251 break;
252
253 default:
254 break;
255 }
256 }
257
258 return true;
259 }
260
261 void
nir_emit_system_values()262 fs_visitor::nir_emit_system_values()
263 {
264 nir_system_values = ralloc_array(mem_ctx, fs_reg, SYSTEM_VALUE_MAX);
265 for (unsigned i = 0; i < SYSTEM_VALUE_MAX; i++) {
266 nir_system_values[i] = fs_reg();
267 }
268
269 /* Always emit SUBGROUP_INVOCATION. Dead code will clean it up if we
270 * never end up using it.
271 */
272 {
273 const fs_builder abld = bld.annotate("gl_SubgroupInvocation", NULL);
274 fs_reg ® = nir_system_values[SYSTEM_VALUE_SUBGROUP_INVOCATION];
275 reg = abld.vgrf(BRW_REGISTER_TYPE_UW);
276
277 const fs_builder allbld8 = abld.group(8, 0).exec_all();
278 allbld8.MOV(reg, brw_imm_v(0x76543210));
279 if (dispatch_width > 8)
280 allbld8.ADD(byte_offset(reg, 16), reg, brw_imm_uw(8u));
281 if (dispatch_width > 16) {
282 const fs_builder allbld16 = abld.group(16, 0).exec_all();
283 allbld16.ADD(byte_offset(reg, 32), reg, brw_imm_uw(16u));
284 }
285 }
286
287 nir_function_impl *impl = nir_shader_get_entrypoint((nir_shader *)nir);
288 nir_foreach_block(block, impl)
289 emit_system_values_block(block, this);
290 }
291
292 /*
293 * Returns a type based on a reference_type (word, float, half-float) and a
294 * given bit_size.
295 *
296 * Reference BRW_REGISTER_TYPE are HF,F,DF,W,D,UW,UD.
297 *
298 * @FIXME: 64-bit return types are always DF on integer types to maintain
299 * compability with uses of DF previously to the introduction of int64
300 * support.
301 */
302 static brw_reg_type
brw_reg_type_from_bit_size(const unsigned bit_size,const brw_reg_type reference_type)303 brw_reg_type_from_bit_size(const unsigned bit_size,
304 const brw_reg_type reference_type)
305 {
306 switch(reference_type) {
307 case BRW_REGISTER_TYPE_HF:
308 case BRW_REGISTER_TYPE_F:
309 case BRW_REGISTER_TYPE_DF:
310 switch(bit_size) {
311 case 16:
312 return BRW_REGISTER_TYPE_HF;
313 case 32:
314 return BRW_REGISTER_TYPE_F;
315 case 64:
316 return BRW_REGISTER_TYPE_DF;
317 default:
318 unreachable("Invalid bit size");
319 }
320 case BRW_REGISTER_TYPE_B:
321 case BRW_REGISTER_TYPE_W:
322 case BRW_REGISTER_TYPE_D:
323 case BRW_REGISTER_TYPE_Q:
324 switch(bit_size) {
325 case 8:
326 return BRW_REGISTER_TYPE_B;
327 case 16:
328 return BRW_REGISTER_TYPE_W;
329 case 32:
330 return BRW_REGISTER_TYPE_D;
331 case 64:
332 return BRW_REGISTER_TYPE_Q;
333 default:
334 unreachable("Invalid bit size");
335 }
336 case BRW_REGISTER_TYPE_UB:
337 case BRW_REGISTER_TYPE_UW:
338 case BRW_REGISTER_TYPE_UD:
339 case BRW_REGISTER_TYPE_UQ:
340 switch(bit_size) {
341 case 8:
342 return BRW_REGISTER_TYPE_UB;
343 case 16:
344 return BRW_REGISTER_TYPE_UW;
345 case 32:
346 return BRW_REGISTER_TYPE_UD;
347 case 64:
348 return BRW_REGISTER_TYPE_UQ;
349 default:
350 unreachable("Invalid bit size");
351 }
352 default:
353 unreachable("Unknown type");
354 }
355 }
356
357 void
nir_emit_impl(nir_function_impl * impl)358 fs_visitor::nir_emit_impl(nir_function_impl *impl)
359 {
360 nir_locals = ralloc_array(mem_ctx, fs_reg, impl->reg_alloc);
361 for (unsigned i = 0; i < impl->reg_alloc; i++) {
362 nir_locals[i] = fs_reg();
363 }
364
365 foreach_list_typed(nir_register, reg, node, &impl->registers) {
366 unsigned array_elems =
367 reg->num_array_elems == 0 ? 1 : reg->num_array_elems;
368 unsigned size = array_elems * reg->num_components;
369 const brw_reg_type reg_type = reg->bit_size == 8 ? BRW_REGISTER_TYPE_B :
370 brw_reg_type_from_bit_size(reg->bit_size, BRW_REGISTER_TYPE_F);
371 nir_locals[reg->index] = bld.vgrf(reg_type, size);
372 }
373
374 nir_ssa_values = reralloc(mem_ctx, nir_ssa_values, fs_reg,
375 impl->ssa_alloc);
376
377 nir_emit_cf_list(&impl->body);
378 }
379
380 void
nir_emit_cf_list(exec_list * list)381 fs_visitor::nir_emit_cf_list(exec_list *list)
382 {
383 exec_list_validate(list);
384 foreach_list_typed(nir_cf_node, node, node, list) {
385 switch (node->type) {
386 case nir_cf_node_if:
387 nir_emit_if(nir_cf_node_as_if(node));
388 break;
389
390 case nir_cf_node_loop:
391 nir_emit_loop(nir_cf_node_as_loop(node));
392 break;
393
394 case nir_cf_node_block:
395 nir_emit_block(nir_cf_node_as_block(node));
396 break;
397
398 default:
399 unreachable("Invalid CFG node block");
400 }
401 }
402 }
403
404 void
nir_emit_if(nir_if * if_stmt)405 fs_visitor::nir_emit_if(nir_if *if_stmt)
406 {
407 bool invert;
408 fs_reg cond_reg;
409
410 /* If the condition has the form !other_condition, use other_condition as
411 * the source, but invert the predicate on the if instruction.
412 */
413 nir_alu_instr *cond = nir_src_as_alu_instr(if_stmt->condition);
414 if (cond != NULL && cond->op == nir_op_inot) {
415 invert = true;
416 cond_reg = get_nir_src(cond->src[0].src);
417 } else {
418 invert = false;
419 cond_reg = get_nir_src(if_stmt->condition);
420 }
421
422 /* first, put the condition into f0 */
423 fs_inst *inst = bld.MOV(bld.null_reg_d(),
424 retype(cond_reg, BRW_REGISTER_TYPE_D));
425 inst->conditional_mod = BRW_CONDITIONAL_NZ;
426
427 bld.IF(BRW_PREDICATE_NORMAL)->predicate_inverse = invert;
428
429 nir_emit_cf_list(&if_stmt->then_list);
430
431 if (!nir_cf_list_is_empty_block(&if_stmt->else_list)) {
432 bld.emit(BRW_OPCODE_ELSE);
433 nir_emit_cf_list(&if_stmt->else_list);
434 }
435
436 bld.emit(BRW_OPCODE_ENDIF);
437
438 if (devinfo->gen < 7)
439 limit_dispatch_width(16, "Non-uniform control flow unsupported "
440 "in SIMD32 mode.");
441 }
442
443 void
nir_emit_loop(nir_loop * loop)444 fs_visitor::nir_emit_loop(nir_loop *loop)
445 {
446 bld.emit(BRW_OPCODE_DO);
447
448 nir_emit_cf_list(&loop->body);
449
450 bld.emit(BRW_OPCODE_WHILE);
451
452 if (devinfo->gen < 7)
453 limit_dispatch_width(16, "Non-uniform control flow unsupported "
454 "in SIMD32 mode.");
455 }
456
457 void
nir_emit_block(nir_block * block)458 fs_visitor::nir_emit_block(nir_block *block)
459 {
460 nir_foreach_instr(instr, block) {
461 nir_emit_instr(instr);
462 }
463 }
464
465 void
nir_emit_instr(nir_instr * instr)466 fs_visitor::nir_emit_instr(nir_instr *instr)
467 {
468 const fs_builder abld = bld.annotate(NULL, instr);
469
470 switch (instr->type) {
471 case nir_instr_type_alu:
472 nir_emit_alu(abld, nir_instr_as_alu(instr), true);
473 break;
474
475 case nir_instr_type_deref:
476 unreachable("All derefs should've been lowered");
477 break;
478
479 case nir_instr_type_intrinsic:
480 switch (stage) {
481 case MESA_SHADER_VERTEX:
482 nir_emit_vs_intrinsic(abld, nir_instr_as_intrinsic(instr));
483 break;
484 case MESA_SHADER_TESS_CTRL:
485 nir_emit_tcs_intrinsic(abld, nir_instr_as_intrinsic(instr));
486 break;
487 case MESA_SHADER_TESS_EVAL:
488 nir_emit_tes_intrinsic(abld, nir_instr_as_intrinsic(instr));
489 break;
490 case MESA_SHADER_GEOMETRY:
491 nir_emit_gs_intrinsic(abld, nir_instr_as_intrinsic(instr));
492 break;
493 case MESA_SHADER_FRAGMENT:
494 nir_emit_fs_intrinsic(abld, nir_instr_as_intrinsic(instr));
495 break;
496 case MESA_SHADER_COMPUTE:
497 case MESA_SHADER_KERNEL:
498 nir_emit_cs_intrinsic(abld, nir_instr_as_intrinsic(instr));
499 break;
500 default:
501 unreachable("unsupported shader stage");
502 }
503 break;
504
505 case nir_instr_type_tex:
506 nir_emit_texture(abld, nir_instr_as_tex(instr));
507 break;
508
509 case nir_instr_type_load_const:
510 nir_emit_load_const(abld, nir_instr_as_load_const(instr));
511 break;
512
513 case nir_instr_type_ssa_undef:
514 /* We create a new VGRF for undefs on every use (by handling
515 * them in get_nir_src()), rather than for each definition.
516 * This helps register coalescing eliminate MOVs from undef.
517 */
518 break;
519
520 case nir_instr_type_jump:
521 nir_emit_jump(abld, nir_instr_as_jump(instr));
522 break;
523
524 default:
525 unreachable("unknown instruction type");
526 }
527 }
528
529 /**
530 * Recognizes a parent instruction of nir_op_extract_* and changes the type to
531 * match instr.
532 */
533 bool
optimize_extract_to_float(nir_alu_instr * instr,const fs_reg & result)534 fs_visitor::optimize_extract_to_float(nir_alu_instr *instr,
535 const fs_reg &result)
536 {
537 if (!instr->src[0].src.is_ssa ||
538 !instr->src[0].src.ssa->parent_instr)
539 return false;
540
541 if (instr->src[0].src.ssa->parent_instr->type != nir_instr_type_alu)
542 return false;
543
544 nir_alu_instr *src0 =
545 nir_instr_as_alu(instr->src[0].src.ssa->parent_instr);
546
547 if (src0->op != nir_op_extract_u8 && src0->op != nir_op_extract_u16 &&
548 src0->op != nir_op_extract_i8 && src0->op != nir_op_extract_i16)
549 return false;
550
551 unsigned element = nir_src_as_uint(src0->src[1].src);
552
553 /* Element type to extract.*/
554 const brw_reg_type type = brw_int_type(
555 src0->op == nir_op_extract_u16 || src0->op == nir_op_extract_i16 ? 2 : 1,
556 src0->op == nir_op_extract_i16 || src0->op == nir_op_extract_i8);
557
558 fs_reg op0 = get_nir_src(src0->src[0].src);
559 op0.type = brw_type_for_nir_type(devinfo,
560 (nir_alu_type)(nir_op_infos[src0->op].input_types[0] |
561 nir_src_bit_size(src0->src[0].src)));
562 op0 = offset(op0, bld, src0->src[0].swizzle[0]);
563
564 bld.MOV(result, subscript(op0, type, element));
565 return true;
566 }
567
568 bool
optimize_frontfacing_ternary(nir_alu_instr * instr,const fs_reg & result)569 fs_visitor::optimize_frontfacing_ternary(nir_alu_instr *instr,
570 const fs_reg &result)
571 {
572 nir_intrinsic_instr *src0 = nir_src_as_intrinsic(instr->src[0].src);
573 if (src0 == NULL || src0->intrinsic != nir_intrinsic_load_front_face)
574 return false;
575
576 if (!nir_src_is_const(instr->src[1].src) ||
577 !nir_src_is_const(instr->src[2].src))
578 return false;
579
580 const float value1 = nir_src_as_float(instr->src[1].src);
581 const float value2 = nir_src_as_float(instr->src[2].src);
582 if (fabsf(value1) != 1.0f || fabsf(value2) != 1.0f)
583 return false;
584
585 /* nir_opt_algebraic should have gotten rid of bcsel(b, a, a) */
586 assert(value1 == -value2);
587
588 fs_reg tmp = vgrf(glsl_type::int_type);
589
590 if (devinfo->gen >= 12) {
591 /* Bit 15 of g1.1 is 0 if the polygon is front facing. */
592 fs_reg g1 = fs_reg(retype(brw_vec1_grf(1, 1), BRW_REGISTER_TYPE_W));
593
594 /* For (gl_FrontFacing ? 1.0 : -1.0), emit:
595 *
596 * or(8) tmp.1<2>W g0.0<0,1,0>W 0x00003f80W
597 * and(8) dst<1>D tmp<8,8,1>D 0xbf800000D
598 *
599 * and negate the result for (gl_FrontFacing ? -1.0 : 1.0).
600 */
601 bld.OR(subscript(tmp, BRW_REGISTER_TYPE_W, 1),
602 g1, brw_imm_uw(0x3f80));
603
604 if (value1 == -1.0f)
605 bld.MOV(tmp, negate(tmp));
606
607 } else if (devinfo->gen >= 6) {
608 /* Bit 15 of g0.0 is 0 if the polygon is front facing. */
609 fs_reg g0 = fs_reg(retype(brw_vec1_grf(0, 0), BRW_REGISTER_TYPE_W));
610
611 /* For (gl_FrontFacing ? 1.0 : -1.0), emit:
612 *
613 * or(8) tmp.1<2>W g0.0<0,1,0>W 0x00003f80W
614 * and(8) dst<1>D tmp<8,8,1>D 0xbf800000D
615 *
616 * and negate g0.0<0,1,0>W for (gl_FrontFacing ? -1.0 : 1.0).
617 *
618 * This negation looks like it's safe in practice, because bits 0:4 will
619 * surely be TRIANGLES
620 */
621
622 if (value1 == -1.0f) {
623 g0.negate = true;
624 }
625
626 bld.OR(subscript(tmp, BRW_REGISTER_TYPE_W, 1),
627 g0, brw_imm_uw(0x3f80));
628 } else {
629 /* Bit 31 of g1.6 is 0 if the polygon is front facing. */
630 fs_reg g1_6 = fs_reg(retype(brw_vec1_grf(1, 6), BRW_REGISTER_TYPE_D));
631
632 /* For (gl_FrontFacing ? 1.0 : -1.0), emit:
633 *
634 * or(8) tmp<1>D g1.6<0,1,0>D 0x3f800000D
635 * and(8) dst<1>D tmp<8,8,1>D 0xbf800000D
636 *
637 * and negate g1.6<0,1,0>D for (gl_FrontFacing ? -1.0 : 1.0).
638 *
639 * This negation looks like it's safe in practice, because bits 0:4 will
640 * surely be TRIANGLES
641 */
642
643 if (value1 == -1.0f) {
644 g1_6.negate = true;
645 }
646
647 bld.OR(tmp, g1_6, brw_imm_d(0x3f800000));
648 }
649 bld.AND(retype(result, BRW_REGISTER_TYPE_D), tmp, brw_imm_d(0xbf800000));
650
651 return true;
652 }
653
654 static void
emit_find_msb_using_lzd(const fs_builder & bld,const fs_reg & result,const fs_reg & src,bool is_signed)655 emit_find_msb_using_lzd(const fs_builder &bld,
656 const fs_reg &result,
657 const fs_reg &src,
658 bool is_signed)
659 {
660 fs_inst *inst;
661 fs_reg temp = src;
662
663 if (is_signed) {
664 /* LZD of an absolute value source almost always does the right
665 * thing. There are two problem values:
666 *
667 * * 0x80000000. Since abs(0x80000000) == 0x80000000, LZD returns
668 * 0. However, findMSB(int(0x80000000)) == 30.
669 *
670 * * 0xffffffff. Since abs(0xffffffff) == 1, LZD returns
671 * 31. Section 8.8 (Integer Functions) of the GLSL 4.50 spec says:
672 *
673 * For a value of zero or negative one, -1 will be returned.
674 *
675 * * Negative powers of two. LZD(abs(-(1<<x))) returns x, but
676 * findMSB(-(1<<x)) should return x-1.
677 *
678 * For all negative number cases, including 0x80000000 and
679 * 0xffffffff, the correct value is obtained from LZD if instead of
680 * negating the (already negative) value the logical-not is used. A
681 * conditonal logical-not can be achieved in two instructions.
682 */
683 temp = bld.vgrf(BRW_REGISTER_TYPE_D);
684
685 bld.ASR(temp, src, brw_imm_d(31));
686 bld.XOR(temp, temp, src);
687 }
688
689 bld.LZD(retype(result, BRW_REGISTER_TYPE_UD),
690 retype(temp, BRW_REGISTER_TYPE_UD));
691
692 /* LZD counts from the MSB side, while GLSL's findMSB() wants the count
693 * from the LSB side. Subtract the result from 31 to convert the MSB
694 * count into an LSB count. If no bits are set, LZD will return 32.
695 * 31-32 = -1, which is exactly what findMSB() is supposed to return.
696 */
697 inst = bld.ADD(result, retype(result, BRW_REGISTER_TYPE_D), brw_imm_d(31));
698 inst->src[0].negate = true;
699 }
700
701 static brw_rnd_mode
brw_rnd_mode_from_nir_op(const nir_op op)702 brw_rnd_mode_from_nir_op (const nir_op op) {
703 switch (op) {
704 case nir_op_f2f16_rtz:
705 return BRW_RND_MODE_RTZ;
706 case nir_op_f2f16_rtne:
707 return BRW_RND_MODE_RTNE;
708 default:
709 unreachable("Operation doesn't support rounding mode");
710 }
711 }
712
713 static brw_rnd_mode
brw_rnd_mode_from_execution_mode(unsigned execution_mode)714 brw_rnd_mode_from_execution_mode(unsigned execution_mode)
715 {
716 if (nir_has_any_rounding_mode_rtne(execution_mode))
717 return BRW_RND_MODE_RTNE;
718 if (nir_has_any_rounding_mode_rtz(execution_mode))
719 return BRW_RND_MODE_RTZ;
720 return BRW_RND_MODE_UNSPECIFIED;
721 }
722
723 fs_reg
prepare_alu_destination_and_sources(const fs_builder & bld,nir_alu_instr * instr,fs_reg * op,bool need_dest)724 fs_visitor::prepare_alu_destination_and_sources(const fs_builder &bld,
725 nir_alu_instr *instr,
726 fs_reg *op,
727 bool need_dest)
728 {
729 fs_reg result =
730 need_dest ? get_nir_dest(instr->dest.dest) : bld.null_reg_ud();
731
732 result.type = brw_type_for_nir_type(devinfo,
733 (nir_alu_type)(nir_op_infos[instr->op].output_type |
734 nir_dest_bit_size(instr->dest.dest)));
735
736 assert(!instr->dest.saturate);
737
738 for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++) {
739 /* We don't lower to source modifiers so they should not exist. */
740 assert(!instr->src[i].abs);
741 assert(!instr->src[i].negate);
742
743 op[i] = get_nir_src(instr->src[i].src);
744 op[i].type = brw_type_for_nir_type(devinfo,
745 (nir_alu_type)(nir_op_infos[instr->op].input_types[i] |
746 nir_src_bit_size(instr->src[i].src)));
747 }
748
749 /* Move and vecN instrutions may still be vectored. Return the raw,
750 * vectored source and destination so that fs_visitor::nir_emit_alu can
751 * handle it. Other callers should not have to handle these kinds of
752 * instructions.
753 */
754 switch (instr->op) {
755 case nir_op_mov:
756 case nir_op_vec2:
757 case nir_op_vec3:
758 case nir_op_vec4:
759 case nir_op_vec8:
760 case nir_op_vec16:
761 return result;
762 default:
763 break;
764 }
765
766 /* At this point, we have dealt with any instruction that operates on
767 * more than a single channel. Therefore, we can just adjust the source
768 * and destination registers for that channel and emit the instruction.
769 */
770 unsigned channel = 0;
771 if (nir_op_infos[instr->op].output_size == 0) {
772 /* Since NIR is doing the scalarizing for us, we should only ever see
773 * vectorized operations with a single channel.
774 */
775 assert(util_bitcount(instr->dest.write_mask) == 1);
776 channel = ffs(instr->dest.write_mask) - 1;
777
778 result = offset(result, bld, channel);
779 }
780
781 for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++) {
782 assert(nir_op_infos[instr->op].input_sizes[i] < 2);
783 op[i] = offset(op[i], bld, instr->src[i].swizzle[channel]);
784 }
785
786 return result;
787 }
788
789 void
resolve_inot_sources(const fs_builder & bld,nir_alu_instr * instr,fs_reg * op)790 fs_visitor::resolve_inot_sources(const fs_builder &bld, nir_alu_instr *instr,
791 fs_reg *op)
792 {
793 for (unsigned i = 0; i < 2; i++) {
794 nir_alu_instr *inot_instr = nir_src_as_alu_instr(instr->src[i].src);
795
796 if (inot_instr != NULL && inot_instr->op == nir_op_inot) {
797 /* The source of the inot is now the source of instr. */
798 prepare_alu_destination_and_sources(bld, inot_instr, &op[i], false);
799
800 assert(!op[i].negate);
801 op[i].negate = true;
802 } else {
803 op[i] = resolve_source_modifiers(op[i]);
804 }
805 }
806 }
807
808 bool
try_emit_b2fi_of_inot(const fs_builder & bld,fs_reg result,nir_alu_instr * instr)809 fs_visitor::try_emit_b2fi_of_inot(const fs_builder &bld,
810 fs_reg result,
811 nir_alu_instr *instr)
812 {
813 if (devinfo->gen < 6 || devinfo->gen >= 12)
814 return false;
815
816 nir_alu_instr *inot_instr = nir_src_as_alu_instr(instr->src[0].src);
817
818 if (inot_instr == NULL || inot_instr->op != nir_op_inot)
819 return false;
820
821 /* HF is also possible as a destination on BDW+. For nir_op_b2i, the set
822 * of valid size-changing combinations is a bit more complex.
823 *
824 * The source restriction is just because I was lazy about generating the
825 * constant below.
826 */
827 if (nir_dest_bit_size(instr->dest.dest) != 32 ||
828 nir_src_bit_size(inot_instr->src[0].src) != 32)
829 return false;
830
831 /* b2[fi](inot(a)) maps a=0 => 1, a=-1 => 0. Since a can only be 0 or -1,
832 * this is float(1 + a).
833 */
834 fs_reg op;
835
836 prepare_alu_destination_and_sources(bld, inot_instr, &op, false);
837
838 /* Ignore the saturate modifier, if there is one. The result of the
839 * arithmetic can only be 0 or 1, so the clamping will do nothing anyway.
840 */
841 bld.ADD(result, op, brw_imm_d(1));
842
843 return true;
844 }
845
846 /**
847 * Emit code for nir_op_fsign possibly fused with a nir_op_fmul
848 *
849 * If \c instr is not the \c nir_op_fsign, then \c fsign_src is the index of
850 * the source of \c instr that is a \c nir_op_fsign.
851 */
852 void
emit_fsign(const fs_builder & bld,const nir_alu_instr * instr,fs_reg result,fs_reg * op,unsigned fsign_src)853 fs_visitor::emit_fsign(const fs_builder &bld, const nir_alu_instr *instr,
854 fs_reg result, fs_reg *op, unsigned fsign_src)
855 {
856 fs_inst *inst;
857
858 assert(instr->op == nir_op_fsign || instr->op == nir_op_fmul);
859 assert(fsign_src < nir_op_infos[instr->op].num_inputs);
860
861 if (instr->op != nir_op_fsign) {
862 const nir_alu_instr *const fsign_instr =
863 nir_src_as_alu_instr(instr->src[fsign_src].src);
864
865 /* op[fsign_src] has the nominal result of the fsign, and op[1 -
866 * fsign_src] has the other multiply source. This must be rearranged so
867 * that op[0] is the source of the fsign op[1] is the other multiply
868 * source.
869 */
870 if (fsign_src != 0)
871 op[1] = op[0];
872
873 op[0] = get_nir_src(fsign_instr->src[0].src);
874
875 const nir_alu_type t =
876 (nir_alu_type)(nir_op_infos[instr->op].input_types[0] |
877 nir_src_bit_size(fsign_instr->src[0].src));
878
879 op[0].type = brw_type_for_nir_type(devinfo, t);
880
881 unsigned channel = 0;
882 if (nir_op_infos[instr->op].output_size == 0) {
883 /* Since NIR is doing the scalarizing for us, we should only ever see
884 * vectorized operations with a single channel.
885 */
886 assert(util_bitcount(instr->dest.write_mask) == 1);
887 channel = ffs(instr->dest.write_mask) - 1;
888 }
889
890 op[0] = offset(op[0], bld, fsign_instr->src[0].swizzle[channel]);
891 }
892
893 if (type_sz(op[0].type) == 2) {
894 /* AND(val, 0x8000) gives the sign bit.
895 *
896 * Predicated OR ORs 1.0 (0x3c00) with the sign bit if val is not zero.
897 */
898 fs_reg zero = retype(brw_imm_uw(0), BRW_REGISTER_TYPE_HF);
899 bld.CMP(bld.null_reg_f(), op[0], zero, BRW_CONDITIONAL_NZ);
900
901 op[0].type = BRW_REGISTER_TYPE_UW;
902 result.type = BRW_REGISTER_TYPE_UW;
903 bld.AND(result, op[0], brw_imm_uw(0x8000u));
904
905 if (instr->op == nir_op_fsign)
906 inst = bld.OR(result, result, brw_imm_uw(0x3c00u));
907 else {
908 /* Use XOR here to get the result sign correct. */
909 inst = bld.XOR(result, result, retype(op[1], BRW_REGISTER_TYPE_UW));
910 }
911
912 inst->predicate = BRW_PREDICATE_NORMAL;
913 } else if (type_sz(op[0].type) == 4) {
914 /* AND(val, 0x80000000) gives the sign bit.
915 *
916 * Predicated OR ORs 1.0 (0x3f800000) with the sign bit if val is not
917 * zero.
918 */
919 bld.CMP(bld.null_reg_f(), op[0], brw_imm_f(0.0f), BRW_CONDITIONAL_NZ);
920
921 op[0].type = BRW_REGISTER_TYPE_UD;
922 result.type = BRW_REGISTER_TYPE_UD;
923 bld.AND(result, op[0], brw_imm_ud(0x80000000u));
924
925 if (instr->op == nir_op_fsign)
926 inst = bld.OR(result, result, brw_imm_ud(0x3f800000u));
927 else {
928 /* Use XOR here to get the result sign correct. */
929 inst = bld.XOR(result, result, retype(op[1], BRW_REGISTER_TYPE_UD));
930 }
931
932 inst->predicate = BRW_PREDICATE_NORMAL;
933 } else {
934 /* For doubles we do the same but we need to consider:
935 *
936 * - 2-src instructions can't operate with 64-bit immediates
937 * - The sign is encoded in the high 32-bit of each DF
938 * - We need to produce a DF result.
939 */
940
941 fs_reg zero = vgrf(glsl_type::double_type);
942 bld.MOV(zero, setup_imm_df(bld, 0.0));
943 bld.CMP(bld.null_reg_df(), op[0], zero, BRW_CONDITIONAL_NZ);
944
945 bld.MOV(result, zero);
946
947 fs_reg r = subscript(result, BRW_REGISTER_TYPE_UD, 1);
948 bld.AND(r, subscript(op[0], BRW_REGISTER_TYPE_UD, 1),
949 brw_imm_ud(0x80000000u));
950
951 if (instr->op == nir_op_fsign) {
952 set_predicate(BRW_PREDICATE_NORMAL,
953 bld.OR(r, r, brw_imm_ud(0x3ff00000u)));
954 } else {
955 /* This could be done better in some cases. If the scale is an
956 * immediate with the low 32-bits all 0, emitting a separate XOR and
957 * OR would allow an algebraic optimization to remove the OR. There
958 * are currently zero instances of fsign(double(x))*IMM in shader-db
959 * or any test suite, so it is hard to care at this time.
960 */
961 fs_reg result_int64 = retype(result, BRW_REGISTER_TYPE_UQ);
962 inst = bld.XOR(result_int64, result_int64,
963 retype(op[1], BRW_REGISTER_TYPE_UQ));
964 }
965 }
966 }
967
968 /**
969 * Deteremine whether sources of a nir_op_fmul can be fused with a nir_op_fsign
970 *
971 * Checks the operands of a \c nir_op_fmul to determine whether or not
972 * \c emit_fsign could fuse the multiplication with the \c sign() calculation.
973 *
974 * \param instr The multiplication instruction
975 *
976 * \param fsign_src The source of \c instr that may or may not be a
977 * \c nir_op_fsign
978 */
979 static bool
can_fuse_fmul_fsign(nir_alu_instr * instr,unsigned fsign_src)980 can_fuse_fmul_fsign(nir_alu_instr *instr, unsigned fsign_src)
981 {
982 assert(instr->op == nir_op_fmul);
983
984 nir_alu_instr *const fsign_instr =
985 nir_src_as_alu_instr(instr->src[fsign_src].src);
986
987 /* Rules:
988 *
989 * 1. instr->src[fsign_src] must be a nir_op_fsign.
990 * 2. The nir_op_fsign can only be used by this multiplication.
991 * 3. The source that is the nir_op_fsign does not have source modifiers.
992 * \c emit_fsign only examines the source modifiers of the source of the
993 * \c nir_op_fsign.
994 *
995 * The nir_op_fsign must also not have the saturate modifier, but steps
996 * have already been taken (in nir_opt_algebraic) to ensure that.
997 */
998 return fsign_instr != NULL && fsign_instr->op == nir_op_fsign &&
999 is_used_once(fsign_instr);
1000 }
1001
1002 void
nir_emit_alu(const fs_builder & bld,nir_alu_instr * instr,bool need_dest)1003 fs_visitor::nir_emit_alu(const fs_builder &bld, nir_alu_instr *instr,
1004 bool need_dest)
1005 {
1006 struct brw_wm_prog_key *fs_key = (struct brw_wm_prog_key *) this->key;
1007 fs_inst *inst;
1008 unsigned execution_mode =
1009 bld.shader->nir->info.float_controls_execution_mode;
1010
1011 fs_reg op[NIR_MAX_VEC_COMPONENTS];
1012 fs_reg result = prepare_alu_destination_and_sources(bld, instr, op, need_dest);
1013
1014 switch (instr->op) {
1015 case nir_op_mov:
1016 case nir_op_vec2:
1017 case nir_op_vec3:
1018 case nir_op_vec4:
1019 case nir_op_vec8:
1020 case nir_op_vec16: {
1021 fs_reg temp = result;
1022 bool need_extra_copy = false;
1023 for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++) {
1024 if (!instr->src[i].src.is_ssa &&
1025 instr->dest.dest.reg.reg == instr->src[i].src.reg.reg) {
1026 need_extra_copy = true;
1027 temp = bld.vgrf(result.type, 4);
1028 break;
1029 }
1030 }
1031
1032 for (unsigned i = 0; i < 4; i++) {
1033 if (!(instr->dest.write_mask & (1 << i)))
1034 continue;
1035
1036 if (instr->op == nir_op_mov) {
1037 bld.MOV(offset(temp, bld, i),
1038 offset(op[0], bld, instr->src[0].swizzle[i]));
1039 } else {
1040 bld.MOV(offset(temp, bld, i),
1041 offset(op[i], bld, instr->src[i].swizzle[0]));
1042 }
1043 }
1044
1045 /* In this case the source and destination registers were the same,
1046 * so we need to insert an extra set of moves in order to deal with
1047 * any swizzling.
1048 */
1049 if (need_extra_copy) {
1050 for (unsigned i = 0; i < 4; i++) {
1051 if (!(instr->dest.write_mask & (1 << i)))
1052 continue;
1053
1054 bld.MOV(offset(result, bld, i), offset(temp, bld, i));
1055 }
1056 }
1057 return;
1058 }
1059
1060 case nir_op_i2f32:
1061 case nir_op_u2f32:
1062 if (optimize_extract_to_float(instr, result))
1063 return;
1064 inst = bld.MOV(result, op[0]);
1065 break;
1066
1067 case nir_op_f2f16_rtne:
1068 case nir_op_f2f16_rtz:
1069 case nir_op_f2f16: {
1070 brw_rnd_mode rnd = BRW_RND_MODE_UNSPECIFIED;
1071
1072 if (nir_op_f2f16 == instr->op)
1073 rnd = brw_rnd_mode_from_execution_mode(execution_mode);
1074 else
1075 rnd = brw_rnd_mode_from_nir_op(instr->op);
1076
1077 if (BRW_RND_MODE_UNSPECIFIED != rnd)
1078 bld.emit(SHADER_OPCODE_RND_MODE, bld.null_reg_ud(), brw_imm_d(rnd));
1079
1080 /* In theory, it would be better to use BRW_OPCODE_F32TO16. Depending
1081 * on the HW gen, it is a special hw opcode or just a MOV, and
1082 * brw_F32TO16 (at brw_eu_emit) would do the work to chose.
1083 *
1084 * But if we want to use that opcode, we need to provide support on
1085 * different optimizations and lowerings. As right now HF support is
1086 * only for gen8+, it will be better to use directly the MOV, and use
1087 * BRW_OPCODE_F32TO16 when/if we work for HF support on gen7.
1088 */
1089 assert(type_sz(op[0].type) < 8); /* brw_nir_lower_conversions */
1090 inst = bld.MOV(result, op[0]);
1091 break;
1092 }
1093
1094 case nir_op_b2i8:
1095 case nir_op_b2i16:
1096 case nir_op_b2i32:
1097 case nir_op_b2i64:
1098 case nir_op_b2f16:
1099 case nir_op_b2f32:
1100 case nir_op_b2f64:
1101 if (try_emit_b2fi_of_inot(bld, result, instr))
1102 break;
1103 op[0].type = BRW_REGISTER_TYPE_D;
1104 op[0].negate = !op[0].negate;
1105 /* fallthrough */
1106 case nir_op_i2f64:
1107 case nir_op_i2i64:
1108 case nir_op_u2f64:
1109 case nir_op_u2u64:
1110 case nir_op_f2f64:
1111 case nir_op_f2i64:
1112 case nir_op_f2u64:
1113 case nir_op_i2i32:
1114 case nir_op_u2u32:
1115 case nir_op_f2i32:
1116 case nir_op_f2u32:
1117 case nir_op_i2f16:
1118 case nir_op_i2i16:
1119 case nir_op_u2f16:
1120 case nir_op_u2u16:
1121 case nir_op_f2i16:
1122 case nir_op_f2u16:
1123 case nir_op_i2i8:
1124 case nir_op_u2u8:
1125 case nir_op_f2i8:
1126 case nir_op_f2u8:
1127 if (result.type == BRW_REGISTER_TYPE_B ||
1128 result.type == BRW_REGISTER_TYPE_UB ||
1129 result.type == BRW_REGISTER_TYPE_HF)
1130 assert(type_sz(op[0].type) < 8); /* brw_nir_lower_conversions */
1131
1132 if (op[0].type == BRW_REGISTER_TYPE_B ||
1133 op[0].type == BRW_REGISTER_TYPE_UB ||
1134 op[0].type == BRW_REGISTER_TYPE_HF)
1135 assert(type_sz(result.type) < 8); /* brw_nir_lower_conversions */
1136
1137 inst = bld.MOV(result, op[0]);
1138 break;
1139
1140 case nir_op_fsat:
1141 inst = bld.MOV(result, op[0]);
1142 inst->saturate = true;
1143 break;
1144
1145 case nir_op_fneg:
1146 case nir_op_ineg:
1147 op[0].negate = true;
1148 inst = bld.MOV(result, op[0]);
1149 break;
1150
1151 case nir_op_fabs:
1152 case nir_op_iabs:
1153 op[0].negate = false;
1154 op[0].abs = true;
1155 inst = bld.MOV(result, op[0]);
1156 break;
1157
1158 case nir_op_f2f32:
1159 if (nir_has_any_rounding_mode_enabled(execution_mode)) {
1160 brw_rnd_mode rnd =
1161 brw_rnd_mode_from_execution_mode(execution_mode);
1162 bld.emit(SHADER_OPCODE_RND_MODE, bld.null_reg_ud(),
1163 brw_imm_d(rnd));
1164 }
1165
1166 if (op[0].type == BRW_REGISTER_TYPE_HF)
1167 assert(type_sz(result.type) < 8); /* brw_nir_lower_conversions */
1168
1169 inst = bld.MOV(result, op[0]);
1170 break;
1171
1172 case nir_op_fsign:
1173 emit_fsign(bld, instr, result, op, 0);
1174 break;
1175
1176 case nir_op_frcp:
1177 inst = bld.emit(SHADER_OPCODE_RCP, result, op[0]);
1178 break;
1179
1180 case nir_op_fexp2:
1181 inst = bld.emit(SHADER_OPCODE_EXP2, result, op[0]);
1182 break;
1183
1184 case nir_op_flog2:
1185 inst = bld.emit(SHADER_OPCODE_LOG2, result, op[0]);
1186 break;
1187
1188 case nir_op_fsin:
1189 inst = bld.emit(SHADER_OPCODE_SIN, result, op[0]);
1190 break;
1191
1192 case nir_op_fcos:
1193 inst = bld.emit(SHADER_OPCODE_COS, result, op[0]);
1194 break;
1195
1196 case nir_op_fddx:
1197 if (fs_key->high_quality_derivatives) {
1198 inst = bld.emit(FS_OPCODE_DDX_FINE, result, op[0]);
1199 } else {
1200 inst = bld.emit(FS_OPCODE_DDX_COARSE, result, op[0]);
1201 }
1202 break;
1203 case nir_op_fddx_fine:
1204 inst = bld.emit(FS_OPCODE_DDX_FINE, result, op[0]);
1205 break;
1206 case nir_op_fddx_coarse:
1207 inst = bld.emit(FS_OPCODE_DDX_COARSE, result, op[0]);
1208 break;
1209 case nir_op_fddy:
1210 if (fs_key->high_quality_derivatives) {
1211 inst = bld.emit(FS_OPCODE_DDY_FINE, result, op[0]);
1212 } else {
1213 inst = bld.emit(FS_OPCODE_DDY_COARSE, result, op[0]);
1214 }
1215 break;
1216 case nir_op_fddy_fine:
1217 inst = bld.emit(FS_OPCODE_DDY_FINE, result, op[0]);
1218 break;
1219 case nir_op_fddy_coarse:
1220 inst = bld.emit(FS_OPCODE_DDY_COARSE, result, op[0]);
1221 break;
1222
1223 case nir_op_fadd:
1224 if (nir_has_any_rounding_mode_enabled(execution_mode)) {
1225 brw_rnd_mode rnd =
1226 brw_rnd_mode_from_execution_mode(execution_mode);
1227 bld.emit(SHADER_OPCODE_RND_MODE, bld.null_reg_ud(),
1228 brw_imm_d(rnd));
1229 }
1230 /* fallthrough */
1231 case nir_op_iadd:
1232 inst = bld.ADD(result, op[0], op[1]);
1233 break;
1234
1235 case nir_op_iadd_sat:
1236 case nir_op_uadd_sat:
1237 inst = bld.ADD(result, op[0], op[1]);
1238 inst->saturate = true;
1239 break;
1240
1241 case nir_op_isub_sat:
1242 bld.emit(SHADER_OPCODE_ISUB_SAT, result, op[0], op[1]);
1243 break;
1244
1245 case nir_op_usub_sat:
1246 bld.emit(SHADER_OPCODE_USUB_SAT, result, op[0], op[1]);
1247 break;
1248
1249 case nir_op_irhadd:
1250 case nir_op_urhadd:
1251 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1252 inst = bld.AVG(result, op[0], op[1]);
1253 break;
1254
1255 case nir_op_ihadd:
1256 case nir_op_uhadd: {
1257 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1258 fs_reg tmp = bld.vgrf(result.type);
1259
1260 if (devinfo->gen >= 8) {
1261 op[0] = resolve_source_modifiers(op[0]);
1262 op[1] = resolve_source_modifiers(op[1]);
1263 }
1264
1265 /* AVG(x, y) - ((x ^ y) & 1) */
1266 bld.XOR(tmp, op[0], op[1]);
1267 bld.AND(tmp, tmp, retype(brw_imm_ud(1), result.type));
1268 bld.AVG(result, op[0], op[1]);
1269 inst = bld.ADD(result, result, tmp);
1270 inst->src[1].negate = true;
1271 break;
1272 }
1273
1274 case nir_op_fmul:
1275 for (unsigned i = 0; i < 2; i++) {
1276 if (can_fuse_fmul_fsign(instr, i)) {
1277 emit_fsign(bld, instr, result, op, i);
1278 return;
1279 }
1280 }
1281
1282 /* We emit the rounding mode after the previous fsign optimization since
1283 * it won't result in a MUL, but will try to negate the value by other
1284 * means.
1285 */
1286 if (nir_has_any_rounding_mode_enabled(execution_mode)) {
1287 brw_rnd_mode rnd =
1288 brw_rnd_mode_from_execution_mode(execution_mode);
1289 bld.emit(SHADER_OPCODE_RND_MODE, bld.null_reg_ud(),
1290 brw_imm_d(rnd));
1291 }
1292
1293 inst = bld.MUL(result, op[0], op[1]);
1294 break;
1295
1296 case nir_op_imul_2x32_64:
1297 case nir_op_umul_2x32_64:
1298 bld.MUL(result, op[0], op[1]);
1299 break;
1300
1301 case nir_op_imul_32x16:
1302 case nir_op_umul_32x16: {
1303 const bool ud = instr->op == nir_op_umul_32x16;
1304
1305 assert(nir_dest_bit_size(instr->dest.dest) == 32);
1306
1307 /* Before Gen7, the order of the 32-bit source and the 16-bit source was
1308 * swapped. The extension isn't enabled on those platforms, so don't
1309 * pretend to support the differences.
1310 */
1311 assert(devinfo->gen >= 7);
1312
1313 if (op[1].file == IMM)
1314 op[1] = ud ? brw_imm_uw(op[1].ud) : brw_imm_w(op[1].d);
1315 else {
1316 const enum brw_reg_type word_type =
1317 ud ? BRW_REGISTER_TYPE_UW : BRW_REGISTER_TYPE_W;
1318
1319 op[1] = subscript(op[1], word_type, 0);
1320 }
1321
1322 const enum brw_reg_type dword_type =
1323 ud ? BRW_REGISTER_TYPE_UD : BRW_REGISTER_TYPE_D;
1324
1325 bld.MUL(result, retype(op[0], dword_type), op[1]);
1326 break;
1327 }
1328
1329 case nir_op_imul:
1330 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1331 bld.MUL(result, op[0], op[1]);
1332 break;
1333
1334 case nir_op_imul_high:
1335 case nir_op_umul_high:
1336 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1337 bld.emit(SHADER_OPCODE_MULH, result, op[0], op[1]);
1338 break;
1339
1340 case nir_op_idiv:
1341 case nir_op_udiv:
1342 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1343 bld.emit(SHADER_OPCODE_INT_QUOTIENT, result, op[0], op[1]);
1344 break;
1345
1346 case nir_op_uadd_carry:
1347 unreachable("Should have been lowered by carry_to_arith().");
1348
1349 case nir_op_usub_borrow:
1350 unreachable("Should have been lowered by borrow_to_arith().");
1351
1352 case nir_op_umod:
1353 case nir_op_irem:
1354 /* According to the sign table for INT DIV in the Ivy Bridge PRM, it
1355 * appears that our hardware just does the right thing for signed
1356 * remainder.
1357 */
1358 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1359 bld.emit(SHADER_OPCODE_INT_REMAINDER, result, op[0], op[1]);
1360 break;
1361
1362 case nir_op_imod: {
1363 /* Get a regular C-style remainder. If a % b == 0, set the predicate. */
1364 bld.emit(SHADER_OPCODE_INT_REMAINDER, result, op[0], op[1]);
1365
1366 /* Math instructions don't support conditional mod */
1367 inst = bld.MOV(bld.null_reg_d(), result);
1368 inst->conditional_mod = BRW_CONDITIONAL_NZ;
1369
1370 /* Now, we need to determine if signs of the sources are different.
1371 * When we XOR the sources, the top bit is 0 if they are the same and 1
1372 * if they are different. We can then use a conditional modifier to
1373 * turn that into a predicate. This leads us to an XOR.l instruction.
1374 *
1375 * Technically, according to the PRM, you're not allowed to use .l on a
1376 * XOR instruction. However, emperical experiments and Curro's reading
1377 * of the simulator source both indicate that it's safe.
1378 */
1379 fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_D);
1380 inst = bld.XOR(tmp, op[0], op[1]);
1381 inst->predicate = BRW_PREDICATE_NORMAL;
1382 inst->conditional_mod = BRW_CONDITIONAL_L;
1383
1384 /* If the result of the initial remainder operation is non-zero and the
1385 * two sources have different signs, add in a copy of op[1] to get the
1386 * final integer modulus value.
1387 */
1388 inst = bld.ADD(result, result, op[1]);
1389 inst->predicate = BRW_PREDICATE_NORMAL;
1390 break;
1391 }
1392
1393 case nir_op_flt32:
1394 case nir_op_fge32:
1395 case nir_op_feq32:
1396 case nir_op_fneu32: {
1397 fs_reg dest = result;
1398
1399 const uint32_t bit_size = nir_src_bit_size(instr->src[0].src);
1400 if (bit_size != 32)
1401 dest = bld.vgrf(op[0].type, 1);
1402
1403 bld.CMP(dest, op[0], op[1], brw_cmod_for_nir_comparison(instr->op));
1404
1405 if (bit_size > 32) {
1406 bld.MOV(result, subscript(dest, BRW_REGISTER_TYPE_UD, 0));
1407 } else if(bit_size < 32) {
1408 /* When we convert the result to 32-bit we need to be careful and do
1409 * it as a signed conversion to get sign extension (for 32-bit true)
1410 */
1411 const brw_reg_type src_type =
1412 brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_D);
1413
1414 bld.MOV(retype(result, BRW_REGISTER_TYPE_D), retype(dest, src_type));
1415 }
1416 break;
1417 }
1418
1419 case nir_op_ilt32:
1420 case nir_op_ult32:
1421 case nir_op_ige32:
1422 case nir_op_uge32:
1423 case nir_op_ieq32:
1424 case nir_op_ine32: {
1425 fs_reg dest = result;
1426
1427 const uint32_t bit_size = type_sz(op[0].type) * 8;
1428 if (bit_size != 32)
1429 dest = bld.vgrf(op[0].type, 1);
1430
1431 bld.CMP(dest, op[0], op[1],
1432 brw_cmod_for_nir_comparison(instr->op));
1433
1434 if (bit_size > 32) {
1435 bld.MOV(result, subscript(dest, BRW_REGISTER_TYPE_UD, 0));
1436 } else if (bit_size < 32) {
1437 /* When we convert the result to 32-bit we need to be careful and do
1438 * it as a signed conversion to get sign extension (for 32-bit true)
1439 */
1440 const brw_reg_type src_type =
1441 brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_D);
1442
1443 bld.MOV(retype(result, BRW_REGISTER_TYPE_D), retype(dest, src_type));
1444 }
1445 break;
1446 }
1447
1448 case nir_op_inot:
1449 if (devinfo->gen >= 8) {
1450 nir_alu_instr *inot_src_instr = nir_src_as_alu_instr(instr->src[0].src);
1451
1452 if (inot_src_instr != NULL &&
1453 (inot_src_instr->op == nir_op_ior ||
1454 inot_src_instr->op == nir_op_ixor ||
1455 inot_src_instr->op == nir_op_iand)) {
1456 /* The sources of the source logical instruction are now the
1457 * sources of the instruction that will be generated.
1458 */
1459 prepare_alu_destination_and_sources(bld, inot_src_instr, op, false);
1460 resolve_inot_sources(bld, inot_src_instr, op);
1461
1462 /* Smash all of the sources and destination to be signed. This
1463 * doesn't matter for the operation of the instruction, but cmod
1464 * propagation fails on unsigned sources with negation (due to
1465 * fs_inst::can_do_cmod returning false).
1466 */
1467 result.type =
1468 brw_type_for_nir_type(devinfo,
1469 (nir_alu_type)(nir_type_int |
1470 nir_dest_bit_size(instr->dest.dest)));
1471 op[0].type =
1472 brw_type_for_nir_type(devinfo,
1473 (nir_alu_type)(nir_type_int |
1474 nir_src_bit_size(inot_src_instr->src[0].src)));
1475 op[1].type =
1476 brw_type_for_nir_type(devinfo,
1477 (nir_alu_type)(nir_type_int |
1478 nir_src_bit_size(inot_src_instr->src[1].src)));
1479
1480 /* For XOR, only invert one of the sources. Arbitrarily choose
1481 * the first source.
1482 */
1483 op[0].negate = !op[0].negate;
1484 if (inot_src_instr->op != nir_op_ixor)
1485 op[1].negate = !op[1].negate;
1486
1487 switch (inot_src_instr->op) {
1488 case nir_op_ior:
1489 bld.AND(result, op[0], op[1]);
1490 return;
1491
1492 case nir_op_iand:
1493 bld.OR(result, op[0], op[1]);
1494 return;
1495
1496 case nir_op_ixor:
1497 bld.XOR(result, op[0], op[1]);
1498 return;
1499
1500 default:
1501 unreachable("impossible opcode");
1502 }
1503 }
1504 op[0] = resolve_source_modifiers(op[0]);
1505 }
1506 bld.NOT(result, op[0]);
1507 break;
1508 case nir_op_ixor:
1509 if (devinfo->gen >= 8) {
1510 resolve_inot_sources(bld, instr, op);
1511 }
1512 bld.XOR(result, op[0], op[1]);
1513 break;
1514 case nir_op_ior:
1515 if (devinfo->gen >= 8) {
1516 resolve_inot_sources(bld, instr, op);
1517 }
1518 bld.OR(result, op[0], op[1]);
1519 break;
1520 case nir_op_iand:
1521 if (devinfo->gen >= 8) {
1522 resolve_inot_sources(bld, instr, op);
1523 }
1524 bld.AND(result, op[0], op[1]);
1525 break;
1526
1527 case nir_op_fdot2:
1528 case nir_op_fdot3:
1529 case nir_op_fdot4:
1530 case nir_op_b32all_fequal2:
1531 case nir_op_b32all_iequal2:
1532 case nir_op_b32all_fequal3:
1533 case nir_op_b32all_iequal3:
1534 case nir_op_b32all_fequal4:
1535 case nir_op_b32all_iequal4:
1536 case nir_op_b32any_fnequal2:
1537 case nir_op_b32any_inequal2:
1538 case nir_op_b32any_fnequal3:
1539 case nir_op_b32any_inequal3:
1540 case nir_op_b32any_fnequal4:
1541 case nir_op_b32any_inequal4:
1542 unreachable("Lowered by nir_lower_alu_reductions");
1543
1544 case nir_op_ldexp:
1545 unreachable("not reached: should be handled by ldexp_to_arith()");
1546
1547 case nir_op_fsqrt:
1548 inst = bld.emit(SHADER_OPCODE_SQRT, result, op[0]);
1549 break;
1550
1551 case nir_op_frsq:
1552 inst = bld.emit(SHADER_OPCODE_RSQ, result, op[0]);
1553 break;
1554
1555 case nir_op_i2b32:
1556 case nir_op_f2b32: {
1557 uint32_t bit_size = nir_src_bit_size(instr->src[0].src);
1558 if (bit_size == 64) {
1559 /* two-argument instructions can't take 64-bit immediates */
1560 fs_reg zero;
1561 fs_reg tmp;
1562
1563 if (instr->op == nir_op_f2b32) {
1564 zero = vgrf(glsl_type::double_type);
1565 tmp = vgrf(glsl_type::double_type);
1566 bld.MOV(zero, setup_imm_df(bld, 0.0));
1567 } else {
1568 zero = vgrf(glsl_type::int64_t_type);
1569 tmp = vgrf(glsl_type::int64_t_type);
1570 bld.MOV(zero, brw_imm_q(0));
1571 }
1572
1573 /* A SIMD16 execution needs to be split in two instructions, so use
1574 * a vgrf instead of the flag register as dst so instruction splitting
1575 * works
1576 */
1577 bld.CMP(tmp, op[0], zero, BRW_CONDITIONAL_NZ);
1578 bld.MOV(result, subscript(tmp, BRW_REGISTER_TYPE_UD, 0));
1579 } else {
1580 fs_reg zero;
1581 if (bit_size == 32) {
1582 zero = instr->op == nir_op_f2b32 ? brw_imm_f(0.0f) : brw_imm_d(0);
1583 } else {
1584 assert(bit_size == 16);
1585 zero = instr->op == nir_op_f2b32 ?
1586 retype(brw_imm_w(0), BRW_REGISTER_TYPE_HF) : brw_imm_w(0);
1587 }
1588 bld.CMP(result, op[0], zero, BRW_CONDITIONAL_NZ);
1589 }
1590 break;
1591 }
1592
1593 case nir_op_ftrunc:
1594 inst = bld.RNDZ(result, op[0]);
1595 if (devinfo->gen < 6) {
1596 set_condmod(BRW_CONDITIONAL_R, inst);
1597 set_predicate(BRW_PREDICATE_NORMAL,
1598 bld.ADD(result, result, brw_imm_f(1.0f)));
1599 inst = bld.MOV(result, result); /* for potential saturation */
1600 }
1601 break;
1602
1603 case nir_op_fceil: {
1604 op[0].negate = !op[0].negate;
1605 fs_reg temp = vgrf(glsl_type::float_type);
1606 bld.RNDD(temp, op[0]);
1607 temp.negate = true;
1608 inst = bld.MOV(result, temp);
1609 break;
1610 }
1611 case nir_op_ffloor:
1612 inst = bld.RNDD(result, op[0]);
1613 break;
1614 case nir_op_ffract:
1615 inst = bld.FRC(result, op[0]);
1616 break;
1617 case nir_op_fround_even:
1618 inst = bld.RNDE(result, op[0]);
1619 if (devinfo->gen < 6) {
1620 set_condmod(BRW_CONDITIONAL_R, inst);
1621 set_predicate(BRW_PREDICATE_NORMAL,
1622 bld.ADD(result, result, brw_imm_f(1.0f)));
1623 inst = bld.MOV(result, result); /* for potential saturation */
1624 }
1625 break;
1626
1627 case nir_op_fquantize2f16: {
1628 fs_reg tmp16 = bld.vgrf(BRW_REGISTER_TYPE_D);
1629 fs_reg tmp32 = bld.vgrf(BRW_REGISTER_TYPE_F);
1630 fs_reg zero = bld.vgrf(BRW_REGISTER_TYPE_F);
1631
1632 /* The destination stride must be at least as big as the source stride. */
1633 tmp16.type = BRW_REGISTER_TYPE_W;
1634 tmp16.stride = 2;
1635
1636 /* Check for denormal */
1637 fs_reg abs_src0 = op[0];
1638 abs_src0.abs = true;
1639 bld.CMP(bld.null_reg_f(), abs_src0, brw_imm_f(ldexpf(1.0, -14)),
1640 BRW_CONDITIONAL_L);
1641 /* Get the appropriately signed zero */
1642 bld.AND(retype(zero, BRW_REGISTER_TYPE_UD),
1643 retype(op[0], BRW_REGISTER_TYPE_UD),
1644 brw_imm_ud(0x80000000));
1645 /* Do the actual F32 -> F16 -> F32 conversion */
1646 bld.emit(BRW_OPCODE_F32TO16, tmp16, op[0]);
1647 bld.emit(BRW_OPCODE_F16TO32, tmp32, tmp16);
1648 /* Select that or zero based on normal status */
1649 inst = bld.SEL(result, zero, tmp32);
1650 inst->predicate = BRW_PREDICATE_NORMAL;
1651 break;
1652 }
1653
1654 case nir_op_imin:
1655 case nir_op_umin:
1656 case nir_op_fmin:
1657 inst = bld.emit_minmax(result, op[0], op[1], BRW_CONDITIONAL_L);
1658 break;
1659
1660 case nir_op_imax:
1661 case nir_op_umax:
1662 case nir_op_fmax:
1663 inst = bld.emit_minmax(result, op[0], op[1], BRW_CONDITIONAL_GE);
1664 break;
1665
1666 case nir_op_pack_snorm_2x16:
1667 case nir_op_pack_snorm_4x8:
1668 case nir_op_pack_unorm_2x16:
1669 case nir_op_pack_unorm_4x8:
1670 case nir_op_unpack_snorm_2x16:
1671 case nir_op_unpack_snorm_4x8:
1672 case nir_op_unpack_unorm_2x16:
1673 case nir_op_unpack_unorm_4x8:
1674 case nir_op_unpack_half_2x16:
1675 case nir_op_pack_half_2x16:
1676 unreachable("not reached: should be handled by lower_packing_builtins");
1677
1678 case nir_op_unpack_half_2x16_split_x_flush_to_zero:
1679 assert(FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16 & execution_mode);
1680 /* Fall-through */
1681 case nir_op_unpack_half_2x16_split_x:
1682 inst = bld.emit(BRW_OPCODE_F16TO32, result,
1683 subscript(op[0], BRW_REGISTER_TYPE_UW, 0));
1684 break;
1685
1686 case nir_op_unpack_half_2x16_split_y_flush_to_zero:
1687 assert(FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16 & execution_mode);
1688 /* Fall-through */
1689 case nir_op_unpack_half_2x16_split_y:
1690 inst = bld.emit(BRW_OPCODE_F16TO32, result,
1691 subscript(op[0], BRW_REGISTER_TYPE_UW, 1));
1692 break;
1693
1694 case nir_op_pack_64_2x32_split:
1695 case nir_op_pack_32_2x16_split:
1696 bld.emit(FS_OPCODE_PACK, result, op[0], op[1]);
1697 break;
1698
1699 case nir_op_unpack_64_2x32_split_x:
1700 case nir_op_unpack_64_2x32_split_y: {
1701 if (instr->op == nir_op_unpack_64_2x32_split_x)
1702 bld.MOV(result, subscript(op[0], BRW_REGISTER_TYPE_UD, 0));
1703 else
1704 bld.MOV(result, subscript(op[0], BRW_REGISTER_TYPE_UD, 1));
1705 break;
1706 }
1707
1708 case nir_op_unpack_32_2x16_split_x:
1709 case nir_op_unpack_32_2x16_split_y: {
1710 if (instr->op == nir_op_unpack_32_2x16_split_x)
1711 bld.MOV(result, subscript(op[0], BRW_REGISTER_TYPE_UW, 0));
1712 else
1713 bld.MOV(result, subscript(op[0], BRW_REGISTER_TYPE_UW, 1));
1714 break;
1715 }
1716
1717 case nir_op_fpow:
1718 inst = bld.emit(SHADER_OPCODE_POW, result, op[0], op[1]);
1719 break;
1720
1721 case nir_op_bitfield_reverse:
1722 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1723 bld.BFREV(result, op[0]);
1724 break;
1725
1726 case nir_op_bit_count:
1727 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1728 bld.CBIT(result, op[0]);
1729 break;
1730
1731 case nir_op_ufind_msb: {
1732 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1733 emit_find_msb_using_lzd(bld, result, op[0], false);
1734 break;
1735 }
1736
1737 case nir_op_uclz:
1738 assert(nir_dest_bit_size(instr->dest.dest) == 32);
1739 bld.LZD(retype(result, BRW_REGISTER_TYPE_UD), op[0]);
1740 break;
1741
1742 case nir_op_ifind_msb: {
1743 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1744
1745 if (devinfo->gen < 7) {
1746 emit_find_msb_using_lzd(bld, result, op[0], true);
1747 } else {
1748 bld.FBH(retype(result, BRW_REGISTER_TYPE_UD), op[0]);
1749
1750 /* FBH counts from the MSB side, while GLSL's findMSB() wants the
1751 * count from the LSB side. If FBH didn't return an error
1752 * (0xFFFFFFFF), then subtract the result from 31 to convert the MSB
1753 * count into an LSB count.
1754 */
1755 bld.CMP(bld.null_reg_d(), result, brw_imm_d(-1), BRW_CONDITIONAL_NZ);
1756
1757 inst = bld.ADD(result, result, brw_imm_d(31));
1758 inst->predicate = BRW_PREDICATE_NORMAL;
1759 inst->src[0].negate = true;
1760 }
1761 break;
1762 }
1763
1764 case nir_op_find_lsb:
1765 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1766
1767 if (devinfo->gen < 7) {
1768 fs_reg temp = vgrf(glsl_type::int_type);
1769
1770 /* (x & -x) generates a value that consists of only the LSB of x.
1771 * For all powers of 2, findMSB(y) == findLSB(y).
1772 */
1773 fs_reg src = retype(op[0], BRW_REGISTER_TYPE_D);
1774 fs_reg negated_src = src;
1775
1776 /* One must be negated, and the other must be non-negated. It
1777 * doesn't matter which is which.
1778 */
1779 negated_src.negate = true;
1780 src.negate = false;
1781
1782 bld.AND(temp, src, negated_src);
1783 emit_find_msb_using_lzd(bld, result, temp, false);
1784 } else {
1785 bld.FBL(result, op[0]);
1786 }
1787 break;
1788
1789 case nir_op_ubitfield_extract:
1790 case nir_op_ibitfield_extract:
1791 unreachable("should have been lowered");
1792 case nir_op_ubfe:
1793 case nir_op_ibfe:
1794 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1795 bld.BFE(result, op[2], op[1], op[0]);
1796 break;
1797 case nir_op_bfm:
1798 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1799 bld.BFI1(result, op[0], op[1]);
1800 break;
1801 case nir_op_bfi:
1802 assert(nir_dest_bit_size(instr->dest.dest) < 64);
1803 bld.BFI2(result, op[0], op[1], op[2]);
1804 break;
1805
1806 case nir_op_bitfield_insert:
1807 unreachable("not reached: should have been lowered");
1808
1809 case nir_op_ishl:
1810 bld.SHL(result, op[0], op[1]);
1811 break;
1812 case nir_op_ishr:
1813 bld.ASR(result, op[0], op[1]);
1814 break;
1815 case nir_op_ushr:
1816 bld.SHR(result, op[0], op[1]);
1817 break;
1818
1819 case nir_op_urol:
1820 bld.ROL(result, op[0], op[1]);
1821 break;
1822 case nir_op_uror:
1823 bld.ROR(result, op[0], op[1]);
1824 break;
1825
1826 case nir_op_pack_half_2x16_split:
1827 bld.emit(FS_OPCODE_PACK_HALF_2x16_SPLIT, result, op[0], op[1]);
1828 break;
1829
1830 case nir_op_ffma:
1831 if (nir_has_any_rounding_mode_enabled(execution_mode)) {
1832 brw_rnd_mode rnd =
1833 brw_rnd_mode_from_execution_mode(execution_mode);
1834 bld.emit(SHADER_OPCODE_RND_MODE, bld.null_reg_ud(),
1835 brw_imm_d(rnd));
1836 }
1837
1838 inst = bld.MAD(result, op[2], op[1], op[0]);
1839 break;
1840
1841 case nir_op_flrp:
1842 if (nir_has_any_rounding_mode_enabled(execution_mode)) {
1843 brw_rnd_mode rnd =
1844 brw_rnd_mode_from_execution_mode(execution_mode);
1845 bld.emit(SHADER_OPCODE_RND_MODE, bld.null_reg_ud(),
1846 brw_imm_d(rnd));
1847 }
1848
1849 inst = bld.LRP(result, op[0], op[1], op[2]);
1850 break;
1851
1852 case nir_op_b32csel:
1853 if (optimize_frontfacing_ternary(instr, result))
1854 return;
1855
1856 bld.CMP(bld.null_reg_d(), op[0], brw_imm_d(0), BRW_CONDITIONAL_NZ);
1857 inst = bld.SEL(result, op[1], op[2]);
1858 inst->predicate = BRW_PREDICATE_NORMAL;
1859 break;
1860
1861 case nir_op_extract_u8:
1862 case nir_op_extract_i8: {
1863 unsigned byte = nir_src_as_uint(instr->src[1].src);
1864
1865 /* The PRMs say:
1866 *
1867 * BDW+
1868 * There is no direct conversion from B/UB to Q/UQ or Q/UQ to B/UB.
1869 * Use two instructions and a word or DWord intermediate integer type.
1870 */
1871 if (nir_dest_bit_size(instr->dest.dest) == 64) {
1872 const brw_reg_type type = brw_int_type(1, instr->op == nir_op_extract_i8);
1873
1874 if (instr->op == nir_op_extract_i8) {
1875 /* If we need to sign extend, extract to a word first */
1876 fs_reg w_temp = bld.vgrf(BRW_REGISTER_TYPE_W);
1877 bld.MOV(w_temp, subscript(op[0], type, byte));
1878 bld.MOV(result, w_temp);
1879 } else if (byte & 1) {
1880 /* Extract the high byte from the word containing the desired byte
1881 * offset.
1882 */
1883 bld.SHR(result,
1884 subscript(op[0], BRW_REGISTER_TYPE_UW, byte / 2),
1885 brw_imm_uw(8));
1886 } else {
1887 /* Otherwise use an AND with 0xff and a word type */
1888 bld.AND(result,
1889 subscript(op[0], BRW_REGISTER_TYPE_UW, byte / 2),
1890 brw_imm_uw(0xff));
1891 }
1892 } else {
1893 const brw_reg_type type = brw_int_type(1, instr->op == nir_op_extract_i8);
1894 bld.MOV(result, subscript(op[0], type, byte));
1895 }
1896 break;
1897 }
1898
1899 case nir_op_extract_u16:
1900 case nir_op_extract_i16: {
1901 const brw_reg_type type = brw_int_type(2, instr->op == nir_op_extract_i16);
1902 unsigned word = nir_src_as_uint(instr->src[1].src);
1903 bld.MOV(result, subscript(op[0], type, word));
1904 break;
1905 }
1906
1907 default:
1908 unreachable("unhandled instruction");
1909 }
1910
1911 /* If we need to do a boolean resolve, replace the result with -(x & 1)
1912 * to sign extend the low bit to 0/~0
1913 */
1914 if (devinfo->gen <= 5 &&
1915 !result.is_null() &&
1916 (instr->instr.pass_flags & BRW_NIR_BOOLEAN_MASK) == BRW_NIR_BOOLEAN_NEEDS_RESOLVE) {
1917 fs_reg masked = vgrf(glsl_type::int_type);
1918 bld.AND(masked, result, brw_imm_d(1));
1919 masked.negate = true;
1920 bld.MOV(retype(result, BRW_REGISTER_TYPE_D), masked);
1921 }
1922 }
1923
1924 void
nir_emit_load_const(const fs_builder & bld,nir_load_const_instr * instr)1925 fs_visitor::nir_emit_load_const(const fs_builder &bld,
1926 nir_load_const_instr *instr)
1927 {
1928 const brw_reg_type reg_type =
1929 brw_reg_type_from_bit_size(instr->def.bit_size, BRW_REGISTER_TYPE_D);
1930 fs_reg reg = bld.vgrf(reg_type, instr->def.num_components);
1931
1932 switch (instr->def.bit_size) {
1933 case 8:
1934 for (unsigned i = 0; i < instr->def.num_components; i++)
1935 bld.MOV(offset(reg, bld, i), setup_imm_b(bld, instr->value[i].i8));
1936 break;
1937
1938 case 16:
1939 for (unsigned i = 0; i < instr->def.num_components; i++)
1940 bld.MOV(offset(reg, bld, i), brw_imm_w(instr->value[i].i16));
1941 break;
1942
1943 case 32:
1944 for (unsigned i = 0; i < instr->def.num_components; i++)
1945 bld.MOV(offset(reg, bld, i), brw_imm_d(instr->value[i].i32));
1946 break;
1947
1948 case 64:
1949 assert(devinfo->gen >= 7);
1950 if (devinfo->gen == 7) {
1951 /* We don't get 64-bit integer types until gen8 */
1952 for (unsigned i = 0; i < instr->def.num_components; i++) {
1953 bld.MOV(retype(offset(reg, bld, i), BRW_REGISTER_TYPE_DF),
1954 setup_imm_df(bld, instr->value[i].f64));
1955 }
1956 } else {
1957 for (unsigned i = 0; i < instr->def.num_components; i++)
1958 bld.MOV(offset(reg, bld, i), brw_imm_q(instr->value[i].i64));
1959 }
1960 break;
1961
1962 default:
1963 unreachable("Invalid bit size");
1964 }
1965
1966 nir_ssa_values[instr->def.index] = reg;
1967 }
1968
1969 fs_reg
get_nir_src(const nir_src & src)1970 fs_visitor::get_nir_src(const nir_src &src)
1971 {
1972 fs_reg reg;
1973 if (src.is_ssa) {
1974 if (src.ssa->parent_instr->type == nir_instr_type_ssa_undef) {
1975 const brw_reg_type reg_type =
1976 brw_reg_type_from_bit_size(src.ssa->bit_size, BRW_REGISTER_TYPE_D);
1977 reg = bld.vgrf(reg_type, src.ssa->num_components);
1978 } else {
1979 reg = nir_ssa_values[src.ssa->index];
1980 }
1981 } else {
1982 /* We don't handle indirects on locals */
1983 assert(src.reg.indirect == NULL);
1984 reg = offset(nir_locals[src.reg.reg->index], bld,
1985 src.reg.base_offset * src.reg.reg->num_components);
1986 }
1987
1988 if (nir_src_bit_size(src) == 64 && devinfo->gen == 7) {
1989 /* The only 64-bit type available on gen7 is DF, so use that. */
1990 reg.type = BRW_REGISTER_TYPE_DF;
1991 } else {
1992 /* To avoid floating-point denorm flushing problems, set the type by
1993 * default to an integer type - instructions that need floating point
1994 * semantics will set this to F if they need to
1995 */
1996 reg.type = brw_reg_type_from_bit_size(nir_src_bit_size(src),
1997 BRW_REGISTER_TYPE_D);
1998 }
1999
2000 return reg;
2001 }
2002
2003 /**
2004 * Return an IMM for constants; otherwise call get_nir_src() as normal.
2005 *
2006 * This function should not be called on any value which may be 64 bits.
2007 * We could theoretically support 64-bit on gen8+ but we choose not to
2008 * because it wouldn't work in general (no gen7 support) and there are
2009 * enough restrictions in 64-bit immediates that you can't take the return
2010 * value and treat it the same as the result of get_nir_src().
2011 */
2012 fs_reg
get_nir_src_imm(const nir_src & src)2013 fs_visitor::get_nir_src_imm(const nir_src &src)
2014 {
2015 assert(nir_src_bit_size(src) == 32);
2016 return nir_src_is_const(src) ?
2017 fs_reg(brw_imm_d(nir_src_as_int(src))) : get_nir_src(src);
2018 }
2019
2020 fs_reg
get_nir_dest(const nir_dest & dest)2021 fs_visitor::get_nir_dest(const nir_dest &dest)
2022 {
2023 if (dest.is_ssa) {
2024 const brw_reg_type reg_type =
2025 brw_reg_type_from_bit_size(dest.ssa.bit_size,
2026 dest.ssa.bit_size == 8 ?
2027 BRW_REGISTER_TYPE_D :
2028 BRW_REGISTER_TYPE_F);
2029 nir_ssa_values[dest.ssa.index] =
2030 bld.vgrf(reg_type, dest.ssa.num_components);
2031 bld.UNDEF(nir_ssa_values[dest.ssa.index]);
2032 return nir_ssa_values[dest.ssa.index];
2033 } else {
2034 /* We don't handle indirects on locals */
2035 assert(dest.reg.indirect == NULL);
2036 return offset(nir_locals[dest.reg.reg->index], bld,
2037 dest.reg.base_offset * dest.reg.reg->num_components);
2038 }
2039 }
2040
2041 void
emit_percomp(const fs_builder & bld,const fs_inst & inst,unsigned wr_mask)2042 fs_visitor::emit_percomp(const fs_builder &bld, const fs_inst &inst,
2043 unsigned wr_mask)
2044 {
2045 for (unsigned i = 0; i < 4; i++) {
2046 if (!((wr_mask >> i) & 1))
2047 continue;
2048
2049 fs_inst *new_inst = new(mem_ctx) fs_inst(inst);
2050 new_inst->dst = offset(new_inst->dst, bld, i);
2051 for (unsigned j = 0; j < new_inst->sources; j++)
2052 if (new_inst->src[j].file == VGRF)
2053 new_inst->src[j] = offset(new_inst->src[j], bld, i);
2054
2055 bld.emit(new_inst);
2056 }
2057 }
2058
2059 static fs_inst *
emit_pixel_interpolater_send(const fs_builder & bld,enum opcode opcode,const fs_reg & dst,const fs_reg & src,const fs_reg & desc,glsl_interp_mode interpolation)2060 emit_pixel_interpolater_send(const fs_builder &bld,
2061 enum opcode opcode,
2062 const fs_reg &dst,
2063 const fs_reg &src,
2064 const fs_reg &desc,
2065 glsl_interp_mode interpolation)
2066 {
2067 struct brw_wm_prog_data *wm_prog_data =
2068 brw_wm_prog_data(bld.shader->stage_prog_data);
2069
2070 fs_inst *inst = bld.emit(opcode, dst, src, desc);
2071 /* 2 floats per slot returned */
2072 inst->size_written = 2 * dst.component_size(inst->exec_size);
2073 inst->pi_noperspective = interpolation == INTERP_MODE_NOPERSPECTIVE;
2074
2075 wm_prog_data->pulls_bary = true;
2076
2077 return inst;
2078 }
2079
2080 /**
2081 * Computes 1 << x, given a D/UD register containing some value x.
2082 */
2083 static fs_reg
intexp2(const fs_builder & bld,const fs_reg & x)2084 intexp2(const fs_builder &bld, const fs_reg &x)
2085 {
2086 assert(x.type == BRW_REGISTER_TYPE_UD || x.type == BRW_REGISTER_TYPE_D);
2087
2088 fs_reg result = bld.vgrf(x.type, 1);
2089 fs_reg one = bld.vgrf(x.type, 1);
2090
2091 bld.MOV(one, retype(brw_imm_d(1), one.type));
2092 bld.SHL(result, one, x);
2093 return result;
2094 }
2095
2096 void
emit_gs_end_primitive(const nir_src & vertex_count_nir_src)2097 fs_visitor::emit_gs_end_primitive(const nir_src &vertex_count_nir_src)
2098 {
2099 assert(stage == MESA_SHADER_GEOMETRY);
2100
2101 struct brw_gs_prog_data *gs_prog_data = brw_gs_prog_data(prog_data);
2102
2103 if (gs_compile->control_data_header_size_bits == 0)
2104 return;
2105
2106 /* We can only do EndPrimitive() functionality when the control data
2107 * consists of cut bits. Fortunately, the only time it isn't is when the
2108 * output type is points, in which case EndPrimitive() is a no-op.
2109 */
2110 if (gs_prog_data->control_data_format !=
2111 GEN7_GS_CONTROL_DATA_FORMAT_GSCTL_CUT) {
2112 return;
2113 }
2114
2115 /* Cut bits use one bit per vertex. */
2116 assert(gs_compile->control_data_bits_per_vertex == 1);
2117
2118 fs_reg vertex_count = get_nir_src(vertex_count_nir_src);
2119 vertex_count.type = BRW_REGISTER_TYPE_UD;
2120
2121 /* Cut bit n should be set to 1 if EndPrimitive() was called after emitting
2122 * vertex n, 0 otherwise. So all we need to do here is mark bit
2123 * (vertex_count - 1) % 32 in the cut_bits register to indicate that
2124 * EndPrimitive() was called after emitting vertex (vertex_count - 1);
2125 * vec4_gs_visitor::emit_control_data_bits() will take care of the rest.
2126 *
2127 * Note that if EndPrimitive() is called before emitting any vertices, this
2128 * will cause us to set bit 31 of the control_data_bits register to 1.
2129 * That's fine because:
2130 *
2131 * - If max_vertices < 32, then vertex number 31 (zero-based) will never be
2132 * output, so the hardware will ignore cut bit 31.
2133 *
2134 * - If max_vertices == 32, then vertex number 31 is guaranteed to be the
2135 * last vertex, so setting cut bit 31 has no effect (since the primitive
2136 * is automatically ended when the GS terminates).
2137 *
2138 * - If max_vertices > 32, then the ir_emit_vertex visitor will reset the
2139 * control_data_bits register to 0 when the first vertex is emitted.
2140 */
2141
2142 const fs_builder abld = bld.annotate("end primitive");
2143
2144 /* control_data_bits |= 1 << ((vertex_count - 1) % 32) */
2145 fs_reg prev_count = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2146 abld.ADD(prev_count, vertex_count, brw_imm_ud(0xffffffffu));
2147 fs_reg mask = intexp2(abld, prev_count);
2148 /* Note: we're relying on the fact that the GEN SHL instruction only pays
2149 * attention to the lower 5 bits of its second source argument, so on this
2150 * architecture, 1 << (vertex_count - 1) is equivalent to 1 <<
2151 * ((vertex_count - 1) % 32).
2152 */
2153 abld.OR(this->control_data_bits, this->control_data_bits, mask);
2154 }
2155
2156 void
emit_gs_control_data_bits(const fs_reg & vertex_count)2157 fs_visitor::emit_gs_control_data_bits(const fs_reg &vertex_count)
2158 {
2159 assert(stage == MESA_SHADER_GEOMETRY);
2160 assert(gs_compile->control_data_bits_per_vertex != 0);
2161
2162 struct brw_gs_prog_data *gs_prog_data = brw_gs_prog_data(prog_data);
2163
2164 const fs_builder abld = bld.annotate("emit control data bits");
2165 const fs_builder fwa_bld = bld.exec_all();
2166
2167 /* We use a single UD register to accumulate control data bits (32 bits
2168 * for each of the SIMD8 channels). So we need to write a DWord (32 bits)
2169 * at a time.
2170 *
2171 * Unfortunately, the URB_WRITE_SIMD8 message uses 128-bit (OWord) offsets.
2172 * We have select a 128-bit group via the Global and Per-Slot Offsets, then
2173 * use the Channel Mask phase to enable/disable which DWord within that
2174 * group to write. (Remember, different SIMD8 channels may have emitted
2175 * different numbers of vertices, so we may need per-slot offsets.)
2176 *
2177 * Channel masking presents an annoying problem: we may have to replicate
2178 * the data up to 4 times:
2179 *
2180 * Msg = Handles, Per-Slot Offsets, Channel Masks, Data, Data, Data, Data.
2181 *
2182 * To avoid penalizing shaders that emit a small number of vertices, we
2183 * can avoid these sometimes: if the size of the control data header is
2184 * <= 128 bits, then there is only 1 OWord. All SIMD8 channels will land
2185 * land in the same 128-bit group, so we can skip per-slot offsets.
2186 *
2187 * Similarly, if the control data header is <= 32 bits, there is only one
2188 * DWord, so we can skip channel masks.
2189 */
2190 enum opcode opcode = SHADER_OPCODE_URB_WRITE_SIMD8;
2191
2192 fs_reg channel_mask, per_slot_offset;
2193
2194 if (gs_compile->control_data_header_size_bits > 32) {
2195 opcode = SHADER_OPCODE_URB_WRITE_SIMD8_MASKED;
2196 channel_mask = vgrf(glsl_type::uint_type);
2197 }
2198
2199 if (gs_compile->control_data_header_size_bits > 128) {
2200 opcode = SHADER_OPCODE_URB_WRITE_SIMD8_MASKED_PER_SLOT;
2201 per_slot_offset = vgrf(glsl_type::uint_type);
2202 }
2203
2204 /* Figure out which DWord we're trying to write to using the formula:
2205 *
2206 * dword_index = (vertex_count - 1) * bits_per_vertex / 32
2207 *
2208 * Since bits_per_vertex is a power of two, and is known at compile
2209 * time, this can be optimized to:
2210 *
2211 * dword_index = (vertex_count - 1) >> (6 - log2(bits_per_vertex))
2212 */
2213 if (opcode != SHADER_OPCODE_URB_WRITE_SIMD8) {
2214 fs_reg dword_index = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2215 fs_reg prev_count = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2216 abld.ADD(prev_count, vertex_count, brw_imm_ud(0xffffffffu));
2217 unsigned log2_bits_per_vertex =
2218 util_last_bit(gs_compile->control_data_bits_per_vertex);
2219 abld.SHR(dword_index, prev_count, brw_imm_ud(6u - log2_bits_per_vertex));
2220
2221 if (per_slot_offset.file != BAD_FILE) {
2222 /* Set the per-slot offset to dword_index / 4, so that we'll write to
2223 * the appropriate OWord within the control data header.
2224 */
2225 abld.SHR(per_slot_offset, dword_index, brw_imm_ud(2u));
2226 }
2227
2228 /* Set the channel masks to 1 << (dword_index % 4), so that we'll
2229 * write to the appropriate DWORD within the OWORD.
2230 */
2231 fs_reg channel = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2232 fwa_bld.AND(channel, dword_index, brw_imm_ud(3u));
2233 channel_mask = intexp2(fwa_bld, channel);
2234 /* Then the channel masks need to be in bits 23:16. */
2235 fwa_bld.SHL(channel_mask, channel_mask, brw_imm_ud(16u));
2236 }
2237
2238 /* Store the control data bits in the message payload and send it. */
2239 unsigned mlen = 2;
2240 if (channel_mask.file != BAD_FILE)
2241 mlen += 4; /* channel masks, plus 3 extra copies of the data */
2242 if (per_slot_offset.file != BAD_FILE)
2243 mlen++;
2244
2245 fs_reg payload = bld.vgrf(BRW_REGISTER_TYPE_UD, mlen);
2246 fs_reg *sources = ralloc_array(mem_ctx, fs_reg, mlen);
2247 unsigned i = 0;
2248 sources[i++] = fs_reg(retype(brw_vec8_grf(1, 0), BRW_REGISTER_TYPE_UD));
2249 if (per_slot_offset.file != BAD_FILE)
2250 sources[i++] = per_slot_offset;
2251 if (channel_mask.file != BAD_FILE)
2252 sources[i++] = channel_mask;
2253 while (i < mlen) {
2254 sources[i++] = this->control_data_bits;
2255 }
2256
2257 abld.LOAD_PAYLOAD(payload, sources, mlen, mlen);
2258 fs_inst *inst = abld.emit(opcode, reg_undef, payload);
2259 inst->mlen = mlen;
2260 /* We need to increment Global Offset by 256-bits to make room for
2261 * Broadwell's extra "Vertex Count" payload at the beginning of the
2262 * URB entry. Since this is an OWord message, Global Offset is counted
2263 * in 128-bit units, so we must set it to 2.
2264 */
2265 if (gs_prog_data->static_vertex_count == -1)
2266 inst->offset = 2;
2267 }
2268
2269 void
set_gs_stream_control_data_bits(const fs_reg & vertex_count,unsigned stream_id)2270 fs_visitor::set_gs_stream_control_data_bits(const fs_reg &vertex_count,
2271 unsigned stream_id)
2272 {
2273 /* control_data_bits |= stream_id << ((2 * (vertex_count - 1)) % 32) */
2274
2275 /* Note: we are calling this *before* increasing vertex_count, so
2276 * this->vertex_count == vertex_count - 1 in the formula above.
2277 */
2278
2279 /* Stream mode uses 2 bits per vertex */
2280 assert(gs_compile->control_data_bits_per_vertex == 2);
2281
2282 /* Must be a valid stream */
2283 assert(stream_id < MAX_VERTEX_STREAMS);
2284
2285 /* Control data bits are initialized to 0 so we don't have to set any
2286 * bits when sending vertices to stream 0.
2287 */
2288 if (stream_id == 0)
2289 return;
2290
2291 const fs_builder abld = bld.annotate("set stream control data bits", NULL);
2292
2293 /* reg::sid = stream_id */
2294 fs_reg sid = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2295 abld.MOV(sid, brw_imm_ud(stream_id));
2296
2297 /* reg:shift_count = 2 * (vertex_count - 1) */
2298 fs_reg shift_count = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2299 abld.SHL(shift_count, vertex_count, brw_imm_ud(1u));
2300
2301 /* Note: we're relying on the fact that the GEN SHL instruction only pays
2302 * attention to the lower 5 bits of its second source argument, so on this
2303 * architecture, stream_id << 2 * (vertex_count - 1) is equivalent to
2304 * stream_id << ((2 * (vertex_count - 1)) % 32).
2305 */
2306 fs_reg mask = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2307 abld.SHL(mask, sid, shift_count);
2308 abld.OR(this->control_data_bits, this->control_data_bits, mask);
2309 }
2310
2311 void
emit_gs_vertex(const nir_src & vertex_count_nir_src,unsigned stream_id)2312 fs_visitor::emit_gs_vertex(const nir_src &vertex_count_nir_src,
2313 unsigned stream_id)
2314 {
2315 assert(stage == MESA_SHADER_GEOMETRY);
2316
2317 struct brw_gs_prog_data *gs_prog_data = brw_gs_prog_data(prog_data);
2318
2319 fs_reg vertex_count = get_nir_src(vertex_count_nir_src);
2320 vertex_count.type = BRW_REGISTER_TYPE_UD;
2321
2322 /* Haswell and later hardware ignores the "Render Stream Select" bits
2323 * from the 3DSTATE_STREAMOUT packet when the SOL stage is disabled,
2324 * and instead sends all primitives down the pipeline for rasterization.
2325 * If the SOL stage is enabled, "Render Stream Select" is honored and
2326 * primitives bound to non-zero streams are discarded after stream output.
2327 *
2328 * Since the only purpose of primives sent to non-zero streams is to
2329 * be recorded by transform feedback, we can simply discard all geometry
2330 * bound to these streams when transform feedback is disabled.
2331 */
2332 if (stream_id > 0 && !nir->info.has_transform_feedback_varyings)
2333 return;
2334
2335 /* If we're outputting 32 control data bits or less, then we can wait
2336 * until the shader is over to output them all. Otherwise we need to
2337 * output them as we go. Now is the time to do it, since we're about to
2338 * output the vertex_count'th vertex, so it's guaranteed that the
2339 * control data bits associated with the (vertex_count - 1)th vertex are
2340 * correct.
2341 */
2342 if (gs_compile->control_data_header_size_bits > 32) {
2343 const fs_builder abld =
2344 bld.annotate("emit vertex: emit control data bits");
2345
2346 /* Only emit control data bits if we've finished accumulating a batch
2347 * of 32 bits. This is the case when:
2348 *
2349 * (vertex_count * bits_per_vertex) % 32 == 0
2350 *
2351 * (in other words, when the last 5 bits of vertex_count *
2352 * bits_per_vertex are 0). Assuming bits_per_vertex == 2^n for some
2353 * integer n (which is always the case, since bits_per_vertex is
2354 * always 1 or 2), this is equivalent to requiring that the last 5-n
2355 * bits of vertex_count are 0:
2356 *
2357 * vertex_count & (2^(5-n) - 1) == 0
2358 *
2359 * 2^(5-n) == 2^5 / 2^n == 32 / bits_per_vertex, so this is
2360 * equivalent to:
2361 *
2362 * vertex_count & (32 / bits_per_vertex - 1) == 0
2363 *
2364 * TODO: If vertex_count is an immediate, we could do some of this math
2365 * at compile time...
2366 */
2367 fs_inst *inst =
2368 abld.AND(bld.null_reg_d(), vertex_count,
2369 brw_imm_ud(32u / gs_compile->control_data_bits_per_vertex - 1u));
2370 inst->conditional_mod = BRW_CONDITIONAL_Z;
2371
2372 abld.IF(BRW_PREDICATE_NORMAL);
2373 /* If vertex_count is 0, then no control data bits have been
2374 * accumulated yet, so we can skip emitting them.
2375 */
2376 abld.CMP(bld.null_reg_d(), vertex_count, brw_imm_ud(0u),
2377 BRW_CONDITIONAL_NEQ);
2378 abld.IF(BRW_PREDICATE_NORMAL);
2379 emit_gs_control_data_bits(vertex_count);
2380 abld.emit(BRW_OPCODE_ENDIF);
2381
2382 /* Reset control_data_bits to 0 so we can start accumulating a new
2383 * batch.
2384 *
2385 * Note: in the case where vertex_count == 0, this neutralizes the
2386 * effect of any call to EndPrimitive() that the shader may have
2387 * made before outputting its first vertex.
2388 */
2389 inst = abld.MOV(this->control_data_bits, brw_imm_ud(0u));
2390 inst->force_writemask_all = true;
2391 abld.emit(BRW_OPCODE_ENDIF);
2392 }
2393
2394 emit_urb_writes(vertex_count);
2395
2396 /* In stream mode we have to set control data bits for all vertices
2397 * unless we have disabled control data bits completely (which we do
2398 * do for GL_POINTS outputs that don't use streams).
2399 */
2400 if (gs_compile->control_data_header_size_bits > 0 &&
2401 gs_prog_data->control_data_format ==
2402 GEN7_GS_CONTROL_DATA_FORMAT_GSCTL_SID) {
2403 set_gs_stream_control_data_bits(vertex_count, stream_id);
2404 }
2405 }
2406
2407 void
emit_gs_input_load(const fs_reg & dst,const nir_src & vertex_src,unsigned base_offset,const nir_src & offset_src,unsigned num_components,unsigned first_component)2408 fs_visitor::emit_gs_input_load(const fs_reg &dst,
2409 const nir_src &vertex_src,
2410 unsigned base_offset,
2411 const nir_src &offset_src,
2412 unsigned num_components,
2413 unsigned first_component)
2414 {
2415 assert(type_sz(dst.type) == 4);
2416 struct brw_gs_prog_data *gs_prog_data = brw_gs_prog_data(prog_data);
2417 const unsigned push_reg_count = gs_prog_data->base.urb_read_length * 8;
2418
2419 /* TODO: figure out push input layout for invocations == 1 */
2420 if (gs_prog_data->invocations == 1 &&
2421 nir_src_is_const(offset_src) && nir_src_is_const(vertex_src) &&
2422 4 * (base_offset + nir_src_as_uint(offset_src)) < push_reg_count) {
2423 int imm_offset = (base_offset + nir_src_as_uint(offset_src)) * 4 +
2424 nir_src_as_uint(vertex_src) * push_reg_count;
2425 for (unsigned i = 0; i < num_components; i++) {
2426 bld.MOV(offset(dst, bld, i),
2427 fs_reg(ATTR, imm_offset + i + first_component, dst.type));
2428 }
2429 return;
2430 }
2431
2432 /* Resort to the pull model. Ensure the VUE handles are provided. */
2433 assert(gs_prog_data->base.include_vue_handles);
2434
2435 unsigned first_icp_handle = gs_prog_data->include_primitive_id ? 3 : 2;
2436 fs_reg icp_handle = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2437
2438 if (gs_prog_data->invocations == 1) {
2439 if (nir_src_is_const(vertex_src)) {
2440 /* The vertex index is constant; just select the proper URB handle. */
2441 icp_handle =
2442 retype(brw_vec8_grf(first_icp_handle + nir_src_as_uint(vertex_src), 0),
2443 BRW_REGISTER_TYPE_UD);
2444 } else {
2445 /* The vertex index is non-constant. We need to use indirect
2446 * addressing to fetch the proper URB handle.
2447 *
2448 * First, we start with the sequence <7, 6, 5, 4, 3, 2, 1, 0>
2449 * indicating that channel <n> should read the handle from
2450 * DWord <n>. We convert that to bytes by multiplying by 4.
2451 *
2452 * Next, we convert the vertex index to bytes by multiplying
2453 * by 32 (shifting by 5), and add the two together. This is
2454 * the final indirect byte offset.
2455 */
2456 fs_reg sequence = bld.vgrf(BRW_REGISTER_TYPE_UW, 1);
2457 fs_reg channel_offsets = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2458 fs_reg vertex_offset_bytes = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2459 fs_reg icp_offset_bytes = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2460
2461 /* sequence = <7, 6, 5, 4, 3, 2, 1, 0> */
2462 bld.MOV(sequence, fs_reg(brw_imm_v(0x76543210)));
2463 /* channel_offsets = 4 * sequence = <28, 24, 20, 16, 12, 8, 4, 0> */
2464 bld.SHL(channel_offsets, sequence, brw_imm_ud(2u));
2465 /* Convert vertex_index to bytes (multiply by 32) */
2466 bld.SHL(vertex_offset_bytes,
2467 retype(get_nir_src(vertex_src), BRW_REGISTER_TYPE_UD),
2468 brw_imm_ud(5u));
2469 bld.ADD(icp_offset_bytes, vertex_offset_bytes, channel_offsets);
2470
2471 /* Use first_icp_handle as the base offset. There is one register
2472 * of URB handles per vertex, so inform the register allocator that
2473 * we might read up to nir->info.gs.vertices_in registers.
2474 */
2475 bld.emit(SHADER_OPCODE_MOV_INDIRECT, icp_handle,
2476 retype(brw_vec8_grf(first_icp_handle, 0), icp_handle.type),
2477 fs_reg(icp_offset_bytes),
2478 brw_imm_ud(nir->info.gs.vertices_in * REG_SIZE));
2479 }
2480 } else {
2481 assert(gs_prog_data->invocations > 1);
2482
2483 if (nir_src_is_const(vertex_src)) {
2484 unsigned vertex = nir_src_as_uint(vertex_src);
2485 assert(devinfo->gen >= 9 || vertex <= 5);
2486 bld.MOV(icp_handle,
2487 retype(brw_vec1_grf(first_icp_handle + vertex / 8, vertex % 8),
2488 BRW_REGISTER_TYPE_UD));
2489 } else {
2490 /* The vertex index is non-constant. We need to use indirect
2491 * addressing to fetch the proper URB handle.
2492 *
2493 */
2494 fs_reg icp_offset_bytes = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2495
2496 /* Convert vertex_index to bytes (multiply by 4) */
2497 bld.SHL(icp_offset_bytes,
2498 retype(get_nir_src(vertex_src), BRW_REGISTER_TYPE_UD),
2499 brw_imm_ud(2u));
2500
2501 /* Use first_icp_handle as the base offset. There is one DWord
2502 * of URB handles per vertex, so inform the register allocator that
2503 * we might read up to ceil(nir->info.gs.vertices_in / 8) registers.
2504 */
2505 bld.emit(SHADER_OPCODE_MOV_INDIRECT, icp_handle,
2506 retype(brw_vec8_grf(first_icp_handle, 0), icp_handle.type),
2507 fs_reg(icp_offset_bytes),
2508 brw_imm_ud(DIV_ROUND_UP(nir->info.gs.vertices_in, 8) *
2509 REG_SIZE));
2510 }
2511 }
2512
2513 fs_inst *inst;
2514 fs_reg indirect_offset = get_nir_src(offset_src);
2515
2516 if (nir_src_is_const(offset_src)) {
2517 /* Constant indexing - use global offset. */
2518 if (first_component != 0) {
2519 unsigned read_components = num_components + first_component;
2520 fs_reg tmp = bld.vgrf(dst.type, read_components);
2521 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, tmp, icp_handle);
2522 inst->size_written = read_components *
2523 tmp.component_size(inst->exec_size);
2524 for (unsigned i = 0; i < num_components; i++) {
2525 bld.MOV(offset(dst, bld, i),
2526 offset(tmp, bld, i + first_component));
2527 }
2528 } else {
2529 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, dst, icp_handle);
2530 inst->size_written = num_components *
2531 dst.component_size(inst->exec_size);
2532 }
2533 inst->offset = base_offset + nir_src_as_uint(offset_src);
2534 inst->mlen = 1;
2535 } else {
2536 /* Indirect indexing - use per-slot offsets as well. */
2537 const fs_reg srcs[] = { icp_handle, indirect_offset };
2538 unsigned read_components = num_components + first_component;
2539 fs_reg tmp = bld.vgrf(dst.type, read_components);
2540 fs_reg payload = bld.vgrf(BRW_REGISTER_TYPE_UD, 2);
2541 bld.LOAD_PAYLOAD(payload, srcs, ARRAY_SIZE(srcs), 0);
2542 if (first_component != 0) {
2543 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, tmp,
2544 payload);
2545 inst->size_written = read_components *
2546 tmp.component_size(inst->exec_size);
2547 for (unsigned i = 0; i < num_components; i++) {
2548 bld.MOV(offset(dst, bld, i),
2549 offset(tmp, bld, i + first_component));
2550 }
2551 } else {
2552 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, dst, payload);
2553 inst->size_written = num_components *
2554 dst.component_size(inst->exec_size);
2555 }
2556 inst->offset = base_offset;
2557 inst->mlen = 2;
2558 }
2559 }
2560
2561 fs_reg
get_indirect_offset(nir_intrinsic_instr * instr)2562 fs_visitor::get_indirect_offset(nir_intrinsic_instr *instr)
2563 {
2564 nir_src *offset_src = nir_get_io_offset_src(instr);
2565
2566 if (nir_src_is_const(*offset_src)) {
2567 /* The only constant offset we should find is 0. brw_nir.c's
2568 * add_const_offset_to_base() will fold other constant offsets
2569 * into instr->const_index[0].
2570 */
2571 assert(nir_src_as_uint(*offset_src) == 0);
2572 return fs_reg();
2573 }
2574
2575 return get_nir_src(*offset_src);
2576 }
2577
2578 void
nir_emit_vs_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)2579 fs_visitor::nir_emit_vs_intrinsic(const fs_builder &bld,
2580 nir_intrinsic_instr *instr)
2581 {
2582 assert(stage == MESA_SHADER_VERTEX);
2583
2584 fs_reg dest;
2585 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
2586 dest = get_nir_dest(instr->dest);
2587
2588 switch (instr->intrinsic) {
2589 case nir_intrinsic_load_vertex_id:
2590 case nir_intrinsic_load_base_vertex:
2591 unreachable("should be lowered by nir_lower_system_values()");
2592
2593 case nir_intrinsic_load_input: {
2594 assert(nir_dest_bit_size(instr->dest) == 32);
2595 fs_reg src = fs_reg(ATTR, nir_intrinsic_base(instr) * 4, dest.type);
2596 src = offset(src, bld, nir_intrinsic_component(instr));
2597 src = offset(src, bld, nir_src_as_uint(instr->src[0]));
2598
2599 for (unsigned i = 0; i < instr->num_components; i++)
2600 bld.MOV(offset(dest, bld, i), offset(src, bld, i));
2601 break;
2602 }
2603
2604 case nir_intrinsic_load_vertex_id_zero_base:
2605 case nir_intrinsic_load_instance_id:
2606 case nir_intrinsic_load_base_instance:
2607 case nir_intrinsic_load_draw_id:
2608 case nir_intrinsic_load_first_vertex:
2609 case nir_intrinsic_load_is_indexed_draw:
2610 unreachable("lowered by brw_nir_lower_vs_inputs");
2611
2612 default:
2613 nir_emit_intrinsic(bld, instr);
2614 break;
2615 }
2616 }
2617
2618 fs_reg
get_tcs_single_patch_icp_handle(const fs_builder & bld,nir_intrinsic_instr * instr)2619 fs_visitor::get_tcs_single_patch_icp_handle(const fs_builder &bld,
2620 nir_intrinsic_instr *instr)
2621 {
2622 struct brw_tcs_prog_data *tcs_prog_data = brw_tcs_prog_data(prog_data);
2623 const nir_src &vertex_src = instr->src[0];
2624 nir_intrinsic_instr *vertex_intrin = nir_src_as_intrinsic(vertex_src);
2625 fs_reg icp_handle;
2626
2627 if (nir_src_is_const(vertex_src)) {
2628 /* Emit a MOV to resolve <0,1,0> regioning. */
2629 icp_handle = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2630 unsigned vertex = nir_src_as_uint(vertex_src);
2631 bld.MOV(icp_handle,
2632 retype(brw_vec1_grf(1 + (vertex >> 3), vertex & 7),
2633 BRW_REGISTER_TYPE_UD));
2634 } else if (tcs_prog_data->instances == 1 && vertex_intrin &&
2635 vertex_intrin->intrinsic == nir_intrinsic_load_invocation_id) {
2636 /* For the common case of only 1 instance, an array index of
2637 * gl_InvocationID means reading g1. Skip all the indirect work.
2638 */
2639 icp_handle = retype(brw_vec8_grf(1, 0), BRW_REGISTER_TYPE_UD);
2640 } else {
2641 /* The vertex index is non-constant. We need to use indirect
2642 * addressing to fetch the proper URB handle.
2643 */
2644 icp_handle = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2645
2646 /* Each ICP handle is a single DWord (4 bytes) */
2647 fs_reg vertex_offset_bytes = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2648 bld.SHL(vertex_offset_bytes,
2649 retype(get_nir_src(vertex_src), BRW_REGISTER_TYPE_UD),
2650 brw_imm_ud(2u));
2651
2652 /* Start at g1. We might read up to 4 registers. */
2653 bld.emit(SHADER_OPCODE_MOV_INDIRECT, icp_handle,
2654 retype(brw_vec8_grf(1, 0), icp_handle.type), vertex_offset_bytes,
2655 brw_imm_ud(4 * REG_SIZE));
2656 }
2657
2658 return icp_handle;
2659 }
2660
2661 fs_reg
get_tcs_eight_patch_icp_handle(const fs_builder & bld,nir_intrinsic_instr * instr)2662 fs_visitor::get_tcs_eight_patch_icp_handle(const fs_builder &bld,
2663 nir_intrinsic_instr *instr)
2664 {
2665 struct brw_tcs_prog_key *tcs_key = (struct brw_tcs_prog_key *) key;
2666 struct brw_tcs_prog_data *tcs_prog_data = brw_tcs_prog_data(prog_data);
2667 const nir_src &vertex_src = instr->src[0];
2668
2669 unsigned first_icp_handle = tcs_prog_data->include_primitive_id ? 3 : 2;
2670
2671 if (nir_src_is_const(vertex_src)) {
2672 return fs_reg(retype(brw_vec8_grf(first_icp_handle +
2673 nir_src_as_uint(vertex_src), 0),
2674 BRW_REGISTER_TYPE_UD));
2675 }
2676
2677 /* The vertex index is non-constant. We need to use indirect
2678 * addressing to fetch the proper URB handle.
2679 *
2680 * First, we start with the sequence <7, 6, 5, 4, 3, 2, 1, 0>
2681 * indicating that channel <n> should read the handle from
2682 * DWord <n>. We convert that to bytes by multiplying by 4.
2683 *
2684 * Next, we convert the vertex index to bytes by multiplying
2685 * by 32 (shifting by 5), and add the two together. This is
2686 * the final indirect byte offset.
2687 */
2688 fs_reg icp_handle = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2689 fs_reg sequence = bld.vgrf(BRW_REGISTER_TYPE_UW, 1);
2690 fs_reg channel_offsets = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2691 fs_reg vertex_offset_bytes = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2692 fs_reg icp_offset_bytes = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2693
2694 /* sequence = <7, 6, 5, 4, 3, 2, 1, 0> */
2695 bld.MOV(sequence, fs_reg(brw_imm_v(0x76543210)));
2696 /* channel_offsets = 4 * sequence = <28, 24, 20, 16, 12, 8, 4, 0> */
2697 bld.SHL(channel_offsets, sequence, brw_imm_ud(2u));
2698 /* Convert vertex_index to bytes (multiply by 32) */
2699 bld.SHL(vertex_offset_bytes,
2700 retype(get_nir_src(vertex_src), BRW_REGISTER_TYPE_UD),
2701 brw_imm_ud(5u));
2702 bld.ADD(icp_offset_bytes, vertex_offset_bytes, channel_offsets);
2703
2704 /* Use first_icp_handle as the base offset. There is one register
2705 * of URB handles per vertex, so inform the register allocator that
2706 * we might read up to nir->info.gs.vertices_in registers.
2707 */
2708 bld.emit(SHADER_OPCODE_MOV_INDIRECT, icp_handle,
2709 retype(brw_vec8_grf(first_icp_handle, 0), icp_handle.type),
2710 icp_offset_bytes, brw_imm_ud(tcs_key->input_vertices * REG_SIZE));
2711
2712 return icp_handle;
2713 }
2714
2715 struct brw_reg
get_tcs_output_urb_handle()2716 fs_visitor::get_tcs_output_urb_handle()
2717 {
2718 struct brw_vue_prog_data *vue_prog_data = brw_vue_prog_data(prog_data);
2719
2720 if (vue_prog_data->dispatch_mode == DISPATCH_MODE_TCS_SINGLE_PATCH) {
2721 return retype(brw_vec1_grf(0, 0), BRW_REGISTER_TYPE_UD);
2722 } else {
2723 assert(vue_prog_data->dispatch_mode == DISPATCH_MODE_TCS_8_PATCH);
2724 return retype(brw_vec8_grf(1, 0), BRW_REGISTER_TYPE_UD);
2725 }
2726 }
2727
2728 void
nir_emit_tcs_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)2729 fs_visitor::nir_emit_tcs_intrinsic(const fs_builder &bld,
2730 nir_intrinsic_instr *instr)
2731 {
2732 assert(stage == MESA_SHADER_TESS_CTRL);
2733 struct brw_tcs_prog_key *tcs_key = (struct brw_tcs_prog_key *) key;
2734 struct brw_tcs_prog_data *tcs_prog_data = brw_tcs_prog_data(prog_data);
2735 struct brw_vue_prog_data *vue_prog_data = &tcs_prog_data->base;
2736
2737 bool eight_patch =
2738 vue_prog_data->dispatch_mode == DISPATCH_MODE_TCS_8_PATCH;
2739
2740 fs_reg dst;
2741 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
2742 dst = get_nir_dest(instr->dest);
2743
2744 switch (instr->intrinsic) {
2745 case nir_intrinsic_load_primitive_id:
2746 bld.MOV(dst, fs_reg(eight_patch ? brw_vec8_grf(2, 0)
2747 : brw_vec1_grf(0, 1)));
2748 break;
2749 case nir_intrinsic_load_invocation_id:
2750 bld.MOV(retype(dst, invocation_id.type), invocation_id);
2751 break;
2752 case nir_intrinsic_load_patch_vertices_in:
2753 bld.MOV(retype(dst, BRW_REGISTER_TYPE_D),
2754 brw_imm_d(tcs_key->input_vertices));
2755 break;
2756
2757 case nir_intrinsic_control_barrier: {
2758 if (tcs_prog_data->instances == 1)
2759 break;
2760
2761 fs_reg m0 = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2762 fs_reg m0_2 = component(m0, 2);
2763
2764 const fs_builder chanbld = bld.exec_all().group(1, 0);
2765
2766 /* Zero the message header */
2767 bld.exec_all().MOV(m0, brw_imm_ud(0u));
2768
2769 if (devinfo->gen < 11) {
2770 /* Copy "Barrier ID" from r0.2, bits 16:13 */
2771 chanbld.AND(m0_2, retype(brw_vec1_grf(0, 2), BRW_REGISTER_TYPE_UD),
2772 brw_imm_ud(INTEL_MASK(16, 13)));
2773
2774 /* Shift it up to bits 27:24. */
2775 chanbld.SHL(m0_2, m0_2, brw_imm_ud(11));
2776 } else {
2777 chanbld.AND(m0_2, retype(brw_vec1_grf(0, 2), BRW_REGISTER_TYPE_UD),
2778 brw_imm_ud(INTEL_MASK(30, 24)));
2779 }
2780
2781 /* Set the Barrier Count and the enable bit */
2782 if (devinfo->gen < 11) {
2783 chanbld.OR(m0_2, m0_2,
2784 brw_imm_ud(tcs_prog_data->instances << 9 | (1 << 15)));
2785 } else {
2786 chanbld.OR(m0_2, m0_2,
2787 brw_imm_ud(tcs_prog_data->instances << 8 | (1 << 15)));
2788 }
2789
2790 bld.emit(SHADER_OPCODE_BARRIER, bld.null_reg_ud(), m0);
2791 break;
2792 }
2793
2794 case nir_intrinsic_load_input:
2795 unreachable("nir_lower_io should never give us these.");
2796 break;
2797
2798 case nir_intrinsic_load_per_vertex_input: {
2799 assert(nir_dest_bit_size(instr->dest) == 32);
2800 fs_reg indirect_offset = get_indirect_offset(instr);
2801 unsigned imm_offset = instr->const_index[0];
2802 fs_inst *inst;
2803
2804 fs_reg icp_handle =
2805 eight_patch ? get_tcs_eight_patch_icp_handle(bld, instr)
2806 : get_tcs_single_patch_icp_handle(bld, instr);
2807
2808 /* We can only read two double components with each URB read, so
2809 * we send two read messages in that case, each one loading up to
2810 * two double components.
2811 */
2812 unsigned num_components = instr->num_components;
2813 unsigned first_component = nir_intrinsic_component(instr);
2814
2815 if (indirect_offset.file == BAD_FILE) {
2816 /* Constant indexing - use global offset. */
2817 if (first_component != 0) {
2818 unsigned read_components = num_components + first_component;
2819 fs_reg tmp = bld.vgrf(dst.type, read_components);
2820 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, tmp, icp_handle);
2821 for (unsigned i = 0; i < num_components; i++) {
2822 bld.MOV(offset(dst, bld, i),
2823 offset(tmp, bld, i + first_component));
2824 }
2825 } else {
2826 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, dst, icp_handle);
2827 }
2828 inst->offset = imm_offset;
2829 inst->mlen = 1;
2830 } else {
2831 /* Indirect indexing - use per-slot offsets as well. */
2832 const fs_reg srcs[] = { icp_handle, indirect_offset };
2833 fs_reg payload = bld.vgrf(BRW_REGISTER_TYPE_UD, 2);
2834 bld.LOAD_PAYLOAD(payload, srcs, ARRAY_SIZE(srcs), 0);
2835 if (first_component != 0) {
2836 unsigned read_components = num_components + first_component;
2837 fs_reg tmp = bld.vgrf(dst.type, read_components);
2838 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, tmp,
2839 payload);
2840 for (unsigned i = 0; i < num_components; i++) {
2841 bld.MOV(offset(dst, bld, i),
2842 offset(tmp, bld, i + first_component));
2843 }
2844 } else {
2845 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, dst,
2846 payload);
2847 }
2848 inst->offset = imm_offset;
2849 inst->mlen = 2;
2850 }
2851 inst->size_written = (num_components + first_component) *
2852 inst->dst.component_size(inst->exec_size);
2853
2854 /* Copy the temporary to the destination to deal with writemasking.
2855 *
2856 * Also attempt to deal with gl_PointSize being in the .w component.
2857 */
2858 if (inst->offset == 0 && indirect_offset.file == BAD_FILE) {
2859 assert(type_sz(dst.type) == 4);
2860 inst->dst = bld.vgrf(dst.type, 4);
2861 inst->size_written = 4 * REG_SIZE;
2862 bld.MOV(dst, offset(inst->dst, bld, 3));
2863 }
2864 break;
2865 }
2866
2867 case nir_intrinsic_load_output:
2868 case nir_intrinsic_load_per_vertex_output: {
2869 assert(nir_dest_bit_size(instr->dest) == 32);
2870 fs_reg indirect_offset = get_indirect_offset(instr);
2871 unsigned imm_offset = instr->const_index[0];
2872 unsigned first_component = nir_intrinsic_component(instr);
2873
2874 struct brw_reg output_handles = get_tcs_output_urb_handle();
2875
2876 fs_inst *inst;
2877 if (indirect_offset.file == BAD_FILE) {
2878 /* This MOV replicates the output handle to all enabled channels
2879 * is SINGLE_PATCH mode.
2880 */
2881 fs_reg patch_handle = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
2882 bld.MOV(patch_handle, output_handles);
2883
2884 {
2885 if (first_component != 0) {
2886 unsigned read_components =
2887 instr->num_components + first_component;
2888 fs_reg tmp = bld.vgrf(dst.type, read_components);
2889 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, tmp,
2890 patch_handle);
2891 inst->size_written = read_components * REG_SIZE;
2892 for (unsigned i = 0; i < instr->num_components; i++) {
2893 bld.MOV(offset(dst, bld, i),
2894 offset(tmp, bld, i + first_component));
2895 }
2896 } else {
2897 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, dst,
2898 patch_handle);
2899 inst->size_written = instr->num_components * REG_SIZE;
2900 }
2901 inst->offset = imm_offset;
2902 inst->mlen = 1;
2903 }
2904 } else {
2905 /* Indirect indexing - use per-slot offsets as well. */
2906 const fs_reg srcs[] = { output_handles, indirect_offset };
2907 fs_reg payload = bld.vgrf(BRW_REGISTER_TYPE_UD, 2);
2908 bld.LOAD_PAYLOAD(payload, srcs, ARRAY_SIZE(srcs), 0);
2909 if (first_component != 0) {
2910 unsigned read_components =
2911 instr->num_components + first_component;
2912 fs_reg tmp = bld.vgrf(dst.type, read_components);
2913 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, tmp,
2914 payload);
2915 inst->size_written = read_components * REG_SIZE;
2916 for (unsigned i = 0; i < instr->num_components; i++) {
2917 bld.MOV(offset(dst, bld, i),
2918 offset(tmp, bld, i + first_component));
2919 }
2920 } else {
2921 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, dst,
2922 payload);
2923 inst->size_written = instr->num_components * REG_SIZE;
2924 }
2925 inst->offset = imm_offset;
2926 inst->mlen = 2;
2927 }
2928 break;
2929 }
2930
2931 case nir_intrinsic_store_output:
2932 case nir_intrinsic_store_per_vertex_output: {
2933 assert(nir_src_bit_size(instr->src[0]) == 32);
2934 fs_reg value = get_nir_src(instr->src[0]);
2935 fs_reg indirect_offset = get_indirect_offset(instr);
2936 unsigned imm_offset = instr->const_index[0];
2937 unsigned mask = instr->const_index[1];
2938 unsigned header_regs = 0;
2939 struct brw_reg output_handles = get_tcs_output_urb_handle();
2940
2941 fs_reg srcs[7];
2942 srcs[header_regs++] = output_handles;
2943
2944 if (indirect_offset.file != BAD_FILE) {
2945 srcs[header_regs++] = indirect_offset;
2946 }
2947
2948 if (mask == 0)
2949 break;
2950
2951 unsigned num_components = util_last_bit(mask);
2952 enum opcode opcode;
2953
2954 /* We can only pack two 64-bit components in a single message, so send
2955 * 2 messages if we have more components
2956 */
2957 unsigned first_component = nir_intrinsic_component(instr);
2958 mask = mask << first_component;
2959
2960 if (mask != WRITEMASK_XYZW) {
2961 srcs[header_regs++] = brw_imm_ud(mask << 16);
2962 opcode = indirect_offset.file != BAD_FILE ?
2963 SHADER_OPCODE_URB_WRITE_SIMD8_MASKED_PER_SLOT :
2964 SHADER_OPCODE_URB_WRITE_SIMD8_MASKED;
2965 } else {
2966 opcode = indirect_offset.file != BAD_FILE ?
2967 SHADER_OPCODE_URB_WRITE_SIMD8_PER_SLOT :
2968 SHADER_OPCODE_URB_WRITE_SIMD8;
2969 }
2970
2971 for (unsigned i = 0; i < num_components; i++) {
2972 if (!(mask & (1 << (i + first_component))))
2973 continue;
2974
2975 srcs[header_regs + i + first_component] = offset(value, bld, i);
2976 }
2977
2978 unsigned mlen = header_regs + num_components + first_component;
2979 fs_reg payload =
2980 bld.vgrf(BRW_REGISTER_TYPE_UD, mlen);
2981 bld.LOAD_PAYLOAD(payload, srcs, mlen, header_regs);
2982
2983 fs_inst *inst = bld.emit(opcode, bld.null_reg_ud(), payload);
2984 inst->offset = imm_offset;
2985 inst->mlen = mlen;
2986 break;
2987 }
2988
2989 default:
2990 nir_emit_intrinsic(bld, instr);
2991 break;
2992 }
2993 }
2994
2995 void
nir_emit_tes_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)2996 fs_visitor::nir_emit_tes_intrinsic(const fs_builder &bld,
2997 nir_intrinsic_instr *instr)
2998 {
2999 assert(stage == MESA_SHADER_TESS_EVAL);
3000 struct brw_tes_prog_data *tes_prog_data = brw_tes_prog_data(prog_data);
3001
3002 fs_reg dest;
3003 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
3004 dest = get_nir_dest(instr->dest);
3005
3006 switch (instr->intrinsic) {
3007 case nir_intrinsic_load_primitive_id:
3008 bld.MOV(dest, fs_reg(brw_vec1_grf(0, 1)));
3009 break;
3010 case nir_intrinsic_load_tess_coord:
3011 /* gl_TessCoord is part of the payload in g1-3 */
3012 for (unsigned i = 0; i < 3; i++) {
3013 bld.MOV(offset(dest, bld, i), fs_reg(brw_vec8_grf(1 + i, 0)));
3014 }
3015 break;
3016
3017 case nir_intrinsic_load_input:
3018 case nir_intrinsic_load_per_vertex_input: {
3019 assert(nir_dest_bit_size(instr->dest) == 32);
3020 fs_reg indirect_offset = get_indirect_offset(instr);
3021 unsigned imm_offset = instr->const_index[0];
3022 unsigned first_component = nir_intrinsic_component(instr);
3023
3024 fs_inst *inst;
3025 if (indirect_offset.file == BAD_FILE) {
3026 /* Arbitrarily only push up to 32 vec4 slots worth of data,
3027 * which is 16 registers (since each holds 2 vec4 slots).
3028 */
3029 const unsigned max_push_slots = 32;
3030 if (imm_offset < max_push_slots) {
3031 fs_reg src = fs_reg(ATTR, imm_offset / 2, dest.type);
3032 for (int i = 0; i < instr->num_components; i++) {
3033 unsigned comp = 4 * (imm_offset % 2) + i + first_component;
3034 bld.MOV(offset(dest, bld, i), component(src, comp));
3035 }
3036
3037 tes_prog_data->base.urb_read_length =
3038 MAX2(tes_prog_data->base.urb_read_length,
3039 (imm_offset / 2) + 1);
3040 } else {
3041 /* Replicate the patch handle to all enabled channels */
3042 const fs_reg srcs[] = {
3043 retype(brw_vec1_grf(0, 0), BRW_REGISTER_TYPE_UD)
3044 };
3045 fs_reg patch_handle = bld.vgrf(BRW_REGISTER_TYPE_UD, 1);
3046 bld.LOAD_PAYLOAD(patch_handle, srcs, ARRAY_SIZE(srcs), 0);
3047
3048 if (first_component != 0) {
3049 unsigned read_components =
3050 instr->num_components + first_component;
3051 fs_reg tmp = bld.vgrf(dest.type, read_components);
3052 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, tmp,
3053 patch_handle);
3054 inst->size_written = read_components * REG_SIZE;
3055 for (unsigned i = 0; i < instr->num_components; i++) {
3056 bld.MOV(offset(dest, bld, i),
3057 offset(tmp, bld, i + first_component));
3058 }
3059 } else {
3060 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8, dest,
3061 patch_handle);
3062 inst->size_written = instr->num_components * REG_SIZE;
3063 }
3064 inst->mlen = 1;
3065 inst->offset = imm_offset;
3066 }
3067 } else {
3068 /* Indirect indexing - use per-slot offsets as well. */
3069
3070 /* We can only read two double components with each URB read, so
3071 * we send two read messages in that case, each one loading up to
3072 * two double components.
3073 */
3074 unsigned num_components = instr->num_components;
3075 const fs_reg srcs[] = {
3076 retype(brw_vec1_grf(0, 0), BRW_REGISTER_TYPE_UD),
3077 indirect_offset
3078 };
3079 fs_reg payload = bld.vgrf(BRW_REGISTER_TYPE_UD, 2);
3080 bld.LOAD_PAYLOAD(payload, srcs, ARRAY_SIZE(srcs), 0);
3081
3082 if (first_component != 0) {
3083 unsigned read_components =
3084 num_components + first_component;
3085 fs_reg tmp = bld.vgrf(dest.type, read_components);
3086 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, tmp,
3087 payload);
3088 for (unsigned i = 0; i < num_components; i++) {
3089 bld.MOV(offset(dest, bld, i),
3090 offset(tmp, bld, i + first_component));
3091 }
3092 } else {
3093 inst = bld.emit(SHADER_OPCODE_URB_READ_SIMD8_PER_SLOT, dest,
3094 payload);
3095 }
3096 inst->mlen = 2;
3097 inst->offset = imm_offset;
3098 inst->size_written = (num_components + first_component) *
3099 inst->dst.component_size(inst->exec_size);
3100 }
3101 break;
3102 }
3103 default:
3104 nir_emit_intrinsic(bld, instr);
3105 break;
3106 }
3107 }
3108
3109 void
nir_emit_gs_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)3110 fs_visitor::nir_emit_gs_intrinsic(const fs_builder &bld,
3111 nir_intrinsic_instr *instr)
3112 {
3113 assert(stage == MESA_SHADER_GEOMETRY);
3114 fs_reg indirect_offset;
3115
3116 fs_reg dest;
3117 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
3118 dest = get_nir_dest(instr->dest);
3119
3120 switch (instr->intrinsic) {
3121 case nir_intrinsic_load_primitive_id:
3122 assert(stage == MESA_SHADER_GEOMETRY);
3123 assert(brw_gs_prog_data(prog_data)->include_primitive_id);
3124 bld.MOV(retype(dest, BRW_REGISTER_TYPE_UD),
3125 retype(fs_reg(brw_vec8_grf(2, 0)), BRW_REGISTER_TYPE_UD));
3126 break;
3127
3128 case nir_intrinsic_load_input:
3129 unreachable("load_input intrinsics are invalid for the GS stage");
3130
3131 case nir_intrinsic_load_per_vertex_input:
3132 emit_gs_input_load(dest, instr->src[0], instr->const_index[0],
3133 instr->src[1], instr->num_components,
3134 nir_intrinsic_component(instr));
3135 break;
3136
3137 case nir_intrinsic_emit_vertex_with_counter:
3138 emit_gs_vertex(instr->src[0], instr->const_index[0]);
3139 break;
3140
3141 case nir_intrinsic_end_primitive_with_counter:
3142 emit_gs_end_primitive(instr->src[0]);
3143 break;
3144
3145 case nir_intrinsic_set_vertex_and_primitive_count:
3146 bld.MOV(this->final_gs_vertex_count, get_nir_src(instr->src[0]));
3147 break;
3148
3149 case nir_intrinsic_load_invocation_id: {
3150 fs_reg val = nir_system_values[SYSTEM_VALUE_INVOCATION_ID];
3151 assert(val.file != BAD_FILE);
3152 dest.type = val.type;
3153 bld.MOV(dest, val);
3154 break;
3155 }
3156
3157 default:
3158 nir_emit_intrinsic(bld, instr);
3159 break;
3160 }
3161 }
3162
3163 /**
3164 * Fetch the current render target layer index.
3165 */
3166 static fs_reg
fetch_render_target_array_index(const fs_builder & bld)3167 fetch_render_target_array_index(const fs_builder &bld)
3168 {
3169 if (bld.shader->devinfo->gen >= 12) {
3170 /* The render target array index is provided in the thread payload as
3171 * bits 26:16 of r1.1.
3172 */
3173 const fs_reg idx = bld.vgrf(BRW_REGISTER_TYPE_UD);
3174 bld.AND(idx, brw_uw1_reg(BRW_GENERAL_REGISTER_FILE, 1, 3),
3175 brw_imm_uw(0x7ff));
3176 return idx;
3177 } else if (bld.shader->devinfo->gen >= 6) {
3178 /* The render target array index is provided in the thread payload as
3179 * bits 26:16 of r0.0.
3180 */
3181 const fs_reg idx = bld.vgrf(BRW_REGISTER_TYPE_UD);
3182 bld.AND(idx, brw_uw1_reg(BRW_GENERAL_REGISTER_FILE, 0, 1),
3183 brw_imm_uw(0x7ff));
3184 return idx;
3185 } else {
3186 /* Pre-SNB we only ever render into the first layer of the framebuffer
3187 * since layered rendering is not implemented.
3188 */
3189 return brw_imm_ud(0);
3190 }
3191 }
3192
3193 /**
3194 * Fake non-coherent framebuffer read implemented using TXF to fetch from the
3195 * framebuffer at the current fragment coordinates and sample index.
3196 */
3197 fs_inst *
emit_non_coherent_fb_read(const fs_builder & bld,const fs_reg & dst,unsigned target)3198 fs_visitor::emit_non_coherent_fb_read(const fs_builder &bld, const fs_reg &dst,
3199 unsigned target)
3200 {
3201 const struct gen_device_info *devinfo = bld.shader->devinfo;
3202
3203 assert(bld.shader->stage == MESA_SHADER_FRAGMENT);
3204 const brw_wm_prog_key *wm_key =
3205 reinterpret_cast<const brw_wm_prog_key *>(key);
3206 assert(!wm_key->coherent_fb_fetch);
3207 const struct brw_wm_prog_data *wm_prog_data =
3208 brw_wm_prog_data(stage_prog_data);
3209
3210 /* Calculate the surface index relative to the start of the texture binding
3211 * table block, since that's what the texturing messages expect.
3212 */
3213 const unsigned surface = target +
3214 wm_prog_data->binding_table.render_target_read_start -
3215 wm_prog_data->base.binding_table.texture_start;
3216
3217 /* Calculate the fragment coordinates. */
3218 const fs_reg coords = bld.vgrf(BRW_REGISTER_TYPE_UD, 3);
3219 bld.MOV(offset(coords, bld, 0), pixel_x);
3220 bld.MOV(offset(coords, bld, 1), pixel_y);
3221 bld.MOV(offset(coords, bld, 2), fetch_render_target_array_index(bld));
3222
3223 /* Calculate the sample index and MCS payload when multisampling. Luckily
3224 * the MCS fetch message behaves deterministically for UMS surfaces, so it
3225 * shouldn't be necessary to recompile based on whether the framebuffer is
3226 * CMS or UMS.
3227 */
3228 if (wm_key->multisample_fbo &&
3229 nir_system_values[SYSTEM_VALUE_SAMPLE_ID].file == BAD_FILE)
3230 nir_system_values[SYSTEM_VALUE_SAMPLE_ID] = *emit_sampleid_setup();
3231
3232 const fs_reg sample = nir_system_values[SYSTEM_VALUE_SAMPLE_ID];
3233 const fs_reg mcs = wm_key->multisample_fbo ?
3234 emit_mcs_fetch(coords, 3, brw_imm_ud(surface), fs_reg()) : fs_reg();
3235
3236 /* Use either a normal or a CMS texel fetch message depending on whether
3237 * the framebuffer is single or multisample. On SKL+ use the wide CMS
3238 * message just in case the framebuffer uses 16x multisampling, it should
3239 * be equivalent to the normal CMS fetch for lower multisampling modes.
3240 */
3241 const opcode op = !wm_key->multisample_fbo ? SHADER_OPCODE_TXF_LOGICAL :
3242 devinfo->gen >= 9 ? SHADER_OPCODE_TXF_CMS_W_LOGICAL :
3243 SHADER_OPCODE_TXF_CMS_LOGICAL;
3244
3245 /* Emit the instruction. */
3246 fs_reg srcs[TEX_LOGICAL_NUM_SRCS];
3247 srcs[TEX_LOGICAL_SRC_COORDINATE] = coords;
3248 srcs[TEX_LOGICAL_SRC_LOD] = brw_imm_ud(0);
3249 srcs[TEX_LOGICAL_SRC_SAMPLE_INDEX] = sample;
3250 srcs[TEX_LOGICAL_SRC_MCS] = mcs;
3251 srcs[TEX_LOGICAL_SRC_SURFACE] = brw_imm_ud(surface);
3252 srcs[TEX_LOGICAL_SRC_SAMPLER] = brw_imm_ud(0);
3253 srcs[TEX_LOGICAL_SRC_COORD_COMPONENTS] = brw_imm_ud(3);
3254 srcs[TEX_LOGICAL_SRC_GRAD_COMPONENTS] = brw_imm_ud(0);
3255
3256 fs_inst *inst = bld.emit(op, dst, srcs, ARRAY_SIZE(srcs));
3257 inst->size_written = 4 * inst->dst.component_size(inst->exec_size);
3258
3259 return inst;
3260 }
3261
3262 /**
3263 * Actual coherent framebuffer read implemented using the native render target
3264 * read message. Requires SKL+.
3265 */
3266 static fs_inst *
emit_coherent_fb_read(const fs_builder & bld,const fs_reg & dst,unsigned target)3267 emit_coherent_fb_read(const fs_builder &bld, const fs_reg &dst, unsigned target)
3268 {
3269 assert(bld.shader->devinfo->gen >= 9);
3270 fs_inst *inst = bld.emit(FS_OPCODE_FB_READ_LOGICAL, dst);
3271 inst->target = target;
3272 inst->size_written = 4 * inst->dst.component_size(inst->exec_size);
3273
3274 return inst;
3275 }
3276
3277 static fs_reg
alloc_temporary(const fs_builder & bld,unsigned size,fs_reg * regs,unsigned n)3278 alloc_temporary(const fs_builder &bld, unsigned size, fs_reg *regs, unsigned n)
3279 {
3280 if (n && regs[0].file != BAD_FILE) {
3281 return regs[0];
3282
3283 } else {
3284 const fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_F, size);
3285
3286 for (unsigned i = 0; i < n; i++)
3287 regs[i] = tmp;
3288
3289 return tmp;
3290 }
3291 }
3292
3293 static fs_reg
alloc_frag_output(fs_visitor * v,unsigned location)3294 alloc_frag_output(fs_visitor *v, unsigned location)
3295 {
3296 assert(v->stage == MESA_SHADER_FRAGMENT);
3297 const brw_wm_prog_key *const key =
3298 reinterpret_cast<const brw_wm_prog_key *>(v->key);
3299 const unsigned l = GET_FIELD(location, BRW_NIR_FRAG_OUTPUT_LOCATION);
3300 const unsigned i = GET_FIELD(location, BRW_NIR_FRAG_OUTPUT_INDEX);
3301
3302 if (i > 0 || (key->force_dual_color_blend && l == FRAG_RESULT_DATA1))
3303 return alloc_temporary(v->bld, 4, &v->dual_src_output, 1);
3304
3305 else if (l == FRAG_RESULT_COLOR)
3306 return alloc_temporary(v->bld, 4, v->outputs,
3307 MAX2(key->nr_color_regions, 1));
3308
3309 else if (l == FRAG_RESULT_DEPTH)
3310 return alloc_temporary(v->bld, 1, &v->frag_depth, 1);
3311
3312 else if (l == FRAG_RESULT_STENCIL)
3313 return alloc_temporary(v->bld, 1, &v->frag_stencil, 1);
3314
3315 else if (l == FRAG_RESULT_SAMPLE_MASK)
3316 return alloc_temporary(v->bld, 1, &v->sample_mask, 1);
3317
3318 else if (l >= FRAG_RESULT_DATA0 &&
3319 l < FRAG_RESULT_DATA0 + BRW_MAX_DRAW_BUFFERS)
3320 return alloc_temporary(v->bld, 4,
3321 &v->outputs[l - FRAG_RESULT_DATA0], 1);
3322
3323 else
3324 unreachable("Invalid location");
3325 }
3326
3327 void
nir_emit_fs_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)3328 fs_visitor::nir_emit_fs_intrinsic(const fs_builder &bld,
3329 nir_intrinsic_instr *instr)
3330 {
3331 assert(stage == MESA_SHADER_FRAGMENT);
3332
3333 fs_reg dest;
3334 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
3335 dest = get_nir_dest(instr->dest);
3336
3337 switch (instr->intrinsic) {
3338 case nir_intrinsic_load_front_face:
3339 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D),
3340 *emit_frontfacing_interpolation());
3341 break;
3342
3343 case nir_intrinsic_load_sample_pos: {
3344 fs_reg sample_pos = nir_system_values[SYSTEM_VALUE_SAMPLE_POS];
3345 assert(sample_pos.file != BAD_FILE);
3346 dest.type = sample_pos.type;
3347 bld.MOV(dest, sample_pos);
3348 bld.MOV(offset(dest, bld, 1), offset(sample_pos, bld, 1));
3349 break;
3350 }
3351
3352 case nir_intrinsic_load_layer_id:
3353 dest.type = BRW_REGISTER_TYPE_UD;
3354 bld.MOV(dest, fetch_render_target_array_index(bld));
3355 break;
3356
3357 case nir_intrinsic_is_helper_invocation: {
3358 /* Unlike the regular gl_HelperInvocation, that is defined at dispatch,
3359 * the helperInvocationEXT() (aka SpvOpIsHelperInvocationEXT) takes into
3360 * consideration demoted invocations. That information is stored in
3361 * f0.1.
3362 */
3363 dest.type = BRW_REGISTER_TYPE_UD;
3364
3365 bld.MOV(dest, brw_imm_ud(0));
3366
3367 fs_inst *mov = bld.MOV(dest, brw_imm_ud(~0));
3368 mov->predicate = BRW_PREDICATE_NORMAL;
3369 mov->predicate_inverse = true;
3370 mov->flag_subreg = sample_mask_flag_subreg(this);
3371 break;
3372 }
3373
3374 case nir_intrinsic_load_helper_invocation:
3375 case nir_intrinsic_load_sample_mask_in:
3376 case nir_intrinsic_load_sample_id: {
3377 gl_system_value sv = nir_system_value_from_intrinsic(instr->intrinsic);
3378 fs_reg val = nir_system_values[sv];
3379 assert(val.file != BAD_FILE);
3380 dest.type = val.type;
3381 bld.MOV(dest, val);
3382 break;
3383 }
3384
3385 case nir_intrinsic_store_output: {
3386 const fs_reg src = get_nir_src(instr->src[0]);
3387 const unsigned store_offset = nir_src_as_uint(instr->src[1]);
3388 const unsigned location = nir_intrinsic_base(instr) +
3389 SET_FIELD(store_offset, BRW_NIR_FRAG_OUTPUT_LOCATION);
3390 const fs_reg new_dest = retype(alloc_frag_output(this, location),
3391 src.type);
3392
3393 for (unsigned j = 0; j < instr->num_components; j++)
3394 bld.MOV(offset(new_dest, bld, nir_intrinsic_component(instr) + j),
3395 offset(src, bld, j));
3396
3397 break;
3398 }
3399
3400 case nir_intrinsic_load_output: {
3401 const unsigned l = GET_FIELD(nir_intrinsic_base(instr),
3402 BRW_NIR_FRAG_OUTPUT_LOCATION);
3403 assert(l >= FRAG_RESULT_DATA0);
3404 const unsigned load_offset = nir_src_as_uint(instr->src[0]);
3405 const unsigned target = l - FRAG_RESULT_DATA0 + load_offset;
3406 const fs_reg tmp = bld.vgrf(dest.type, 4);
3407
3408 if (reinterpret_cast<const brw_wm_prog_key *>(key)->coherent_fb_fetch)
3409 emit_coherent_fb_read(bld, tmp, target);
3410 else
3411 emit_non_coherent_fb_read(bld, tmp, target);
3412
3413 for (unsigned j = 0; j < instr->num_components; j++) {
3414 bld.MOV(offset(dest, bld, j),
3415 offset(tmp, bld, nir_intrinsic_component(instr) + j));
3416 }
3417
3418 break;
3419 }
3420
3421 case nir_intrinsic_demote:
3422 case nir_intrinsic_discard:
3423 case nir_intrinsic_terminate:
3424 case nir_intrinsic_demote_if:
3425 case nir_intrinsic_discard_if:
3426 case nir_intrinsic_terminate_if: {
3427 /* We track our discarded pixels in f0.1/f1.0. By predicating on it, we
3428 * can update just the flag bits that aren't yet discarded. If there's
3429 * no condition, we emit a CMP of g0 != g0, so all currently executing
3430 * channels will get turned off.
3431 */
3432 fs_inst *cmp = NULL;
3433 if (instr->intrinsic == nir_intrinsic_demote_if ||
3434 instr->intrinsic == nir_intrinsic_discard_if ||
3435 instr->intrinsic == nir_intrinsic_terminate_if) {
3436 nir_alu_instr *alu = nir_src_as_alu_instr(instr->src[0]);
3437
3438 if (alu != NULL &&
3439 alu->op != nir_op_bcsel &&
3440 (devinfo->gen > 5 ||
3441 (alu->instr.pass_flags & BRW_NIR_BOOLEAN_MASK) != BRW_NIR_BOOLEAN_NEEDS_RESOLVE ||
3442 alu->op == nir_op_fneu32 || alu->op == nir_op_feq32 ||
3443 alu->op == nir_op_flt32 || alu->op == nir_op_fge32 ||
3444 alu->op == nir_op_ine32 || alu->op == nir_op_ieq32 ||
3445 alu->op == nir_op_ilt32 || alu->op == nir_op_ige32 ||
3446 alu->op == nir_op_ult32 || alu->op == nir_op_uge32)) {
3447 /* Re-emit the instruction that generated the Boolean value, but
3448 * do not store it. Since this instruction will be conditional,
3449 * other instructions that want to use the real Boolean value may
3450 * get garbage. This was a problem for piglit's fs-discard-exit-2
3451 * test.
3452 *
3453 * Ideally we'd detect that the instruction cannot have a
3454 * conditional modifier before emitting the instructions. Alas,
3455 * that is nigh impossible. Instead, we're going to assume the
3456 * instruction (or last instruction) generated can have a
3457 * conditional modifier. If it cannot, fallback to the old-style
3458 * compare, and hope dead code elimination will clean up the
3459 * extra instructions generated.
3460 */
3461 nir_emit_alu(bld, alu, false);
3462
3463 cmp = (fs_inst *) instructions.get_tail();
3464 if (cmp->conditional_mod == BRW_CONDITIONAL_NONE) {
3465 if (cmp->can_do_cmod())
3466 cmp->conditional_mod = BRW_CONDITIONAL_Z;
3467 else
3468 cmp = NULL;
3469 } else {
3470 /* The old sequence that would have been generated is,
3471 * basically, bool_result == false. This is equivalent to
3472 * !bool_result, so negate the old modifier.
3473 */
3474 cmp->conditional_mod = brw_negate_cmod(cmp->conditional_mod);
3475 }
3476 }
3477
3478 if (cmp == NULL) {
3479 cmp = bld.CMP(bld.null_reg_f(), get_nir_src(instr->src[0]),
3480 brw_imm_d(0), BRW_CONDITIONAL_Z);
3481 }
3482 } else {
3483 fs_reg some_reg = fs_reg(retype(brw_vec8_grf(0, 0),
3484 BRW_REGISTER_TYPE_UW));
3485 cmp = bld.CMP(bld.null_reg_f(), some_reg, some_reg, BRW_CONDITIONAL_NZ);
3486 }
3487
3488 cmp->predicate = BRW_PREDICATE_NORMAL;
3489 cmp->flag_subreg = sample_mask_flag_subreg(this);
3490
3491 fs_inst *jump = bld.emit(FS_OPCODE_DISCARD_JUMP);
3492 jump->flag_subreg = sample_mask_flag_subreg(this);
3493 jump->predicate_inverse = true;
3494
3495 if (instr->intrinsic == nir_intrinsic_terminate ||
3496 instr->intrinsic == nir_intrinsic_terminate_if) {
3497 jump->predicate = BRW_PREDICATE_NORMAL;
3498 } else {
3499 /* Only jump when the whole quad is demoted. For historical
3500 * reasons this is also used for discard.
3501 */
3502 jump->predicate = BRW_PREDICATE_ALIGN1_ANY4H;
3503 }
3504
3505 if (devinfo->gen < 7)
3506 limit_dispatch_width(
3507 16, "Fragment discard/demote not implemented in SIMD32 mode.\n");
3508 break;
3509 }
3510
3511 case nir_intrinsic_load_input: {
3512 /* load_input is only used for flat inputs */
3513 assert(nir_dest_bit_size(instr->dest) == 32);
3514 unsigned base = nir_intrinsic_base(instr);
3515 unsigned comp = nir_intrinsic_component(instr);
3516 unsigned num_components = instr->num_components;
3517
3518 /* Special case fields in the VUE header */
3519 if (base == VARYING_SLOT_LAYER)
3520 comp = 1;
3521 else if (base == VARYING_SLOT_VIEWPORT)
3522 comp = 2;
3523
3524 for (unsigned int i = 0; i < num_components; i++) {
3525 bld.MOV(offset(dest, bld, i),
3526 retype(component(interp_reg(base, comp + i), 3), dest.type));
3527 }
3528 break;
3529 }
3530
3531 case nir_intrinsic_load_fs_input_interp_deltas: {
3532 assert(stage == MESA_SHADER_FRAGMENT);
3533 assert(nir_src_as_uint(instr->src[0]) == 0);
3534 fs_reg interp = interp_reg(nir_intrinsic_base(instr),
3535 nir_intrinsic_component(instr));
3536 dest.type = BRW_REGISTER_TYPE_F;
3537 bld.MOV(offset(dest, bld, 0), component(interp, 3));
3538 bld.MOV(offset(dest, bld, 1), component(interp, 1));
3539 bld.MOV(offset(dest, bld, 2), component(interp, 0));
3540 break;
3541 }
3542
3543 case nir_intrinsic_load_barycentric_pixel:
3544 case nir_intrinsic_load_barycentric_centroid:
3545 case nir_intrinsic_load_barycentric_sample: {
3546 /* Use the delta_xy values computed from the payload */
3547 const glsl_interp_mode interp_mode =
3548 (enum glsl_interp_mode) nir_intrinsic_interp_mode(instr);
3549 enum brw_barycentric_mode bary =
3550 brw_barycentric_mode(interp_mode, instr->intrinsic);
3551 const fs_reg srcs[] = { offset(this->delta_xy[bary], bld, 0),
3552 offset(this->delta_xy[bary], bld, 1) };
3553 bld.LOAD_PAYLOAD(dest, srcs, ARRAY_SIZE(srcs), 0);
3554 break;
3555 }
3556
3557 case nir_intrinsic_load_barycentric_at_sample: {
3558 const glsl_interp_mode interpolation =
3559 (enum glsl_interp_mode) nir_intrinsic_interp_mode(instr);
3560
3561 if (nir_src_is_const(instr->src[0])) {
3562 unsigned msg_data = nir_src_as_uint(instr->src[0]) << 4;
3563
3564 emit_pixel_interpolater_send(bld,
3565 FS_OPCODE_INTERPOLATE_AT_SAMPLE,
3566 dest,
3567 fs_reg(), /* src */
3568 brw_imm_ud(msg_data),
3569 interpolation);
3570 } else {
3571 const fs_reg sample_src = retype(get_nir_src(instr->src[0]),
3572 BRW_REGISTER_TYPE_UD);
3573
3574 if (nir_src_is_dynamically_uniform(instr->src[0])) {
3575 const fs_reg sample_id = bld.emit_uniformize(sample_src);
3576 const fs_reg msg_data = vgrf(glsl_type::uint_type);
3577 bld.exec_all().group(1, 0)
3578 .SHL(msg_data, sample_id, brw_imm_ud(4u));
3579 emit_pixel_interpolater_send(bld,
3580 FS_OPCODE_INTERPOLATE_AT_SAMPLE,
3581 dest,
3582 fs_reg(), /* src */
3583 component(msg_data, 0),
3584 interpolation);
3585 } else {
3586 /* Make a loop that sends a message to the pixel interpolater
3587 * for the sample number in each live channel. If there are
3588 * multiple channels with the same sample number then these
3589 * will be handled simultaneously with a single interation of
3590 * the loop.
3591 */
3592 bld.emit(BRW_OPCODE_DO);
3593
3594 /* Get the next live sample number into sample_id_reg */
3595 const fs_reg sample_id = bld.emit_uniformize(sample_src);
3596
3597 /* Set the flag register so that we can perform the send
3598 * message on all channels that have the same sample number
3599 */
3600 bld.CMP(bld.null_reg_ud(),
3601 sample_src, sample_id,
3602 BRW_CONDITIONAL_EQ);
3603 const fs_reg msg_data = vgrf(glsl_type::uint_type);
3604 bld.exec_all().group(1, 0)
3605 .SHL(msg_data, sample_id, brw_imm_ud(4u));
3606 fs_inst *inst =
3607 emit_pixel_interpolater_send(bld,
3608 FS_OPCODE_INTERPOLATE_AT_SAMPLE,
3609 dest,
3610 fs_reg(), /* src */
3611 component(msg_data, 0),
3612 interpolation);
3613 set_predicate(BRW_PREDICATE_NORMAL, inst);
3614
3615 /* Continue the loop if there are any live channels left */
3616 set_predicate_inv(BRW_PREDICATE_NORMAL,
3617 true, /* inverse */
3618 bld.emit(BRW_OPCODE_WHILE));
3619 }
3620 }
3621 break;
3622 }
3623
3624 case nir_intrinsic_load_barycentric_at_offset: {
3625 const glsl_interp_mode interpolation =
3626 (enum glsl_interp_mode) nir_intrinsic_interp_mode(instr);
3627
3628 nir_const_value *const_offset = nir_src_as_const_value(instr->src[0]);
3629
3630 if (const_offset) {
3631 assert(nir_src_bit_size(instr->src[0]) == 32);
3632 unsigned off_x = MIN2((int)(const_offset[0].f32 * 16), 7) & 0xf;
3633 unsigned off_y = MIN2((int)(const_offset[1].f32 * 16), 7) & 0xf;
3634
3635 emit_pixel_interpolater_send(bld,
3636 FS_OPCODE_INTERPOLATE_AT_SHARED_OFFSET,
3637 dest,
3638 fs_reg(), /* src */
3639 brw_imm_ud(off_x | (off_y << 4)),
3640 interpolation);
3641 } else {
3642 fs_reg src = vgrf(glsl_type::ivec2_type);
3643 fs_reg offset_src = retype(get_nir_src(instr->src[0]),
3644 BRW_REGISTER_TYPE_F);
3645 for (int i = 0; i < 2; i++) {
3646 fs_reg temp = vgrf(glsl_type::float_type);
3647 bld.MUL(temp, offset(offset_src, bld, i), brw_imm_f(16.0f));
3648 fs_reg itemp = vgrf(glsl_type::int_type);
3649 /* float to int */
3650 bld.MOV(itemp, temp);
3651
3652 /* Clamp the upper end of the range to +7/16.
3653 * ARB_gpu_shader5 requires that we support a maximum offset
3654 * of +0.5, which isn't representable in a S0.4 value -- if
3655 * we didn't clamp it, we'd end up with -8/16, which is the
3656 * opposite of what the shader author wanted.
3657 *
3658 * This is legal due to ARB_gpu_shader5's quantization
3659 * rules:
3660 *
3661 * "Not all values of <offset> may be supported; x and y
3662 * offsets may be rounded to fixed-point values with the
3663 * number of fraction bits given by the
3664 * implementation-dependent constant
3665 * FRAGMENT_INTERPOLATION_OFFSET_BITS"
3666 */
3667 set_condmod(BRW_CONDITIONAL_L,
3668 bld.SEL(offset(src, bld, i), itemp, brw_imm_d(7)));
3669 }
3670
3671 const enum opcode opcode = FS_OPCODE_INTERPOLATE_AT_PER_SLOT_OFFSET;
3672 emit_pixel_interpolater_send(bld,
3673 opcode,
3674 dest,
3675 src,
3676 brw_imm_ud(0u),
3677 interpolation);
3678 }
3679 break;
3680 }
3681
3682 case nir_intrinsic_load_frag_coord:
3683 emit_fragcoord_interpolation(dest);
3684 break;
3685
3686 case nir_intrinsic_load_interpolated_input: {
3687 assert(instr->src[0].ssa &&
3688 instr->src[0].ssa->parent_instr->type == nir_instr_type_intrinsic);
3689 nir_intrinsic_instr *bary_intrinsic =
3690 nir_instr_as_intrinsic(instr->src[0].ssa->parent_instr);
3691 nir_intrinsic_op bary_intrin = bary_intrinsic->intrinsic;
3692 enum glsl_interp_mode interp_mode =
3693 (enum glsl_interp_mode) nir_intrinsic_interp_mode(bary_intrinsic);
3694 fs_reg dst_xy;
3695
3696 if (bary_intrin == nir_intrinsic_load_barycentric_at_offset ||
3697 bary_intrin == nir_intrinsic_load_barycentric_at_sample) {
3698 /* Use the result of the PI message. */
3699 dst_xy = retype(get_nir_src(instr->src[0]), BRW_REGISTER_TYPE_F);
3700 } else {
3701 /* Use the delta_xy values computed from the payload */
3702 enum brw_barycentric_mode bary =
3703 brw_barycentric_mode(interp_mode, bary_intrin);
3704 dst_xy = this->delta_xy[bary];
3705 }
3706
3707 for (unsigned int i = 0; i < instr->num_components; i++) {
3708 fs_reg interp =
3709 component(interp_reg(nir_intrinsic_base(instr),
3710 nir_intrinsic_component(instr) + i), 0);
3711 interp.type = BRW_REGISTER_TYPE_F;
3712 dest.type = BRW_REGISTER_TYPE_F;
3713
3714 if (devinfo->gen < 6 && interp_mode == INTERP_MODE_SMOOTH) {
3715 fs_reg tmp = vgrf(glsl_type::float_type);
3716 bld.emit(FS_OPCODE_LINTERP, tmp, dst_xy, interp);
3717 bld.MUL(offset(dest, bld, i), tmp, this->pixel_w);
3718 } else {
3719 bld.emit(FS_OPCODE_LINTERP, offset(dest, bld, i), dst_xy, interp);
3720 }
3721 }
3722 break;
3723 }
3724
3725 default:
3726 nir_emit_intrinsic(bld, instr);
3727 break;
3728 }
3729 }
3730
3731 void
nir_emit_cs_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)3732 fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
3733 nir_intrinsic_instr *instr)
3734 {
3735 assert(stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL);
3736 struct brw_cs_prog_data *cs_prog_data = brw_cs_prog_data(prog_data);
3737
3738 fs_reg dest;
3739 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
3740 dest = get_nir_dest(instr->dest);
3741
3742 switch (instr->intrinsic) {
3743 case nir_intrinsic_control_barrier:
3744 /* The whole workgroup fits in a single HW thread, so all the
3745 * invocations are already executed lock-step. Instead of an actual
3746 * barrier just emit a scheduling fence, that will generate no code.
3747 */
3748 if (!nir->info.cs.local_size_variable &&
3749 workgroup_size() <= dispatch_width) {
3750 bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
3751 break;
3752 }
3753
3754 emit_barrier();
3755 cs_prog_data->uses_barrier = true;
3756 break;
3757
3758 case nir_intrinsic_load_subgroup_id:
3759 bld.MOV(retype(dest, BRW_REGISTER_TYPE_UD), subgroup_id);
3760 break;
3761
3762 case nir_intrinsic_load_local_invocation_id:
3763 case nir_intrinsic_load_work_group_id: {
3764 gl_system_value sv = nir_system_value_from_intrinsic(instr->intrinsic);
3765 fs_reg val = nir_system_values[sv];
3766 assert(val.file != BAD_FILE);
3767 dest.type = val.type;
3768 for (unsigned i = 0; i < 3; i++)
3769 bld.MOV(offset(dest, bld, i), offset(val, bld, i));
3770 break;
3771 }
3772
3773 case nir_intrinsic_load_num_work_groups: {
3774 assert(nir_dest_bit_size(instr->dest) == 32);
3775 const unsigned surface =
3776 cs_prog_data->binding_table.work_groups_start;
3777
3778 cs_prog_data->uses_num_work_groups = true;
3779
3780 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
3781 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(surface);
3782 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
3783 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(3); /* num components */
3784 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = brw_imm_ud(0);
3785 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
3786 fs_inst *inst =
3787 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_READ_LOGICAL,
3788 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
3789 inst->size_written = 3 * dispatch_width * 4;
3790 break;
3791 }
3792
3793 case nir_intrinsic_shared_atomic_add:
3794 case nir_intrinsic_shared_atomic_imin:
3795 case nir_intrinsic_shared_atomic_umin:
3796 case nir_intrinsic_shared_atomic_imax:
3797 case nir_intrinsic_shared_atomic_umax:
3798 case nir_intrinsic_shared_atomic_and:
3799 case nir_intrinsic_shared_atomic_or:
3800 case nir_intrinsic_shared_atomic_xor:
3801 case nir_intrinsic_shared_atomic_exchange:
3802 case nir_intrinsic_shared_atomic_comp_swap:
3803 nir_emit_shared_atomic(bld, brw_aop_for_nir_intrinsic(instr), instr);
3804 break;
3805 case nir_intrinsic_shared_atomic_fmin:
3806 case nir_intrinsic_shared_atomic_fmax:
3807 case nir_intrinsic_shared_atomic_fcomp_swap:
3808 nir_emit_shared_atomic_float(bld, brw_aop_for_nir_intrinsic(instr), instr);
3809 break;
3810
3811 case nir_intrinsic_load_shared: {
3812 assert(devinfo->gen >= 7);
3813 assert(stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL);
3814
3815 const unsigned bit_size = nir_dest_bit_size(instr->dest);
3816 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
3817 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(GEN7_BTI_SLM);
3818 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[0]);
3819 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
3820 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
3821
3822 /* Make dest unsigned because that's what the temporary will be */
3823 dest.type = brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
3824
3825 /* Read the vector */
3826 assert(nir_dest_bit_size(instr->dest) <= 32);
3827 assert(nir_intrinsic_align(instr) > 0);
3828 if (nir_dest_bit_size(instr->dest) == 32 &&
3829 nir_intrinsic_align(instr) >= 4) {
3830 assert(nir_dest_num_components(instr->dest) <= 4);
3831 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
3832 fs_inst *inst =
3833 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_READ_LOGICAL,
3834 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
3835 inst->size_written = instr->num_components * dispatch_width * 4;
3836 } else {
3837 assert(nir_dest_num_components(instr->dest) == 1);
3838 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(bit_size);
3839
3840 fs_reg read_result = bld.vgrf(BRW_REGISTER_TYPE_UD);
3841 bld.emit(SHADER_OPCODE_BYTE_SCATTERED_READ_LOGICAL,
3842 read_result, srcs, SURFACE_LOGICAL_NUM_SRCS);
3843 bld.MOV(dest, subscript(read_result, dest.type, 0));
3844 }
3845 break;
3846 }
3847
3848 case nir_intrinsic_store_shared: {
3849 assert(devinfo->gen >= 7);
3850 assert(stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL);
3851
3852 const unsigned bit_size = nir_src_bit_size(instr->src[0]);
3853 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
3854 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(GEN7_BTI_SLM);
3855 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
3856 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
3857 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
3858
3859 fs_reg data = get_nir_src(instr->src[0]);
3860 data.type = brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
3861
3862 assert(nir_src_bit_size(instr->src[0]) <= 32);
3863 assert(nir_intrinsic_write_mask(instr) ==
3864 (1u << instr->num_components) - 1);
3865 assert(nir_intrinsic_align(instr) > 0);
3866 if (nir_src_bit_size(instr->src[0]) == 32 &&
3867 nir_intrinsic_align(instr) >= 4) {
3868 assert(nir_src_num_components(instr->src[0]) <= 4);
3869 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
3870 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
3871 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_WRITE_LOGICAL,
3872 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
3873 } else {
3874 assert(nir_src_num_components(instr->src[0]) == 1);
3875 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(bit_size);
3876
3877 srcs[SURFACE_LOGICAL_SRC_DATA] = bld.vgrf(BRW_REGISTER_TYPE_UD);
3878 bld.MOV(srcs[SURFACE_LOGICAL_SRC_DATA], data);
3879
3880 bld.emit(SHADER_OPCODE_BYTE_SCATTERED_WRITE_LOGICAL,
3881 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
3882 }
3883 break;
3884 }
3885
3886 case nir_intrinsic_load_local_group_size: {
3887 assert(compiler->lower_variable_group_size);
3888 assert(nir->info.cs.local_size_variable);
3889 for (unsigned i = 0; i < 3; i++) {
3890 bld.MOV(retype(offset(dest, bld, i), BRW_REGISTER_TYPE_UD),
3891 group_size[i]);
3892 }
3893 break;
3894 }
3895
3896 default:
3897 nir_emit_intrinsic(bld, instr);
3898 break;
3899 }
3900 }
3901
3902 static fs_reg
brw_nir_reduction_op_identity(const fs_builder & bld,nir_op op,brw_reg_type type)3903 brw_nir_reduction_op_identity(const fs_builder &bld,
3904 nir_op op, brw_reg_type type)
3905 {
3906 nir_const_value value = nir_alu_binop_identity(op, type_sz(type) * 8);
3907 switch (type_sz(type)) {
3908 case 1:
3909 if (type == BRW_REGISTER_TYPE_UB) {
3910 return brw_imm_uw(value.u8);
3911 } else {
3912 assert(type == BRW_REGISTER_TYPE_B);
3913 return brw_imm_w(value.i8);
3914 }
3915 case 2:
3916 return retype(brw_imm_uw(value.u16), type);
3917 case 4:
3918 return retype(brw_imm_ud(value.u32), type);
3919 case 8:
3920 if (type == BRW_REGISTER_TYPE_DF)
3921 return setup_imm_df(bld, value.f64);
3922 else
3923 return retype(brw_imm_u64(value.u64), type);
3924 default:
3925 unreachable("Invalid type size");
3926 }
3927 }
3928
3929 static opcode
brw_op_for_nir_reduction_op(nir_op op)3930 brw_op_for_nir_reduction_op(nir_op op)
3931 {
3932 switch (op) {
3933 case nir_op_iadd: return BRW_OPCODE_ADD;
3934 case nir_op_fadd: return BRW_OPCODE_ADD;
3935 case nir_op_imul: return BRW_OPCODE_MUL;
3936 case nir_op_fmul: return BRW_OPCODE_MUL;
3937 case nir_op_imin: return BRW_OPCODE_SEL;
3938 case nir_op_umin: return BRW_OPCODE_SEL;
3939 case nir_op_fmin: return BRW_OPCODE_SEL;
3940 case nir_op_imax: return BRW_OPCODE_SEL;
3941 case nir_op_umax: return BRW_OPCODE_SEL;
3942 case nir_op_fmax: return BRW_OPCODE_SEL;
3943 case nir_op_iand: return BRW_OPCODE_AND;
3944 case nir_op_ior: return BRW_OPCODE_OR;
3945 case nir_op_ixor: return BRW_OPCODE_XOR;
3946 default:
3947 unreachable("Invalid reduction operation");
3948 }
3949 }
3950
3951 static brw_conditional_mod
brw_cond_mod_for_nir_reduction_op(nir_op op)3952 brw_cond_mod_for_nir_reduction_op(nir_op op)
3953 {
3954 switch (op) {
3955 case nir_op_iadd: return BRW_CONDITIONAL_NONE;
3956 case nir_op_fadd: return BRW_CONDITIONAL_NONE;
3957 case nir_op_imul: return BRW_CONDITIONAL_NONE;
3958 case nir_op_fmul: return BRW_CONDITIONAL_NONE;
3959 case nir_op_imin: return BRW_CONDITIONAL_L;
3960 case nir_op_umin: return BRW_CONDITIONAL_L;
3961 case nir_op_fmin: return BRW_CONDITIONAL_L;
3962 case nir_op_imax: return BRW_CONDITIONAL_GE;
3963 case nir_op_umax: return BRW_CONDITIONAL_GE;
3964 case nir_op_fmax: return BRW_CONDITIONAL_GE;
3965 case nir_op_iand: return BRW_CONDITIONAL_NONE;
3966 case nir_op_ior: return BRW_CONDITIONAL_NONE;
3967 case nir_op_ixor: return BRW_CONDITIONAL_NONE;
3968 default:
3969 unreachable("Invalid reduction operation");
3970 }
3971 }
3972
3973 fs_reg
get_nir_image_intrinsic_image(const brw::fs_builder & bld,nir_intrinsic_instr * instr)3974 fs_visitor::get_nir_image_intrinsic_image(const brw::fs_builder &bld,
3975 nir_intrinsic_instr *instr)
3976 {
3977 fs_reg image = retype(get_nir_src_imm(instr->src[0]), BRW_REGISTER_TYPE_UD);
3978 fs_reg surf_index = image;
3979
3980 if (stage_prog_data->binding_table.image_start > 0) {
3981 if (image.file == BRW_IMMEDIATE_VALUE) {
3982 surf_index =
3983 brw_imm_ud(image.d + stage_prog_data->binding_table.image_start);
3984 } else {
3985 surf_index = vgrf(glsl_type::uint_type);
3986 bld.ADD(surf_index, image,
3987 brw_imm_d(stage_prog_data->binding_table.image_start));
3988 }
3989 }
3990
3991 return bld.emit_uniformize(surf_index);
3992 }
3993
3994 fs_reg
get_nir_ssbo_intrinsic_index(const brw::fs_builder & bld,nir_intrinsic_instr * instr)3995 fs_visitor::get_nir_ssbo_intrinsic_index(const brw::fs_builder &bld,
3996 nir_intrinsic_instr *instr)
3997 {
3998 /* SSBO stores are weird in that their index is in src[1] */
3999 const bool is_store =
4000 instr->intrinsic == nir_intrinsic_store_ssbo ||
4001 instr->intrinsic == nir_intrinsic_store_ssbo_block_intel;
4002 const unsigned src = is_store ? 1 : 0;
4003
4004 if (nir_src_is_const(instr->src[src])) {
4005 unsigned index = stage_prog_data->binding_table.ssbo_start +
4006 nir_src_as_uint(instr->src[src]);
4007 return brw_imm_ud(index);
4008 } else {
4009 fs_reg surf_index = vgrf(glsl_type::uint_type);
4010 bld.ADD(surf_index, get_nir_src(instr->src[src]),
4011 brw_imm_ud(stage_prog_data->binding_table.ssbo_start));
4012 return bld.emit_uniformize(surf_index);
4013 }
4014 }
4015
4016 /**
4017 * The offsets we get from NIR act as if each SIMD channel has it's own blob
4018 * of contiguous space. However, if we actually place each SIMD channel in
4019 * it's own space, we end up with terrible cache performance because each SIMD
4020 * channel accesses a different cache line even when they're all accessing the
4021 * same byte offset. To deal with this problem, we swizzle the address using
4022 * a simple algorithm which ensures that any time a SIMD message reads or
4023 * writes the same address, it's all in the same cache line. We have to keep
4024 * the bottom two bits fixed so that we can read/write up to a dword at a time
4025 * and the individual element is contiguous. We do this by splitting the
4026 * address as follows:
4027 *
4028 * 31 4-6 2 0
4029 * +-------------------------------+------------+----------+
4030 * | Hi address bits | chan index | addr low |
4031 * +-------------------------------+------------+----------+
4032 *
4033 * In other words, the bottom two address bits stay, and the top 30 get
4034 * shifted up so that we can stick the SIMD channel index in the middle. This
4035 * way, we can access 8, 16, or 32-bit elements and, when accessing a 32-bit
4036 * at the same logical offset, the scratch read/write instruction acts on
4037 * continuous elements and we get good cache locality.
4038 */
4039 fs_reg
swizzle_nir_scratch_addr(const brw::fs_builder & bld,const fs_reg & nir_addr,bool in_dwords)4040 fs_visitor::swizzle_nir_scratch_addr(const brw::fs_builder &bld,
4041 const fs_reg &nir_addr,
4042 bool in_dwords)
4043 {
4044 const fs_reg &chan_index =
4045 nir_system_values[SYSTEM_VALUE_SUBGROUP_INVOCATION];
4046 const unsigned chan_index_bits = ffs(dispatch_width) - 1;
4047
4048 fs_reg addr = bld.vgrf(BRW_REGISTER_TYPE_UD);
4049 if (in_dwords) {
4050 /* In this case, we know the address is aligned to a DWORD and we want
4051 * the final address in DWORDs.
4052 */
4053 bld.SHL(addr, nir_addr, brw_imm_ud(chan_index_bits - 2));
4054 bld.OR(addr, addr, chan_index);
4055 } else {
4056 /* This case substantially more annoying because we have to pay
4057 * attention to those pesky two bottom bits.
4058 */
4059 fs_reg addr_hi = bld.vgrf(BRW_REGISTER_TYPE_UD);
4060 bld.AND(addr_hi, nir_addr, brw_imm_ud(~0x3u));
4061 bld.SHL(addr_hi, addr_hi, brw_imm_ud(chan_index_bits));
4062 fs_reg chan_addr = bld.vgrf(BRW_REGISTER_TYPE_UD);
4063 bld.SHL(chan_addr, chan_index, brw_imm_ud(2));
4064 bld.AND(addr, nir_addr, brw_imm_ud(0x3u));
4065 bld.OR(addr, addr, addr_hi);
4066 bld.OR(addr, addr, chan_addr);
4067 }
4068 return addr;
4069 }
4070
4071 static unsigned
choose_oword_block_size_dwords(unsigned dwords)4072 choose_oword_block_size_dwords(unsigned dwords)
4073 {
4074 unsigned block;
4075 if (dwords >= 32) {
4076 block = 32;
4077 } else if (dwords >= 16) {
4078 block = 16;
4079 } else {
4080 block = 8;
4081 }
4082 assert(block <= dwords);
4083 return block;
4084 }
4085
4086 static void
increment_a64_address(const fs_builder & bld,fs_reg address,uint32_t v)4087 increment_a64_address(const fs_builder &bld, fs_reg address, uint32_t v)
4088 {
4089 if (bld.shader->devinfo->has_64bit_int) {
4090 bld.ADD(address, address, brw_imm_ud(v));
4091 } else {
4092 fs_reg low = retype(address, BRW_REGISTER_TYPE_UD);
4093 fs_reg high = offset(low, bld, 1);
4094
4095 /* Add low and if that overflows, add carry to high. */
4096 bld.ADD(low, low, brw_imm_ud(v))->conditional_mod = BRW_CONDITIONAL_O;
4097 bld.ADD(high, high, brw_imm_ud(0x1))->predicate = BRW_PREDICATE_NORMAL;
4098 }
4099 }
4100
4101 void
nir_emit_intrinsic(const fs_builder & bld,nir_intrinsic_instr * instr)4102 fs_visitor::nir_emit_intrinsic(const fs_builder &bld, nir_intrinsic_instr *instr)
4103 {
4104 fs_reg dest;
4105 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
4106 dest = get_nir_dest(instr->dest);
4107
4108 switch (instr->intrinsic) {
4109 case nir_intrinsic_image_load:
4110 case nir_intrinsic_image_store:
4111 case nir_intrinsic_image_atomic_add:
4112 case nir_intrinsic_image_atomic_imin:
4113 case nir_intrinsic_image_atomic_umin:
4114 case nir_intrinsic_image_atomic_imax:
4115 case nir_intrinsic_image_atomic_umax:
4116 case nir_intrinsic_image_atomic_and:
4117 case nir_intrinsic_image_atomic_or:
4118 case nir_intrinsic_image_atomic_xor:
4119 case nir_intrinsic_image_atomic_exchange:
4120 case nir_intrinsic_image_atomic_comp_swap:
4121 case nir_intrinsic_bindless_image_load:
4122 case nir_intrinsic_bindless_image_store:
4123 case nir_intrinsic_bindless_image_atomic_add:
4124 case nir_intrinsic_bindless_image_atomic_imin:
4125 case nir_intrinsic_bindless_image_atomic_umin:
4126 case nir_intrinsic_bindless_image_atomic_imax:
4127 case nir_intrinsic_bindless_image_atomic_umax:
4128 case nir_intrinsic_bindless_image_atomic_and:
4129 case nir_intrinsic_bindless_image_atomic_or:
4130 case nir_intrinsic_bindless_image_atomic_xor:
4131 case nir_intrinsic_bindless_image_atomic_exchange:
4132 case nir_intrinsic_bindless_image_atomic_comp_swap: {
4133 /* Get some metadata from the image intrinsic. */
4134 const nir_intrinsic_info *info = &nir_intrinsic_infos[instr->intrinsic];
4135
4136 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4137
4138 switch (instr->intrinsic) {
4139 case nir_intrinsic_image_load:
4140 case nir_intrinsic_image_store:
4141 case nir_intrinsic_image_atomic_add:
4142 case nir_intrinsic_image_atomic_imin:
4143 case nir_intrinsic_image_atomic_umin:
4144 case nir_intrinsic_image_atomic_imax:
4145 case nir_intrinsic_image_atomic_umax:
4146 case nir_intrinsic_image_atomic_and:
4147 case nir_intrinsic_image_atomic_or:
4148 case nir_intrinsic_image_atomic_xor:
4149 case nir_intrinsic_image_atomic_exchange:
4150 case nir_intrinsic_image_atomic_comp_swap:
4151 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4152 get_nir_image_intrinsic_image(bld, instr);
4153 break;
4154
4155 default:
4156 /* Bindless */
4157 srcs[SURFACE_LOGICAL_SRC_SURFACE_HANDLE] =
4158 bld.emit_uniformize(get_nir_src(instr->src[0]));
4159 break;
4160 }
4161
4162 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
4163 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] =
4164 brw_imm_ud(nir_image_intrinsic_coord_components(instr));
4165
4166 /* Emit an image load, store or atomic op. */
4167 if (instr->intrinsic == nir_intrinsic_image_load ||
4168 instr->intrinsic == nir_intrinsic_bindless_image_load) {
4169 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
4170 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
4171 fs_inst *inst =
4172 bld.emit(SHADER_OPCODE_TYPED_SURFACE_READ_LOGICAL,
4173 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
4174 inst->size_written = instr->num_components * dispatch_width * 4;
4175 } else if (instr->intrinsic == nir_intrinsic_image_store ||
4176 instr->intrinsic == nir_intrinsic_bindless_image_store) {
4177 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
4178 srcs[SURFACE_LOGICAL_SRC_DATA] = get_nir_src(instr->src[3]);
4179 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
4180 bld.emit(SHADER_OPCODE_TYPED_SURFACE_WRITE_LOGICAL,
4181 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
4182 } else {
4183 unsigned num_srcs = info->num_srcs;
4184 int op = brw_aop_for_nir_intrinsic(instr);
4185 if (op == BRW_AOP_INC || op == BRW_AOP_DEC) {
4186 assert(num_srcs == 4);
4187 num_srcs = 3;
4188 }
4189
4190 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(op);
4191
4192 fs_reg data;
4193 if (num_srcs >= 4)
4194 data = get_nir_src(instr->src[3]);
4195 if (num_srcs >= 5) {
4196 fs_reg tmp = bld.vgrf(data.type, 2);
4197 fs_reg sources[2] = { data, get_nir_src(instr->src[4]) };
4198 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
4199 data = tmp;
4200 }
4201 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
4202 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
4203
4204 bld.emit(SHADER_OPCODE_TYPED_ATOMIC_LOGICAL,
4205 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
4206 }
4207 break;
4208 }
4209
4210 case nir_intrinsic_image_size:
4211 case nir_intrinsic_bindless_image_size: {
4212 /* Unlike the [un]typed load and store opcodes, the TXS that this turns
4213 * into will handle the binding table index for us in the geneerator.
4214 * Incidentally, this means that we can handle bindless with exactly the
4215 * same code.
4216 */
4217 fs_reg image = retype(get_nir_src_imm(instr->src[0]),
4218 BRW_REGISTER_TYPE_UD);
4219 image = bld.emit_uniformize(image);
4220
4221 assert(nir_src_as_uint(instr->src[1]) == 0);
4222
4223 fs_reg srcs[TEX_LOGICAL_NUM_SRCS];
4224 if (instr->intrinsic == nir_intrinsic_image_size)
4225 srcs[TEX_LOGICAL_SRC_SURFACE] = image;
4226 else
4227 srcs[TEX_LOGICAL_SRC_SURFACE_HANDLE] = image;
4228 srcs[TEX_LOGICAL_SRC_SAMPLER] = brw_imm_d(0);
4229 srcs[TEX_LOGICAL_SRC_COORD_COMPONENTS] = brw_imm_d(0);
4230 srcs[TEX_LOGICAL_SRC_GRAD_COMPONENTS] = brw_imm_d(0);
4231
4232 /* Since the image size is always uniform, we can just emit a SIMD8
4233 * query instruction and splat the result out.
4234 */
4235 const fs_builder ubld = bld.exec_all().group(8, 0);
4236
4237 fs_reg tmp = ubld.vgrf(BRW_REGISTER_TYPE_UD, 4);
4238 fs_inst *inst = ubld.emit(SHADER_OPCODE_IMAGE_SIZE_LOGICAL,
4239 tmp, srcs, ARRAY_SIZE(srcs));
4240 inst->size_written = 4 * REG_SIZE;
4241
4242 for (unsigned c = 0; c < instr->dest.ssa.num_components; ++c) {
4243 if (c == 2 && nir_intrinsic_image_dim(instr) == GLSL_SAMPLER_DIM_CUBE) {
4244 bld.emit(SHADER_OPCODE_INT_QUOTIENT,
4245 offset(retype(dest, tmp.type), bld, c),
4246 component(offset(tmp, ubld, c), 0), brw_imm_ud(6));
4247 } else {
4248 bld.MOV(offset(retype(dest, tmp.type), bld, c),
4249 component(offset(tmp, ubld, c), 0));
4250 }
4251 }
4252 break;
4253 }
4254
4255 case nir_intrinsic_image_load_raw_intel: {
4256 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4257 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4258 get_nir_image_intrinsic_image(bld, instr);
4259 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
4260 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
4261 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
4262 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
4263
4264 fs_inst *inst =
4265 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_READ_LOGICAL,
4266 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
4267 inst->size_written = instr->num_components * dispatch_width * 4;
4268 break;
4269 }
4270
4271 case nir_intrinsic_image_store_raw_intel: {
4272 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4273 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4274 get_nir_image_intrinsic_image(bld, instr);
4275 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
4276 srcs[SURFACE_LOGICAL_SRC_DATA] = get_nir_src(instr->src[2]);
4277 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
4278 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
4279 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
4280
4281 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_WRITE_LOGICAL,
4282 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
4283 break;
4284 }
4285
4286 case nir_intrinsic_scoped_barrier:
4287 assert(nir_intrinsic_execution_scope(instr) == NIR_SCOPE_NONE);
4288 /* Fall through. */
4289 case nir_intrinsic_group_memory_barrier:
4290 case nir_intrinsic_memory_barrier_shared:
4291 case nir_intrinsic_memory_barrier_buffer:
4292 case nir_intrinsic_memory_barrier_image:
4293 case nir_intrinsic_memory_barrier:
4294 case nir_intrinsic_begin_invocation_interlock:
4295 case nir_intrinsic_end_invocation_interlock: {
4296 bool l3_fence, slm_fence;
4297 const enum opcode opcode =
4298 instr->intrinsic == nir_intrinsic_begin_invocation_interlock ?
4299 SHADER_OPCODE_INTERLOCK : SHADER_OPCODE_MEMORY_FENCE;
4300
4301 switch (instr->intrinsic) {
4302 case nir_intrinsic_scoped_barrier: {
4303 nir_variable_mode modes = nir_intrinsic_memory_modes(instr);
4304 l3_fence = modes & (nir_var_shader_out |
4305 nir_var_mem_ssbo |
4306 nir_var_mem_global);
4307 slm_fence = modes & nir_var_mem_shared;
4308 break;
4309 }
4310
4311 case nir_intrinsic_begin_invocation_interlock:
4312 case nir_intrinsic_end_invocation_interlock:
4313 /* For beginInvocationInterlockARB(), we will generate a memory fence
4314 * but with a different opcode so that generator can pick SENDC
4315 * instead of SEND.
4316 *
4317 * For endInvocationInterlockARB(), we need to insert a memory fence which
4318 * stalls in the shader until the memory transactions prior to that
4319 * fence are complete. This ensures that the shader does not end before
4320 * any writes from its critical section have landed. Otherwise, you can
4321 * end up with a case where the next invocation on that pixel properly
4322 * stalls for previous FS invocation on its pixel to complete but
4323 * doesn't actually wait for the dataport memory transactions from that
4324 * thread to land before submitting its own.
4325 *
4326 * Handling them here will allow the logic for IVB render cache (see
4327 * below) to be reused.
4328 */
4329 l3_fence = true;
4330 slm_fence = false;
4331 break;
4332
4333 default:
4334 l3_fence = instr->intrinsic != nir_intrinsic_memory_barrier_shared;
4335 slm_fence = instr->intrinsic == nir_intrinsic_group_memory_barrier ||
4336 instr->intrinsic == nir_intrinsic_memory_barrier ||
4337 instr->intrinsic == nir_intrinsic_memory_barrier_shared;
4338 break;
4339 }
4340
4341 if (stage != MESA_SHADER_COMPUTE && stage != MESA_SHADER_KERNEL)
4342 slm_fence = false;
4343
4344 /* If the workgroup fits in a single HW thread, the messages for SLM are
4345 * processed in-order and the shader itself is already synchronized so
4346 * the memory fence is not necessary.
4347 *
4348 * TODO: Check if applies for many HW threads sharing same Data Port.
4349 */
4350 if (!nir->info.cs.local_size_variable &&
4351 slm_fence && workgroup_size() <= dispatch_width)
4352 slm_fence = false;
4353
4354 /* Prior to Gen11, there's only L3 fence, so emit that instead. */
4355 if (slm_fence && devinfo->gen < 11) {
4356 slm_fence = false;
4357 l3_fence = true;
4358 }
4359
4360 /* IVB does typed surface access through the render cache, so we need
4361 * to flush it too.
4362 */
4363 const bool needs_render_fence =
4364 devinfo->gen == 7 && !devinfo->is_haswell;
4365
4366 /* Be conservative in Gen11+ and always stall in a fence. Since there
4367 * are two different fences, and shader might want to synchronize
4368 * between them.
4369 *
4370 * TODO: Use scope and visibility information for the barriers from NIR
4371 * to make a better decision on whether we need to stall.
4372 */
4373 const bool stall = devinfo->gen >= 11 || needs_render_fence ||
4374 instr->intrinsic == nir_intrinsic_end_invocation_interlock;
4375
4376 const bool commit_enable = stall ||
4377 devinfo->gen >= 10; /* HSD ES # 1404612949 */
4378
4379 unsigned fence_regs_count = 0;
4380 fs_reg fence_regs[2] = {};
4381
4382 const fs_builder ubld = bld.group(8, 0);
4383
4384 if (l3_fence) {
4385 fs_inst *fence =
4386 ubld.emit(opcode,
4387 ubld.vgrf(BRW_REGISTER_TYPE_UD),
4388 brw_vec8_grf(0, 0),
4389 brw_imm_ud(commit_enable),
4390 brw_imm_ud(/* bti */ 0));
4391 fence->sfid = GEN7_SFID_DATAPORT_DATA_CACHE;
4392
4393 fence_regs[fence_regs_count++] = fence->dst;
4394
4395 if (needs_render_fence) {
4396 fs_inst *render_fence =
4397 ubld.emit(opcode,
4398 ubld.vgrf(BRW_REGISTER_TYPE_UD),
4399 brw_vec8_grf(0, 0),
4400 brw_imm_ud(commit_enable),
4401 brw_imm_ud(/* bti */ 0));
4402 render_fence->sfid = GEN6_SFID_DATAPORT_RENDER_CACHE;
4403
4404 fence_regs[fence_regs_count++] = render_fence->dst;
4405 }
4406 }
4407
4408 if (slm_fence) {
4409 assert(opcode == SHADER_OPCODE_MEMORY_FENCE);
4410 fs_inst *fence =
4411 ubld.emit(opcode,
4412 ubld.vgrf(BRW_REGISTER_TYPE_UD),
4413 brw_vec8_grf(0, 0),
4414 brw_imm_ud(commit_enable),
4415 brw_imm_ud(GEN7_BTI_SLM));
4416 fence->sfid = GEN7_SFID_DATAPORT_DATA_CACHE;
4417
4418 fence_regs[fence_regs_count++] = fence->dst;
4419 }
4420
4421 assert(fence_regs_count <= 2);
4422
4423 if (stall || fence_regs_count == 0) {
4424 ubld.exec_all().group(1, 0).emit(
4425 FS_OPCODE_SCHEDULING_FENCE, ubld.null_reg_ud(),
4426 fence_regs, fence_regs_count);
4427 }
4428
4429 break;
4430 }
4431
4432 case nir_intrinsic_memory_barrier_tcs_patch:
4433 break;
4434
4435 case nir_intrinsic_shader_clock: {
4436 /* We cannot do anything if there is an event, so ignore it for now */
4437 const fs_reg shader_clock = get_timestamp(bld);
4438 const fs_reg srcs[] = { component(shader_clock, 0),
4439 component(shader_clock, 1) };
4440 bld.LOAD_PAYLOAD(dest, srcs, ARRAY_SIZE(srcs), 0);
4441 break;
4442 }
4443
4444 case nir_intrinsic_image_samples:
4445 /* The driver does not support multi-sampled images. */
4446 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D), brw_imm_d(1));
4447 break;
4448
4449 case nir_intrinsic_load_reloc_const_intel: {
4450 uint32_t id = nir_intrinsic_param_idx(instr);
4451 bld.emit(SHADER_OPCODE_MOV_RELOC_IMM,
4452 dest, brw_imm_ud(id));
4453 break;
4454 }
4455
4456 case nir_intrinsic_load_uniform: {
4457 /* Offsets are in bytes but they should always aligned to
4458 * the type size
4459 */
4460 assert(instr->const_index[0] % 4 == 0 ||
4461 instr->const_index[0] % type_sz(dest.type) == 0);
4462
4463 fs_reg src(UNIFORM, instr->const_index[0] / 4, dest.type);
4464
4465 if (nir_src_is_const(instr->src[0])) {
4466 unsigned load_offset = nir_src_as_uint(instr->src[0]);
4467 assert(load_offset % type_sz(dest.type) == 0);
4468 /* For 16-bit types we add the module of the const_index[0]
4469 * offset to access to not 32-bit aligned element
4470 */
4471 src.offset = load_offset + instr->const_index[0] % 4;
4472
4473 for (unsigned j = 0; j < instr->num_components; j++) {
4474 bld.MOV(offset(dest, bld, j), offset(src, bld, j));
4475 }
4476 } else {
4477 fs_reg indirect = retype(get_nir_src(instr->src[0]),
4478 BRW_REGISTER_TYPE_UD);
4479
4480 /* We need to pass a size to the MOV_INDIRECT but we don't want it to
4481 * go past the end of the uniform. In order to keep the n'th
4482 * component from running past, we subtract off the size of all but
4483 * one component of the vector.
4484 */
4485 assert(instr->const_index[1] >=
4486 instr->num_components * (int) type_sz(dest.type));
4487 unsigned read_size = instr->const_index[1] -
4488 (instr->num_components - 1) * type_sz(dest.type);
4489
4490 bool supports_64bit_indirects =
4491 !devinfo->is_cherryview && !gen_device_info_is_9lp(devinfo);
4492
4493 if (type_sz(dest.type) != 8 || supports_64bit_indirects) {
4494 for (unsigned j = 0; j < instr->num_components; j++) {
4495 bld.emit(SHADER_OPCODE_MOV_INDIRECT,
4496 offset(dest, bld, j), offset(src, bld, j),
4497 indirect, brw_imm_ud(read_size));
4498 }
4499 } else {
4500 const unsigned num_mov_indirects =
4501 type_sz(dest.type) / type_sz(BRW_REGISTER_TYPE_UD);
4502 /* We read a little bit less per MOV INDIRECT, as they are now
4503 * 32-bits ones instead of 64-bit. Fix read_size then.
4504 */
4505 const unsigned read_size_32bit = read_size -
4506 (num_mov_indirects - 1) * type_sz(BRW_REGISTER_TYPE_UD);
4507 for (unsigned j = 0; j < instr->num_components; j++) {
4508 for (unsigned i = 0; i < num_mov_indirects; i++) {
4509 bld.emit(SHADER_OPCODE_MOV_INDIRECT,
4510 subscript(offset(dest, bld, j), BRW_REGISTER_TYPE_UD, i),
4511 subscript(offset(src, bld, j), BRW_REGISTER_TYPE_UD, i),
4512 indirect, brw_imm_ud(read_size_32bit));
4513 }
4514 }
4515 }
4516 }
4517 break;
4518 }
4519
4520 case nir_intrinsic_load_ubo: {
4521 fs_reg surf_index;
4522 if (nir_src_is_const(instr->src[0])) {
4523 const unsigned index = stage_prog_data->binding_table.ubo_start +
4524 nir_src_as_uint(instr->src[0]);
4525 surf_index = brw_imm_ud(index);
4526 } else {
4527 /* The block index is not a constant. Evaluate the index expression
4528 * per-channel and add the base UBO index; we have to select a value
4529 * from any live channel.
4530 */
4531 surf_index = vgrf(glsl_type::uint_type);
4532 bld.ADD(surf_index, get_nir_src(instr->src[0]),
4533 brw_imm_ud(stage_prog_data->binding_table.ubo_start));
4534 surf_index = bld.emit_uniformize(surf_index);
4535 }
4536
4537 if (!nir_src_is_const(instr->src[1])) {
4538 fs_reg base_offset = retype(get_nir_src(instr->src[1]),
4539 BRW_REGISTER_TYPE_UD);
4540
4541 for (int i = 0; i < instr->num_components; i++)
4542 VARYING_PULL_CONSTANT_LOAD(bld, offset(dest, bld, i), surf_index,
4543 base_offset, i * type_sz(dest.type),
4544 nir_dest_bit_size(instr->dest) / 8);
4545
4546 prog_data->has_ubo_pull = true;
4547 } else {
4548 /* Even if we are loading doubles, a pull constant load will load
4549 * a 32-bit vec4, so should only reserve vgrf space for that. If we
4550 * need to load a full dvec4 we will have to emit 2 loads. This is
4551 * similar to demote_pull_constants(), except that in that case we
4552 * see individual accesses to each component of the vector and then
4553 * we let CSE deal with duplicate loads. Here we see a vector access
4554 * and we have to split it if necessary.
4555 */
4556 const unsigned type_size = type_sz(dest.type);
4557 const unsigned load_offset = nir_src_as_uint(instr->src[1]);
4558
4559 /* See if we've selected this as a push constant candidate */
4560 if (nir_src_is_const(instr->src[0])) {
4561 const unsigned ubo_block = nir_src_as_uint(instr->src[0]);
4562 const unsigned offset_256b = load_offset / 32;
4563
4564 fs_reg push_reg;
4565 for (int i = 0; i < 4; i++) {
4566 const struct brw_ubo_range *range = &prog_data->ubo_ranges[i];
4567 if (range->block == ubo_block &&
4568 offset_256b >= range->start &&
4569 offset_256b < range->start + range->length) {
4570
4571 push_reg = fs_reg(UNIFORM, UBO_START + i, dest.type);
4572 push_reg.offset = load_offset - 32 * range->start;
4573 break;
4574 }
4575 }
4576
4577 if (push_reg.file != BAD_FILE) {
4578 for (unsigned i = 0; i < instr->num_components; i++) {
4579 bld.MOV(offset(dest, bld, i),
4580 byte_offset(push_reg, i * type_size));
4581 }
4582 break;
4583 }
4584 }
4585
4586 prog_data->has_ubo_pull = true;
4587
4588 const unsigned block_sz = 64; /* Fetch one cacheline at a time. */
4589 const fs_builder ubld = bld.exec_all().group(block_sz / 4, 0);
4590 const fs_reg packed_consts = ubld.vgrf(BRW_REGISTER_TYPE_UD);
4591
4592 for (unsigned c = 0; c < instr->num_components;) {
4593 const unsigned base = load_offset + c * type_size;
4594 /* Number of usable components in the next block-aligned load. */
4595 const unsigned count = MIN2(instr->num_components - c,
4596 (block_sz - base % block_sz) / type_size);
4597
4598 ubld.emit(FS_OPCODE_UNIFORM_PULL_CONSTANT_LOAD,
4599 packed_consts, surf_index,
4600 brw_imm_ud(base & ~(block_sz - 1)));
4601
4602 const fs_reg consts =
4603 retype(byte_offset(packed_consts, base & (block_sz - 1)),
4604 dest.type);
4605
4606 for (unsigned d = 0; d < count; d++)
4607 bld.MOV(offset(dest, bld, c + d), component(consts, d));
4608
4609 c += count;
4610 }
4611 }
4612 break;
4613 }
4614
4615 case nir_intrinsic_load_global:
4616 case nir_intrinsic_load_global_constant: {
4617 assert(devinfo->gen >= 8);
4618
4619 assert(nir_dest_bit_size(instr->dest) <= 32);
4620 assert(nir_intrinsic_align(instr) > 0);
4621 if (nir_dest_bit_size(instr->dest) == 32 &&
4622 nir_intrinsic_align(instr) >= 4) {
4623 assert(nir_dest_num_components(instr->dest) <= 4);
4624 fs_inst *inst = bld.emit(SHADER_OPCODE_A64_UNTYPED_READ_LOGICAL,
4625 dest,
4626 get_nir_src(instr->src[0]), /* Address */
4627 fs_reg(), /* No source data */
4628 brw_imm_ud(instr->num_components));
4629 inst->size_written = instr->num_components *
4630 inst->dst.component_size(inst->exec_size);
4631 } else {
4632 const unsigned bit_size = nir_dest_bit_size(instr->dest);
4633 assert(nir_dest_num_components(instr->dest) == 1);
4634 fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_UD);
4635 bld.emit(SHADER_OPCODE_A64_BYTE_SCATTERED_READ_LOGICAL,
4636 tmp,
4637 get_nir_src(instr->src[0]), /* Address */
4638 fs_reg(), /* No source data */
4639 brw_imm_ud(bit_size));
4640 bld.MOV(dest, subscript(tmp, dest.type, 0));
4641 }
4642 break;
4643 }
4644
4645 case nir_intrinsic_store_global:
4646 assert(devinfo->gen >= 8);
4647
4648 assert(nir_src_bit_size(instr->src[0]) <= 32);
4649 assert(nir_intrinsic_write_mask(instr) ==
4650 (1u << instr->num_components) - 1);
4651 assert(nir_intrinsic_align(instr) > 0);
4652 if (nir_src_bit_size(instr->src[0]) == 32 &&
4653 nir_intrinsic_align(instr) >= 4) {
4654 assert(nir_src_num_components(instr->src[0]) <= 4);
4655 bld.emit(SHADER_OPCODE_A64_UNTYPED_WRITE_LOGICAL,
4656 fs_reg(),
4657 get_nir_src(instr->src[1]), /* Address */
4658 get_nir_src(instr->src[0]), /* Data */
4659 brw_imm_ud(instr->num_components));
4660 } else {
4661 assert(nir_src_num_components(instr->src[0]) == 1);
4662 const unsigned bit_size = nir_src_bit_size(instr->src[0]);
4663 brw_reg_type data_type =
4664 brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
4665 fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_UD);
4666 bld.MOV(tmp, retype(get_nir_src(instr->src[0]), data_type));
4667 bld.emit(SHADER_OPCODE_A64_BYTE_SCATTERED_WRITE_LOGICAL,
4668 fs_reg(),
4669 get_nir_src(instr->src[1]), /* Address */
4670 tmp, /* Data */
4671 brw_imm_ud(nir_src_bit_size(instr->src[0])));
4672 }
4673 break;
4674
4675 case nir_intrinsic_global_atomic_add:
4676 case nir_intrinsic_global_atomic_imin:
4677 case nir_intrinsic_global_atomic_umin:
4678 case nir_intrinsic_global_atomic_imax:
4679 case nir_intrinsic_global_atomic_umax:
4680 case nir_intrinsic_global_atomic_and:
4681 case nir_intrinsic_global_atomic_or:
4682 case nir_intrinsic_global_atomic_xor:
4683 case nir_intrinsic_global_atomic_exchange:
4684 case nir_intrinsic_global_atomic_comp_swap:
4685 nir_emit_global_atomic(bld, brw_aop_for_nir_intrinsic(instr), instr);
4686 break;
4687 case nir_intrinsic_global_atomic_fmin:
4688 case nir_intrinsic_global_atomic_fmax:
4689 case nir_intrinsic_global_atomic_fcomp_swap:
4690 nir_emit_global_atomic_float(bld, brw_aop_for_nir_intrinsic(instr), instr);
4691 break;
4692
4693 case nir_intrinsic_load_ssbo: {
4694 assert(devinfo->gen >= 7);
4695
4696 const unsigned bit_size = nir_dest_bit_size(instr->dest);
4697 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4698 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4699 get_nir_ssbo_intrinsic_index(bld, instr);
4700 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
4701 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
4702 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
4703
4704 /* Make dest unsigned because that's what the temporary will be */
4705 dest.type = brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
4706
4707 /* Read the vector */
4708 assert(nir_dest_bit_size(instr->dest) <= 32);
4709 assert(nir_intrinsic_align(instr) > 0);
4710 if (nir_dest_bit_size(instr->dest) == 32 &&
4711 nir_intrinsic_align(instr) >= 4) {
4712 assert(nir_dest_num_components(instr->dest) <= 4);
4713 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
4714 fs_inst *inst =
4715 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_READ_LOGICAL,
4716 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
4717 inst->size_written = instr->num_components * dispatch_width * 4;
4718 } else {
4719 assert(nir_dest_num_components(instr->dest) == 1);
4720 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(bit_size);
4721
4722 fs_reg read_result = bld.vgrf(BRW_REGISTER_TYPE_UD);
4723 bld.emit(SHADER_OPCODE_BYTE_SCATTERED_READ_LOGICAL,
4724 read_result, srcs, SURFACE_LOGICAL_NUM_SRCS);
4725 bld.MOV(dest, subscript(read_result, dest.type, 0));
4726 }
4727 break;
4728 }
4729
4730 case nir_intrinsic_store_ssbo: {
4731 assert(devinfo->gen >= 7);
4732
4733 const unsigned bit_size = nir_src_bit_size(instr->src[0]);
4734 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4735 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4736 get_nir_ssbo_intrinsic_index(bld, instr);
4737 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[2]);
4738 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
4739 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
4740
4741 fs_reg data = get_nir_src(instr->src[0]);
4742 data.type = brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
4743
4744 assert(nir_src_bit_size(instr->src[0]) <= 32);
4745 assert(nir_intrinsic_write_mask(instr) ==
4746 (1u << instr->num_components) - 1);
4747 assert(nir_intrinsic_align(instr) > 0);
4748 if (nir_src_bit_size(instr->src[0]) == 32 &&
4749 nir_intrinsic_align(instr) >= 4) {
4750 assert(nir_src_num_components(instr->src[0]) <= 4);
4751 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
4752 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(instr->num_components);
4753 bld.emit(SHADER_OPCODE_UNTYPED_SURFACE_WRITE_LOGICAL,
4754 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
4755 } else {
4756 assert(nir_src_num_components(instr->src[0]) == 1);
4757 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(bit_size);
4758
4759 srcs[SURFACE_LOGICAL_SRC_DATA] = bld.vgrf(BRW_REGISTER_TYPE_UD);
4760 bld.MOV(srcs[SURFACE_LOGICAL_SRC_DATA], data);
4761
4762 bld.emit(SHADER_OPCODE_BYTE_SCATTERED_WRITE_LOGICAL,
4763 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
4764 }
4765 break;
4766 }
4767
4768 case nir_intrinsic_store_output: {
4769 assert(nir_src_bit_size(instr->src[0]) == 32);
4770 fs_reg src = get_nir_src(instr->src[0]);
4771
4772 unsigned store_offset = nir_src_as_uint(instr->src[1]);
4773 unsigned num_components = instr->num_components;
4774 unsigned first_component = nir_intrinsic_component(instr);
4775
4776 fs_reg new_dest = retype(offset(outputs[instr->const_index[0]], bld,
4777 4 * store_offset), src.type);
4778 for (unsigned j = 0; j < num_components; j++) {
4779 bld.MOV(offset(new_dest, bld, j + first_component),
4780 offset(src, bld, j));
4781 }
4782 break;
4783 }
4784
4785 case nir_intrinsic_ssbo_atomic_add:
4786 case nir_intrinsic_ssbo_atomic_imin:
4787 case nir_intrinsic_ssbo_atomic_umin:
4788 case nir_intrinsic_ssbo_atomic_imax:
4789 case nir_intrinsic_ssbo_atomic_umax:
4790 case nir_intrinsic_ssbo_atomic_and:
4791 case nir_intrinsic_ssbo_atomic_or:
4792 case nir_intrinsic_ssbo_atomic_xor:
4793 case nir_intrinsic_ssbo_atomic_exchange:
4794 case nir_intrinsic_ssbo_atomic_comp_swap:
4795 nir_emit_ssbo_atomic(bld, brw_aop_for_nir_intrinsic(instr), instr);
4796 break;
4797 case nir_intrinsic_ssbo_atomic_fmin:
4798 case nir_intrinsic_ssbo_atomic_fmax:
4799 case nir_intrinsic_ssbo_atomic_fcomp_swap:
4800 nir_emit_ssbo_atomic_float(bld, brw_aop_for_nir_intrinsic(instr), instr);
4801 break;
4802
4803 case nir_intrinsic_get_ssbo_size: {
4804 assert(nir_src_num_components(instr->src[0]) == 1);
4805 unsigned ssbo_index = nir_src_is_const(instr->src[0]) ?
4806 nir_src_as_uint(instr->src[0]) : 0;
4807
4808 /* A resinfo's sampler message is used to get the buffer size. The
4809 * SIMD8's writeback message consists of four registers and SIMD16's
4810 * writeback message consists of 8 destination registers (two per each
4811 * component). Because we are only interested on the first channel of
4812 * the first returned component, where resinfo returns the buffer size
4813 * for SURFTYPE_BUFFER, we can just use the SIMD8 variant regardless of
4814 * the dispatch width.
4815 */
4816 const fs_builder ubld = bld.exec_all().group(8, 0);
4817 fs_reg src_payload = ubld.vgrf(BRW_REGISTER_TYPE_UD);
4818 fs_reg ret_payload = ubld.vgrf(BRW_REGISTER_TYPE_UD, 4);
4819
4820 /* Set LOD = 0 */
4821 ubld.MOV(src_payload, brw_imm_d(0));
4822
4823 const unsigned index = prog_data->binding_table.ssbo_start + ssbo_index;
4824 fs_inst *inst = ubld.emit(SHADER_OPCODE_GET_BUFFER_SIZE, ret_payload,
4825 src_payload, brw_imm_ud(index));
4826 inst->header_size = 0;
4827 inst->mlen = 1;
4828 inst->size_written = 4 * REG_SIZE;
4829
4830 /* SKL PRM, vol07, 3D Media GPGPU Engine, Bounds Checking and Faulting:
4831 *
4832 * "Out-of-bounds checking is always performed at a DWord granularity. If
4833 * any part of the DWord is out-of-bounds then the whole DWord is
4834 * considered out-of-bounds."
4835 *
4836 * This implies that types with size smaller than 4-bytes need to be
4837 * padded if they don't complete the last dword of the buffer. But as we
4838 * need to maintain the original size we need to reverse the padding
4839 * calculation to return the correct size to know the number of elements
4840 * of an unsized array. As we stored in the last two bits of the surface
4841 * size the needed padding for the buffer, we calculate here the
4842 * original buffer_size reversing the surface_size calculation:
4843 *
4844 * surface_size = isl_align(buffer_size, 4) +
4845 * (isl_align(buffer_size) - buffer_size)
4846 *
4847 * buffer_size = surface_size & ~3 - surface_size & 3
4848 */
4849
4850 fs_reg size_aligned4 = ubld.vgrf(BRW_REGISTER_TYPE_UD);
4851 fs_reg size_padding = ubld.vgrf(BRW_REGISTER_TYPE_UD);
4852 fs_reg buffer_size = ubld.vgrf(BRW_REGISTER_TYPE_UD);
4853
4854 ubld.AND(size_padding, ret_payload, brw_imm_ud(3));
4855 ubld.AND(size_aligned4, ret_payload, brw_imm_ud(~3));
4856 ubld.ADD(buffer_size, size_aligned4, negate(size_padding));
4857
4858 bld.MOV(retype(dest, ret_payload.type), component(buffer_size, 0));
4859 break;
4860 }
4861
4862 case nir_intrinsic_load_scratch: {
4863 assert(devinfo->gen >= 7);
4864
4865 assert(nir_dest_num_components(instr->dest) == 1);
4866 const unsigned bit_size = nir_dest_bit_size(instr->dest);
4867 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4868
4869 if (devinfo->gen >= 8) {
4870 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4871 brw_imm_ud(GEN8_BTI_STATELESS_NON_COHERENT);
4872 } else {
4873 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(BRW_BTI_STATELESS);
4874 }
4875
4876 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
4877 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(bit_size);
4878 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
4879 const fs_reg nir_addr = get_nir_src(instr->src[0]);
4880
4881 /* Make dest unsigned because that's what the temporary will be */
4882 dest.type = brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
4883
4884 /* Read the vector */
4885 assert(nir_dest_num_components(instr->dest) == 1);
4886 assert(nir_dest_bit_size(instr->dest) <= 32);
4887 assert(nir_intrinsic_align(instr) > 0);
4888 if (nir_dest_bit_size(instr->dest) >= 4 &&
4889 nir_intrinsic_align(instr) >= 4) {
4890 /* The offset for a DWORD scattered message is in dwords. */
4891 srcs[SURFACE_LOGICAL_SRC_ADDRESS] =
4892 swizzle_nir_scratch_addr(bld, nir_addr, true);
4893
4894 bld.emit(SHADER_OPCODE_DWORD_SCATTERED_READ_LOGICAL,
4895 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
4896 } else {
4897 srcs[SURFACE_LOGICAL_SRC_ADDRESS] =
4898 swizzle_nir_scratch_addr(bld, nir_addr, false);
4899
4900 fs_reg read_result = bld.vgrf(BRW_REGISTER_TYPE_UD);
4901 bld.emit(SHADER_OPCODE_BYTE_SCATTERED_READ_LOGICAL,
4902 read_result, srcs, SURFACE_LOGICAL_NUM_SRCS);
4903 bld.MOV(dest, read_result);
4904 }
4905 break;
4906 }
4907
4908 case nir_intrinsic_store_scratch: {
4909 assert(devinfo->gen >= 7);
4910
4911 assert(nir_src_num_components(instr->src[0]) == 1);
4912 const unsigned bit_size = nir_src_bit_size(instr->src[0]);
4913 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
4914
4915 if (devinfo->gen >= 8) {
4916 srcs[SURFACE_LOGICAL_SRC_SURFACE] =
4917 brw_imm_ud(GEN8_BTI_STATELESS_NON_COHERENT);
4918 } else {
4919 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(BRW_BTI_STATELESS);
4920 }
4921
4922 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
4923 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(bit_size);
4924 /**
4925 * While this instruction has side-effects, it should not be predicated
4926 * on sample mask, because otherwise fs helper invocations would
4927 * load undefined values from scratch memory. And scratch memory
4928 * load-stores are produced from operations without side-effects, thus
4929 * they should not have different behaviour in the helper invocations.
4930 */
4931 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(0);
4932 const fs_reg nir_addr = get_nir_src(instr->src[1]);
4933
4934 fs_reg data = get_nir_src(instr->src[0]);
4935 data.type = brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_UD);
4936
4937 assert(nir_src_num_components(instr->src[0]) == 1);
4938 assert(nir_src_bit_size(instr->src[0]) <= 32);
4939 assert(nir_intrinsic_write_mask(instr) == 1);
4940 assert(nir_intrinsic_align(instr) > 0);
4941 if (nir_src_bit_size(instr->src[0]) == 32 &&
4942 nir_intrinsic_align(instr) >= 4) {
4943 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
4944
4945 /* The offset for a DWORD scattered message is in dwords. */
4946 srcs[SURFACE_LOGICAL_SRC_ADDRESS] =
4947 swizzle_nir_scratch_addr(bld, nir_addr, true);
4948
4949 bld.emit(SHADER_OPCODE_DWORD_SCATTERED_WRITE_LOGICAL,
4950 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
4951 } else {
4952 srcs[SURFACE_LOGICAL_SRC_DATA] = bld.vgrf(BRW_REGISTER_TYPE_UD);
4953 bld.MOV(srcs[SURFACE_LOGICAL_SRC_DATA], data);
4954
4955 srcs[SURFACE_LOGICAL_SRC_ADDRESS] =
4956 swizzle_nir_scratch_addr(bld, nir_addr, false);
4957
4958 bld.emit(SHADER_OPCODE_BYTE_SCATTERED_WRITE_LOGICAL,
4959 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
4960 }
4961 break;
4962 }
4963
4964 case nir_intrinsic_load_subgroup_size:
4965 /* This should only happen for fragment shaders because every other case
4966 * is lowered in NIR so we can optimize on it.
4967 */
4968 assert(stage == MESA_SHADER_FRAGMENT);
4969 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D), brw_imm_d(dispatch_width));
4970 break;
4971
4972 case nir_intrinsic_load_subgroup_invocation:
4973 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D),
4974 nir_system_values[SYSTEM_VALUE_SUBGROUP_INVOCATION]);
4975 break;
4976
4977 case nir_intrinsic_load_subgroup_eq_mask:
4978 case nir_intrinsic_load_subgroup_ge_mask:
4979 case nir_intrinsic_load_subgroup_gt_mask:
4980 case nir_intrinsic_load_subgroup_le_mask:
4981 case nir_intrinsic_load_subgroup_lt_mask:
4982 unreachable("not reached");
4983
4984 case nir_intrinsic_vote_any: {
4985 const fs_builder ubld = bld.exec_all().group(1, 0);
4986
4987 /* The any/all predicates do not consider channel enables. To prevent
4988 * dead channels from affecting the result, we initialize the flag with
4989 * with the identity value for the logical operation.
4990 */
4991 if (dispatch_width == 32) {
4992 /* For SIMD32, we use a UD type so we fill both f0.0 and f0.1. */
4993 ubld.MOV(retype(brw_flag_reg(0, 0), BRW_REGISTER_TYPE_UD),
4994 brw_imm_ud(0));
4995 } else {
4996 ubld.MOV(brw_flag_reg(0, 0), brw_imm_uw(0));
4997 }
4998 bld.CMP(bld.null_reg_d(), get_nir_src(instr->src[0]), brw_imm_d(0), BRW_CONDITIONAL_NZ);
4999
5000 /* For some reason, the any/all predicates don't work properly with
5001 * SIMD32. In particular, it appears that a SEL with a QtrCtrl of 2H
5002 * doesn't read the correct subset of the flag register and you end up
5003 * getting garbage in the second half. Work around this by using a pair
5004 * of 1-wide MOVs and scattering the result.
5005 */
5006 fs_reg res1 = ubld.vgrf(BRW_REGISTER_TYPE_D);
5007 ubld.MOV(res1, brw_imm_d(0));
5008 set_predicate(dispatch_width == 8 ? BRW_PREDICATE_ALIGN1_ANY8H :
5009 dispatch_width == 16 ? BRW_PREDICATE_ALIGN1_ANY16H :
5010 BRW_PREDICATE_ALIGN1_ANY32H,
5011 ubld.MOV(res1, brw_imm_d(-1)));
5012
5013 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D), component(res1, 0));
5014 break;
5015 }
5016 case nir_intrinsic_vote_all: {
5017 const fs_builder ubld = bld.exec_all().group(1, 0);
5018
5019 /* The any/all predicates do not consider channel enables. To prevent
5020 * dead channels from affecting the result, we initialize the flag with
5021 * with the identity value for the logical operation.
5022 */
5023 if (dispatch_width == 32) {
5024 /* For SIMD32, we use a UD type so we fill both f0.0 and f0.1. */
5025 ubld.MOV(retype(brw_flag_reg(0, 0), BRW_REGISTER_TYPE_UD),
5026 brw_imm_ud(0xffffffff));
5027 } else {
5028 ubld.MOV(brw_flag_reg(0, 0), brw_imm_uw(0xffff));
5029 }
5030 bld.CMP(bld.null_reg_d(), get_nir_src(instr->src[0]), brw_imm_d(0), BRW_CONDITIONAL_NZ);
5031
5032 /* For some reason, the any/all predicates don't work properly with
5033 * SIMD32. In particular, it appears that a SEL with a QtrCtrl of 2H
5034 * doesn't read the correct subset of the flag register and you end up
5035 * getting garbage in the second half. Work around this by using a pair
5036 * of 1-wide MOVs and scattering the result.
5037 */
5038 fs_reg res1 = ubld.vgrf(BRW_REGISTER_TYPE_D);
5039 ubld.MOV(res1, brw_imm_d(0));
5040 set_predicate(dispatch_width == 8 ? BRW_PREDICATE_ALIGN1_ALL8H :
5041 dispatch_width == 16 ? BRW_PREDICATE_ALIGN1_ALL16H :
5042 BRW_PREDICATE_ALIGN1_ALL32H,
5043 ubld.MOV(res1, brw_imm_d(-1)));
5044
5045 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D), component(res1, 0));
5046 break;
5047 }
5048 case nir_intrinsic_vote_feq:
5049 case nir_intrinsic_vote_ieq: {
5050 fs_reg value = get_nir_src(instr->src[0]);
5051 if (instr->intrinsic == nir_intrinsic_vote_feq) {
5052 const unsigned bit_size = nir_src_bit_size(instr->src[0]);
5053 value.type = bit_size == 8 ? BRW_REGISTER_TYPE_B :
5054 brw_reg_type_from_bit_size(bit_size, BRW_REGISTER_TYPE_F);
5055 }
5056
5057 fs_reg uniformized = bld.emit_uniformize(value);
5058 const fs_builder ubld = bld.exec_all().group(1, 0);
5059
5060 /* The any/all predicates do not consider channel enables. To prevent
5061 * dead channels from affecting the result, we initialize the flag with
5062 * with the identity value for the logical operation.
5063 */
5064 if (dispatch_width == 32) {
5065 /* For SIMD32, we use a UD type so we fill both f0.0 and f0.1. */
5066 ubld.MOV(retype(brw_flag_reg(0, 0), BRW_REGISTER_TYPE_UD),
5067 brw_imm_ud(0xffffffff));
5068 } else {
5069 ubld.MOV(brw_flag_reg(0, 0), brw_imm_uw(0xffff));
5070 }
5071 bld.CMP(bld.null_reg_d(), value, uniformized, BRW_CONDITIONAL_Z);
5072
5073 /* For some reason, the any/all predicates don't work properly with
5074 * SIMD32. In particular, it appears that a SEL with a QtrCtrl of 2H
5075 * doesn't read the correct subset of the flag register and you end up
5076 * getting garbage in the second half. Work around this by using a pair
5077 * of 1-wide MOVs and scattering the result.
5078 */
5079 fs_reg res1 = ubld.vgrf(BRW_REGISTER_TYPE_D);
5080 ubld.MOV(res1, brw_imm_d(0));
5081 set_predicate(dispatch_width == 8 ? BRW_PREDICATE_ALIGN1_ALL8H :
5082 dispatch_width == 16 ? BRW_PREDICATE_ALIGN1_ALL16H :
5083 BRW_PREDICATE_ALIGN1_ALL32H,
5084 ubld.MOV(res1, brw_imm_d(-1)));
5085
5086 bld.MOV(retype(dest, BRW_REGISTER_TYPE_D), component(res1, 0));
5087 break;
5088 }
5089
5090 case nir_intrinsic_ballot: {
5091 const fs_reg value = retype(get_nir_src(instr->src[0]),
5092 BRW_REGISTER_TYPE_UD);
5093 struct brw_reg flag = brw_flag_reg(0, 0);
5094 /* FIXME: For SIMD32 programs, this causes us to stomp on f0.1 as well
5095 * as f0.0. This is a problem for fragment programs as we currently use
5096 * f0.1 for discards. Fortunately, we don't support SIMD32 fragment
5097 * programs yet so this isn't a problem. When we do, something will
5098 * have to change.
5099 */
5100 if (dispatch_width == 32)
5101 flag.type = BRW_REGISTER_TYPE_UD;
5102
5103 bld.exec_all().group(1, 0).MOV(flag, brw_imm_ud(0u));
5104 bld.CMP(bld.null_reg_ud(), value, brw_imm_ud(0u), BRW_CONDITIONAL_NZ);
5105
5106 if (instr->dest.ssa.bit_size > 32) {
5107 dest.type = BRW_REGISTER_TYPE_UQ;
5108 } else {
5109 dest.type = BRW_REGISTER_TYPE_UD;
5110 }
5111 bld.MOV(dest, flag);
5112 break;
5113 }
5114
5115 case nir_intrinsic_read_invocation: {
5116 const fs_reg value = get_nir_src(instr->src[0]);
5117 const fs_reg invocation = get_nir_src(instr->src[1]);
5118 fs_reg tmp = bld.vgrf(value.type);
5119
5120 bld.exec_all().emit(SHADER_OPCODE_BROADCAST, tmp, value,
5121 bld.emit_uniformize(invocation));
5122
5123 bld.MOV(retype(dest, value.type), fs_reg(component(tmp, 0)));
5124 break;
5125 }
5126
5127 case nir_intrinsic_read_first_invocation: {
5128 const fs_reg value = get_nir_src(instr->src[0]);
5129 bld.MOV(retype(dest, value.type), bld.emit_uniformize(value));
5130 break;
5131 }
5132
5133 case nir_intrinsic_shuffle: {
5134 const fs_reg value = get_nir_src(instr->src[0]);
5135 const fs_reg index = get_nir_src(instr->src[1]);
5136
5137 bld.emit(SHADER_OPCODE_SHUFFLE, retype(dest, value.type), value, index);
5138 break;
5139 }
5140
5141 case nir_intrinsic_first_invocation: {
5142 fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_UD);
5143 bld.exec_all().emit(SHADER_OPCODE_FIND_LIVE_CHANNEL, tmp);
5144 bld.MOV(retype(dest, BRW_REGISTER_TYPE_UD),
5145 fs_reg(component(tmp, 0)));
5146 break;
5147 }
5148
5149 case nir_intrinsic_quad_broadcast: {
5150 const fs_reg value = get_nir_src(instr->src[0]);
5151 const unsigned index = nir_src_as_uint(instr->src[1]);
5152
5153 bld.emit(SHADER_OPCODE_CLUSTER_BROADCAST, retype(dest, value.type),
5154 value, brw_imm_ud(index), brw_imm_ud(4));
5155 break;
5156 }
5157
5158 case nir_intrinsic_quad_swap_horizontal: {
5159 const fs_reg value = get_nir_src(instr->src[0]);
5160 const fs_reg tmp = bld.vgrf(value.type);
5161 if (devinfo->gen <= 7) {
5162 /* The hardware doesn't seem to support these crazy regions with
5163 * compressed instructions on gen7 and earlier so we fall back to
5164 * using quad swizzles. Fortunately, we don't support 64-bit
5165 * anything in Vulkan on gen7.
5166 */
5167 assert(nir_src_bit_size(instr->src[0]) == 32);
5168 const fs_builder ubld = bld.exec_all();
5169 ubld.emit(SHADER_OPCODE_QUAD_SWIZZLE, tmp, value,
5170 brw_imm_ud(BRW_SWIZZLE4(1,0,3,2)));
5171 bld.MOV(retype(dest, value.type), tmp);
5172 } else {
5173 const fs_builder ubld = bld.exec_all().group(dispatch_width / 2, 0);
5174
5175 const fs_reg src_left = horiz_stride(value, 2);
5176 const fs_reg src_right = horiz_stride(horiz_offset(value, 1), 2);
5177 const fs_reg tmp_left = horiz_stride(tmp, 2);
5178 const fs_reg tmp_right = horiz_stride(horiz_offset(tmp, 1), 2);
5179
5180 ubld.MOV(tmp_left, src_right);
5181 ubld.MOV(tmp_right, src_left);
5182
5183 }
5184 bld.MOV(retype(dest, value.type), tmp);
5185 break;
5186 }
5187
5188 case nir_intrinsic_quad_swap_vertical: {
5189 const fs_reg value = get_nir_src(instr->src[0]);
5190 if (nir_src_bit_size(instr->src[0]) == 32) {
5191 /* For 32-bit, we can use a SIMD4x2 instruction to do this easily */
5192 const fs_reg tmp = bld.vgrf(value.type);
5193 const fs_builder ubld = bld.exec_all();
5194 ubld.emit(SHADER_OPCODE_QUAD_SWIZZLE, tmp, value,
5195 brw_imm_ud(BRW_SWIZZLE4(2,3,0,1)));
5196 bld.MOV(retype(dest, value.type), tmp);
5197 } else {
5198 /* For larger data types, we have to either emit dispatch_width many
5199 * MOVs or else fall back to doing indirects.
5200 */
5201 fs_reg idx = bld.vgrf(BRW_REGISTER_TYPE_W);
5202 bld.XOR(idx, nir_system_values[SYSTEM_VALUE_SUBGROUP_INVOCATION],
5203 brw_imm_w(0x2));
5204 bld.emit(SHADER_OPCODE_SHUFFLE, retype(dest, value.type), value, idx);
5205 }
5206 break;
5207 }
5208
5209 case nir_intrinsic_quad_swap_diagonal: {
5210 const fs_reg value = get_nir_src(instr->src[0]);
5211 if (nir_src_bit_size(instr->src[0]) == 32) {
5212 /* For 32-bit, we can use a SIMD4x2 instruction to do this easily */
5213 const fs_reg tmp = bld.vgrf(value.type);
5214 const fs_builder ubld = bld.exec_all();
5215 ubld.emit(SHADER_OPCODE_QUAD_SWIZZLE, tmp, value,
5216 brw_imm_ud(BRW_SWIZZLE4(3,2,1,0)));
5217 bld.MOV(retype(dest, value.type), tmp);
5218 } else {
5219 /* For larger data types, we have to either emit dispatch_width many
5220 * MOVs or else fall back to doing indirects.
5221 */
5222 fs_reg idx = bld.vgrf(BRW_REGISTER_TYPE_W);
5223 bld.XOR(idx, nir_system_values[SYSTEM_VALUE_SUBGROUP_INVOCATION],
5224 brw_imm_w(0x3));
5225 bld.emit(SHADER_OPCODE_SHUFFLE, retype(dest, value.type), value, idx);
5226 }
5227 break;
5228 }
5229
5230 case nir_intrinsic_reduce: {
5231 fs_reg src = get_nir_src(instr->src[0]);
5232 nir_op redop = (nir_op)nir_intrinsic_reduction_op(instr);
5233 unsigned cluster_size = nir_intrinsic_cluster_size(instr);
5234 if (cluster_size == 0 || cluster_size > dispatch_width)
5235 cluster_size = dispatch_width;
5236
5237 /* Figure out the source type */
5238 src.type = brw_type_for_nir_type(devinfo,
5239 (nir_alu_type)(nir_op_infos[redop].input_types[0] |
5240 nir_src_bit_size(instr->src[0])));
5241
5242 fs_reg identity = brw_nir_reduction_op_identity(bld, redop, src.type);
5243 opcode brw_op = brw_op_for_nir_reduction_op(redop);
5244 brw_conditional_mod cond_mod = brw_cond_mod_for_nir_reduction_op(redop);
5245
5246 /* Set up a register for all of our scratching around and initialize it
5247 * to reduction operation's identity value.
5248 */
5249 fs_reg scan = bld.vgrf(src.type);
5250 bld.exec_all().emit(SHADER_OPCODE_SEL_EXEC, scan, src, identity);
5251
5252 bld.emit_scan(brw_op, scan, cluster_size, cond_mod);
5253
5254 dest.type = src.type;
5255 if (cluster_size * type_sz(src.type) >= REG_SIZE * 2) {
5256 /* In this case, CLUSTER_BROADCAST instruction isn't needed because
5257 * the distance between clusters is at least 2 GRFs. In this case,
5258 * we don't need the weird striding of the CLUSTER_BROADCAST
5259 * instruction and can just do regular MOVs.
5260 */
5261 assert((cluster_size * type_sz(src.type)) % (REG_SIZE * 2) == 0);
5262 const unsigned groups =
5263 (dispatch_width * type_sz(src.type)) / (REG_SIZE * 2);
5264 const unsigned group_size = dispatch_width / groups;
5265 for (unsigned i = 0; i < groups; i++) {
5266 const unsigned cluster = (i * group_size) / cluster_size;
5267 const unsigned comp = cluster * cluster_size + (cluster_size - 1);
5268 bld.group(group_size, i).MOV(horiz_offset(dest, i * group_size),
5269 component(scan, comp));
5270 }
5271 } else {
5272 bld.emit(SHADER_OPCODE_CLUSTER_BROADCAST, dest, scan,
5273 brw_imm_ud(cluster_size - 1), brw_imm_ud(cluster_size));
5274 }
5275 break;
5276 }
5277
5278 case nir_intrinsic_inclusive_scan:
5279 case nir_intrinsic_exclusive_scan: {
5280 fs_reg src = get_nir_src(instr->src[0]);
5281 nir_op redop = (nir_op)nir_intrinsic_reduction_op(instr);
5282
5283 /* Figure out the source type */
5284 src.type = brw_type_for_nir_type(devinfo,
5285 (nir_alu_type)(nir_op_infos[redop].input_types[0] |
5286 nir_src_bit_size(instr->src[0])));
5287
5288 fs_reg identity = brw_nir_reduction_op_identity(bld, redop, src.type);
5289 opcode brw_op = brw_op_for_nir_reduction_op(redop);
5290 brw_conditional_mod cond_mod = brw_cond_mod_for_nir_reduction_op(redop);
5291
5292 /* Set up a register for all of our scratching around and initialize it
5293 * to reduction operation's identity value.
5294 */
5295 fs_reg scan = bld.vgrf(src.type);
5296 const fs_builder allbld = bld.exec_all();
5297 allbld.emit(SHADER_OPCODE_SEL_EXEC, scan, src, identity);
5298
5299 if (instr->intrinsic == nir_intrinsic_exclusive_scan) {
5300 /* Exclusive scan is a bit harder because we have to do an annoying
5301 * shift of the contents before we can begin. To make things worse,
5302 * we can't do this with a normal stride; we have to use indirects.
5303 */
5304 fs_reg shifted = bld.vgrf(src.type);
5305 fs_reg idx = bld.vgrf(BRW_REGISTER_TYPE_W);
5306 allbld.ADD(idx, nir_system_values[SYSTEM_VALUE_SUBGROUP_INVOCATION],
5307 brw_imm_w(-1));
5308 allbld.emit(SHADER_OPCODE_SHUFFLE, shifted, scan, idx);
5309 allbld.group(1, 0).MOV(component(shifted, 0), identity);
5310 scan = shifted;
5311 }
5312
5313 bld.emit_scan(brw_op, scan, dispatch_width, cond_mod);
5314
5315 bld.MOV(retype(dest, src.type), scan);
5316 break;
5317 }
5318
5319 case nir_intrinsic_load_global_block_intel: {
5320 assert(nir_dest_bit_size(instr->dest) == 32);
5321
5322 fs_reg address = bld.emit_uniformize(get_nir_src(instr->src[0]));
5323
5324 const fs_builder ubld1 = bld.exec_all().group(1, 0);
5325 const fs_builder ubld8 = bld.exec_all().group(8, 0);
5326 const fs_builder ubld16 = bld.exec_all().group(16, 0);
5327
5328 const unsigned total = instr->num_components * dispatch_width;
5329 unsigned loaded = 0;
5330
5331 while (loaded < total) {
5332 const unsigned block =
5333 choose_oword_block_size_dwords(total - loaded);
5334 const unsigned block_bytes = block * 4;
5335
5336 const fs_builder &ubld = block == 8 ? ubld8 : ubld16;
5337 ubld.emit(SHADER_OPCODE_A64_UNALIGNED_OWORD_BLOCK_READ_LOGICAL,
5338 retype(byte_offset(dest, loaded * 4), BRW_REGISTER_TYPE_UD),
5339 address,
5340 fs_reg(), /* No source data */
5341 brw_imm_ud(block))->size_written = block_bytes;
5342
5343 increment_a64_address(ubld1, address, block_bytes);
5344 loaded += block;
5345 }
5346
5347 assert(loaded == total);
5348 break;
5349 }
5350
5351 case nir_intrinsic_store_global_block_intel: {
5352 assert(nir_src_bit_size(instr->src[0]) == 32);
5353
5354 fs_reg address = bld.emit_uniformize(get_nir_src(instr->src[1]));
5355 fs_reg src = get_nir_src(instr->src[0]);
5356
5357 const fs_builder ubld1 = bld.exec_all().group(1, 0);
5358 const fs_builder ubld8 = bld.exec_all().group(8, 0);
5359 const fs_builder ubld16 = bld.exec_all().group(16, 0);
5360
5361 const unsigned total = instr->num_components * dispatch_width;
5362 unsigned written = 0;
5363
5364 while (written < total) {
5365 const unsigned block =
5366 choose_oword_block_size_dwords(total - written);
5367
5368 const fs_builder &ubld = block == 8 ? ubld8 : ubld16;
5369 ubld.emit(SHADER_OPCODE_A64_OWORD_BLOCK_WRITE_LOGICAL,
5370 fs_reg(),
5371 address,
5372 retype(byte_offset(src, written * 4), BRW_REGISTER_TYPE_UD),
5373 brw_imm_ud(block));
5374
5375 const unsigned block_bytes = block * 4;
5376 increment_a64_address(ubld1, address, block_bytes);
5377 written += block;
5378 }
5379
5380 assert(written == total);
5381 break;
5382 }
5383
5384 case nir_intrinsic_load_shared_block_intel:
5385 case nir_intrinsic_load_ssbo_block_intel: {
5386 assert(nir_dest_bit_size(instr->dest) == 32);
5387
5388 const bool is_ssbo =
5389 instr->intrinsic == nir_intrinsic_load_ssbo_block_intel;
5390 fs_reg address = bld.emit_uniformize(get_nir_src(instr->src[is_ssbo ? 1 : 0]));
5391
5392 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
5393 srcs[SURFACE_LOGICAL_SRC_SURFACE] = is_ssbo ?
5394 get_nir_ssbo_intrinsic_index(bld, instr) : fs_reg(brw_imm_ud(GEN7_BTI_SLM));
5395 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = address;
5396
5397 const fs_builder ubld1 = bld.exec_all().group(1, 0);
5398 const fs_builder ubld8 = bld.exec_all().group(8, 0);
5399 const fs_builder ubld16 = bld.exec_all().group(16, 0);
5400
5401 const unsigned total = instr->num_components * dispatch_width;
5402 unsigned loaded = 0;
5403
5404 while (loaded < total) {
5405 const unsigned block =
5406 choose_oword_block_size_dwords(total - loaded);
5407 const unsigned block_bytes = block * 4;
5408
5409 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(block);
5410
5411 const fs_builder &ubld = block == 8 ? ubld8 : ubld16;
5412 ubld.emit(SHADER_OPCODE_UNALIGNED_OWORD_BLOCK_READ_LOGICAL,
5413 retype(byte_offset(dest, loaded * 4), BRW_REGISTER_TYPE_UD),
5414 srcs, SURFACE_LOGICAL_NUM_SRCS)->size_written = block_bytes;
5415
5416 ubld1.ADD(address, address, brw_imm_ud(block_bytes));
5417 loaded += block;
5418 }
5419
5420 assert(loaded == total);
5421 break;
5422 }
5423
5424 case nir_intrinsic_store_shared_block_intel:
5425 case nir_intrinsic_store_ssbo_block_intel: {
5426 assert(nir_src_bit_size(instr->src[0]) == 32);
5427
5428 const bool is_ssbo =
5429 instr->intrinsic == nir_intrinsic_store_ssbo_block_intel;
5430
5431 fs_reg address = bld.emit_uniformize(get_nir_src(instr->src[is_ssbo ? 2 : 1]));
5432 fs_reg src = get_nir_src(instr->src[0]);
5433
5434 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
5435 srcs[SURFACE_LOGICAL_SRC_SURFACE] = is_ssbo ?
5436 get_nir_ssbo_intrinsic_index(bld, instr) : fs_reg(brw_imm_ud(GEN7_BTI_SLM));
5437 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = address;
5438
5439 const fs_builder ubld1 = bld.exec_all().group(1, 0);
5440 const fs_builder ubld8 = bld.exec_all().group(8, 0);
5441 const fs_builder ubld16 = bld.exec_all().group(16, 0);
5442
5443 const unsigned total = instr->num_components * dispatch_width;
5444 unsigned written = 0;
5445
5446 while (written < total) {
5447 const unsigned block =
5448 choose_oword_block_size_dwords(total - written);
5449
5450 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(block);
5451 srcs[SURFACE_LOGICAL_SRC_DATA] =
5452 retype(byte_offset(src, written * 4), BRW_REGISTER_TYPE_UD);
5453
5454 const fs_builder &ubld = block == 8 ? ubld8 : ubld16;
5455 ubld.emit(SHADER_OPCODE_OWORD_BLOCK_WRITE_LOGICAL,
5456 fs_reg(), srcs, SURFACE_LOGICAL_NUM_SRCS);
5457
5458 const unsigned block_bytes = block * 4;
5459 ubld1.ADD(address, address, brw_imm_ud(block_bytes));
5460 written += block;
5461 }
5462
5463 assert(written == total);
5464 break;
5465 }
5466
5467 default:
5468 unreachable("unknown intrinsic");
5469 }
5470 }
5471
5472 void
nir_emit_ssbo_atomic(const fs_builder & bld,int op,nir_intrinsic_instr * instr)5473 fs_visitor::nir_emit_ssbo_atomic(const fs_builder &bld,
5474 int op, nir_intrinsic_instr *instr)
5475 {
5476 /* The BTI untyped atomic messages only support 32-bit atomics. If you
5477 * just look at the big table of messages in the Vol 7 of the SKL PRM, they
5478 * appear to exist. However, if you look at Vol 2a, there are no message
5479 * descriptors provided for Qword atomic ops except for A64 messages.
5480 */
5481 assert(nir_dest_bit_size(instr->dest) == 32);
5482
5483 fs_reg dest;
5484 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
5485 dest = get_nir_dest(instr->dest);
5486
5487 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
5488 srcs[SURFACE_LOGICAL_SRC_SURFACE] = get_nir_ssbo_intrinsic_index(bld, instr);
5489 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
5490 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
5491 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(op);
5492 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
5493
5494 fs_reg data;
5495 if (op != BRW_AOP_INC && op != BRW_AOP_DEC && op != BRW_AOP_PREDEC)
5496 data = get_nir_src(instr->src[2]);
5497
5498 if (op == BRW_AOP_CMPWR) {
5499 fs_reg tmp = bld.vgrf(data.type, 2);
5500 fs_reg sources[2] = { data, get_nir_src(instr->src[3]) };
5501 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
5502 data = tmp;
5503 }
5504 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
5505
5506 /* Emit the actual atomic operation */
5507
5508 bld.emit(SHADER_OPCODE_UNTYPED_ATOMIC_LOGICAL,
5509 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
5510 }
5511
5512 void
nir_emit_ssbo_atomic_float(const fs_builder & bld,int op,nir_intrinsic_instr * instr)5513 fs_visitor::nir_emit_ssbo_atomic_float(const fs_builder &bld,
5514 int op, nir_intrinsic_instr *instr)
5515 {
5516 fs_reg dest;
5517 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
5518 dest = get_nir_dest(instr->dest);
5519
5520 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
5521 srcs[SURFACE_LOGICAL_SRC_SURFACE] = get_nir_ssbo_intrinsic_index(bld, instr);
5522 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = get_nir_src(instr->src[1]);
5523 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
5524 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(op);
5525 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
5526
5527 fs_reg data = get_nir_src(instr->src[2]);
5528 if (op == BRW_AOP_FCMPWR) {
5529 fs_reg tmp = bld.vgrf(data.type, 2);
5530 fs_reg sources[2] = { data, get_nir_src(instr->src[3]) };
5531 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
5532 data = tmp;
5533 }
5534 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
5535
5536 /* Emit the actual atomic operation */
5537
5538 bld.emit(SHADER_OPCODE_UNTYPED_ATOMIC_FLOAT_LOGICAL,
5539 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
5540 }
5541
5542 void
nir_emit_shared_atomic(const fs_builder & bld,int op,nir_intrinsic_instr * instr)5543 fs_visitor::nir_emit_shared_atomic(const fs_builder &bld,
5544 int op, nir_intrinsic_instr *instr)
5545 {
5546 fs_reg dest;
5547 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
5548 dest = get_nir_dest(instr->dest);
5549
5550 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
5551 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(GEN7_BTI_SLM);
5552 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
5553 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(op);
5554 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
5555
5556 fs_reg data;
5557 if (op != BRW_AOP_INC && op != BRW_AOP_DEC && op != BRW_AOP_PREDEC)
5558 data = get_nir_src(instr->src[1]);
5559 if (op == BRW_AOP_CMPWR) {
5560 fs_reg tmp = bld.vgrf(data.type, 2);
5561 fs_reg sources[2] = { data, get_nir_src(instr->src[2]) };
5562 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
5563 data = tmp;
5564 }
5565 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
5566
5567 /* Get the offset */
5568 if (nir_src_is_const(instr->src[0])) {
5569 srcs[SURFACE_LOGICAL_SRC_ADDRESS] =
5570 brw_imm_ud(instr->const_index[0] + nir_src_as_uint(instr->src[0]));
5571 } else {
5572 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = vgrf(glsl_type::uint_type);
5573 bld.ADD(srcs[SURFACE_LOGICAL_SRC_ADDRESS],
5574 retype(get_nir_src(instr->src[0]), BRW_REGISTER_TYPE_UD),
5575 brw_imm_ud(instr->const_index[0]));
5576 }
5577
5578 /* Emit the actual atomic operation operation */
5579
5580 bld.emit(SHADER_OPCODE_UNTYPED_ATOMIC_LOGICAL,
5581 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
5582 }
5583
5584 void
nir_emit_shared_atomic_float(const fs_builder & bld,int op,nir_intrinsic_instr * instr)5585 fs_visitor::nir_emit_shared_atomic_float(const fs_builder &bld,
5586 int op, nir_intrinsic_instr *instr)
5587 {
5588 fs_reg dest;
5589 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
5590 dest = get_nir_dest(instr->dest);
5591
5592 fs_reg srcs[SURFACE_LOGICAL_NUM_SRCS];
5593 srcs[SURFACE_LOGICAL_SRC_SURFACE] = brw_imm_ud(GEN7_BTI_SLM);
5594 srcs[SURFACE_LOGICAL_SRC_IMM_DIMS] = brw_imm_ud(1);
5595 srcs[SURFACE_LOGICAL_SRC_IMM_ARG] = brw_imm_ud(op);
5596 srcs[SURFACE_LOGICAL_SRC_ALLOW_SAMPLE_MASK] = brw_imm_ud(1);
5597
5598 fs_reg data = get_nir_src(instr->src[1]);
5599 if (op == BRW_AOP_FCMPWR) {
5600 fs_reg tmp = bld.vgrf(data.type, 2);
5601 fs_reg sources[2] = { data, get_nir_src(instr->src[2]) };
5602 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
5603 data = tmp;
5604 }
5605 srcs[SURFACE_LOGICAL_SRC_DATA] = data;
5606
5607 /* Get the offset */
5608 if (nir_src_is_const(instr->src[0])) {
5609 srcs[SURFACE_LOGICAL_SRC_ADDRESS] =
5610 brw_imm_ud(instr->const_index[0] + nir_src_as_uint(instr->src[0]));
5611 } else {
5612 srcs[SURFACE_LOGICAL_SRC_ADDRESS] = vgrf(glsl_type::uint_type);
5613 bld.ADD(srcs[SURFACE_LOGICAL_SRC_ADDRESS],
5614 retype(get_nir_src(instr->src[0]), BRW_REGISTER_TYPE_UD),
5615 brw_imm_ud(instr->const_index[0]));
5616 }
5617
5618 /* Emit the actual atomic operation operation */
5619
5620 bld.emit(SHADER_OPCODE_UNTYPED_ATOMIC_FLOAT_LOGICAL,
5621 dest, srcs, SURFACE_LOGICAL_NUM_SRCS);
5622 }
5623
5624 void
nir_emit_global_atomic(const fs_builder & bld,int op,nir_intrinsic_instr * instr)5625 fs_visitor::nir_emit_global_atomic(const fs_builder &bld,
5626 int op, nir_intrinsic_instr *instr)
5627 {
5628 fs_reg dest;
5629 if (nir_intrinsic_infos[instr->intrinsic].has_dest)
5630 dest = get_nir_dest(instr->dest);
5631
5632 fs_reg addr = get_nir_src(instr->src[0]);
5633
5634 fs_reg data;
5635 if (op != BRW_AOP_INC && op != BRW_AOP_DEC && op != BRW_AOP_PREDEC)
5636 data = get_nir_src(instr->src[1]);
5637
5638 if (op == BRW_AOP_CMPWR) {
5639 fs_reg tmp = bld.vgrf(data.type, 2);
5640 fs_reg sources[2] = { data, get_nir_src(instr->src[2]) };
5641 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
5642 data = tmp;
5643 }
5644
5645 if (nir_dest_bit_size(instr->dest) == 64) {
5646 bld.emit(SHADER_OPCODE_A64_UNTYPED_ATOMIC_INT64_LOGICAL,
5647 dest, addr, data, brw_imm_ud(op));
5648 } else {
5649 assert(nir_dest_bit_size(instr->dest) == 32);
5650 bld.emit(SHADER_OPCODE_A64_UNTYPED_ATOMIC_LOGICAL,
5651 dest, addr, data, brw_imm_ud(op));
5652 }
5653 }
5654
5655 void
nir_emit_global_atomic_float(const fs_builder & bld,int op,nir_intrinsic_instr * instr)5656 fs_visitor::nir_emit_global_atomic_float(const fs_builder &bld,
5657 int op, nir_intrinsic_instr *instr)
5658 {
5659 assert(nir_intrinsic_infos[instr->intrinsic].has_dest);
5660 fs_reg dest = get_nir_dest(instr->dest);
5661
5662 fs_reg addr = get_nir_src(instr->src[0]);
5663
5664 assert(op != BRW_AOP_INC && op != BRW_AOP_DEC && op != BRW_AOP_PREDEC);
5665 fs_reg data = get_nir_src(instr->src[1]);
5666
5667 if (op == BRW_AOP_FCMPWR) {
5668 fs_reg tmp = bld.vgrf(data.type, 2);
5669 fs_reg sources[2] = { data, get_nir_src(instr->src[2]) };
5670 bld.LOAD_PAYLOAD(tmp, sources, 2, 0);
5671 data = tmp;
5672 }
5673
5674 bld.emit(SHADER_OPCODE_A64_UNTYPED_ATOMIC_FLOAT_LOGICAL,
5675 dest, addr, data, brw_imm_ud(op));
5676 }
5677
5678 void
nir_emit_texture(const fs_builder & bld,nir_tex_instr * instr)5679 fs_visitor::nir_emit_texture(const fs_builder &bld, nir_tex_instr *instr)
5680 {
5681 unsigned texture = instr->texture_index;
5682 unsigned sampler = instr->sampler_index;
5683
5684 fs_reg srcs[TEX_LOGICAL_NUM_SRCS];
5685
5686 srcs[TEX_LOGICAL_SRC_SURFACE] = brw_imm_ud(texture);
5687 srcs[TEX_LOGICAL_SRC_SAMPLER] = brw_imm_ud(sampler);
5688
5689 int lod_components = 0;
5690
5691 /* The hardware requires a LOD for buffer textures */
5692 if (instr->sampler_dim == GLSL_SAMPLER_DIM_BUF)
5693 srcs[TEX_LOGICAL_SRC_LOD] = brw_imm_d(0);
5694
5695 uint32_t header_bits = 0;
5696 for (unsigned i = 0; i < instr->num_srcs; i++) {
5697 fs_reg src = get_nir_src(instr->src[i].src);
5698 switch (instr->src[i].src_type) {
5699 case nir_tex_src_bias:
5700 srcs[TEX_LOGICAL_SRC_LOD] =
5701 retype(get_nir_src_imm(instr->src[i].src), BRW_REGISTER_TYPE_F);
5702 break;
5703 case nir_tex_src_comparator:
5704 srcs[TEX_LOGICAL_SRC_SHADOW_C] = retype(src, BRW_REGISTER_TYPE_F);
5705 break;
5706 case nir_tex_src_coord:
5707 switch (instr->op) {
5708 case nir_texop_txf:
5709 case nir_texop_txf_ms:
5710 case nir_texop_txf_ms_mcs:
5711 case nir_texop_samples_identical:
5712 srcs[TEX_LOGICAL_SRC_COORDINATE] = retype(src, BRW_REGISTER_TYPE_D);
5713 break;
5714 default:
5715 srcs[TEX_LOGICAL_SRC_COORDINATE] = retype(src, BRW_REGISTER_TYPE_F);
5716 break;
5717 }
5718 break;
5719 case nir_tex_src_ddx:
5720 srcs[TEX_LOGICAL_SRC_LOD] = retype(src, BRW_REGISTER_TYPE_F);
5721 lod_components = nir_tex_instr_src_size(instr, i);
5722 break;
5723 case nir_tex_src_ddy:
5724 srcs[TEX_LOGICAL_SRC_LOD2] = retype(src, BRW_REGISTER_TYPE_F);
5725 break;
5726 case nir_tex_src_lod:
5727 switch (instr->op) {
5728 case nir_texop_txs:
5729 srcs[TEX_LOGICAL_SRC_LOD] =
5730 retype(get_nir_src_imm(instr->src[i].src), BRW_REGISTER_TYPE_UD);
5731 break;
5732 case nir_texop_txf:
5733 srcs[TEX_LOGICAL_SRC_LOD] =
5734 retype(get_nir_src_imm(instr->src[i].src), BRW_REGISTER_TYPE_D);
5735 break;
5736 default:
5737 srcs[TEX_LOGICAL_SRC_LOD] =
5738 retype(get_nir_src_imm(instr->src[i].src), BRW_REGISTER_TYPE_F);
5739 break;
5740 }
5741 break;
5742 case nir_tex_src_min_lod:
5743 srcs[TEX_LOGICAL_SRC_MIN_LOD] =
5744 retype(get_nir_src_imm(instr->src[i].src), BRW_REGISTER_TYPE_F);
5745 break;
5746 case nir_tex_src_ms_index:
5747 srcs[TEX_LOGICAL_SRC_SAMPLE_INDEX] = retype(src, BRW_REGISTER_TYPE_UD);
5748 break;
5749
5750 case nir_tex_src_offset: {
5751 uint32_t offset_bits = 0;
5752 if (brw_texture_offset(instr, i, &offset_bits)) {
5753 header_bits |= offset_bits;
5754 } else {
5755 srcs[TEX_LOGICAL_SRC_TG4_OFFSET] =
5756 retype(src, BRW_REGISTER_TYPE_D);
5757 }
5758 break;
5759 }
5760
5761 case nir_tex_src_projector:
5762 unreachable("should be lowered");
5763
5764 case nir_tex_src_texture_offset: {
5765 /* Emit code to evaluate the actual indexing expression */
5766 fs_reg tmp = vgrf(glsl_type::uint_type);
5767 bld.ADD(tmp, src, brw_imm_ud(texture));
5768 srcs[TEX_LOGICAL_SRC_SURFACE] = bld.emit_uniformize(tmp);
5769 break;
5770 }
5771
5772 case nir_tex_src_sampler_offset: {
5773 /* Emit code to evaluate the actual indexing expression */
5774 fs_reg tmp = vgrf(glsl_type::uint_type);
5775 bld.ADD(tmp, src, brw_imm_ud(sampler));
5776 srcs[TEX_LOGICAL_SRC_SAMPLER] = bld.emit_uniformize(tmp);
5777 break;
5778 }
5779
5780 case nir_tex_src_texture_handle:
5781 assert(nir_tex_instr_src_index(instr, nir_tex_src_texture_offset) == -1);
5782 srcs[TEX_LOGICAL_SRC_SURFACE] = fs_reg();
5783 srcs[TEX_LOGICAL_SRC_SURFACE_HANDLE] = bld.emit_uniformize(src);
5784 break;
5785
5786 case nir_tex_src_sampler_handle:
5787 assert(nir_tex_instr_src_index(instr, nir_tex_src_sampler_offset) == -1);
5788 srcs[TEX_LOGICAL_SRC_SAMPLER] = fs_reg();
5789 srcs[TEX_LOGICAL_SRC_SAMPLER_HANDLE] = bld.emit_uniformize(src);
5790 break;
5791
5792 case nir_tex_src_ms_mcs:
5793 assert(instr->op == nir_texop_txf_ms);
5794 srcs[TEX_LOGICAL_SRC_MCS] = retype(src, BRW_REGISTER_TYPE_D);
5795 break;
5796
5797 case nir_tex_src_plane: {
5798 const uint32_t plane = nir_src_as_uint(instr->src[i].src);
5799 const uint32_t texture_index =
5800 instr->texture_index +
5801 stage_prog_data->binding_table.plane_start[plane] -
5802 stage_prog_data->binding_table.texture_start;
5803
5804 srcs[TEX_LOGICAL_SRC_SURFACE] = brw_imm_ud(texture_index);
5805 break;
5806 }
5807
5808 default:
5809 unreachable("unknown texture source");
5810 }
5811 }
5812
5813 if (srcs[TEX_LOGICAL_SRC_MCS].file == BAD_FILE &&
5814 (instr->op == nir_texop_txf_ms ||
5815 instr->op == nir_texop_samples_identical)) {
5816 if (devinfo->gen >= 7 &&
5817 key_tex->compressed_multisample_layout_mask & (1 << texture)) {
5818 srcs[TEX_LOGICAL_SRC_MCS] =
5819 emit_mcs_fetch(srcs[TEX_LOGICAL_SRC_COORDINATE],
5820 instr->coord_components,
5821 srcs[TEX_LOGICAL_SRC_SURFACE],
5822 srcs[TEX_LOGICAL_SRC_SURFACE_HANDLE]);
5823 } else {
5824 srcs[TEX_LOGICAL_SRC_MCS] = brw_imm_ud(0u);
5825 }
5826 }
5827
5828 srcs[TEX_LOGICAL_SRC_COORD_COMPONENTS] = brw_imm_d(instr->coord_components);
5829 srcs[TEX_LOGICAL_SRC_GRAD_COMPONENTS] = brw_imm_d(lod_components);
5830
5831 enum opcode opcode;
5832 switch (instr->op) {
5833 case nir_texop_tex:
5834 opcode = SHADER_OPCODE_TEX_LOGICAL;
5835 break;
5836 case nir_texop_txb:
5837 opcode = FS_OPCODE_TXB_LOGICAL;
5838 break;
5839 case nir_texop_txl:
5840 opcode = SHADER_OPCODE_TXL_LOGICAL;
5841 break;
5842 case nir_texop_txd:
5843 opcode = SHADER_OPCODE_TXD_LOGICAL;
5844 break;
5845 case nir_texop_txf:
5846 opcode = SHADER_OPCODE_TXF_LOGICAL;
5847 break;
5848 case nir_texop_txf_ms:
5849 if ((key_tex->msaa_16 & (1 << sampler)))
5850 opcode = SHADER_OPCODE_TXF_CMS_W_LOGICAL;
5851 else
5852 opcode = SHADER_OPCODE_TXF_CMS_LOGICAL;
5853 break;
5854 case nir_texop_txf_ms_mcs:
5855 opcode = SHADER_OPCODE_TXF_MCS_LOGICAL;
5856 break;
5857 case nir_texop_query_levels:
5858 case nir_texop_txs:
5859 opcode = SHADER_OPCODE_TXS_LOGICAL;
5860 break;
5861 case nir_texop_lod:
5862 opcode = SHADER_OPCODE_LOD_LOGICAL;
5863 break;
5864 case nir_texop_tg4:
5865 if (srcs[TEX_LOGICAL_SRC_TG4_OFFSET].file != BAD_FILE)
5866 opcode = SHADER_OPCODE_TG4_OFFSET_LOGICAL;
5867 else
5868 opcode = SHADER_OPCODE_TG4_LOGICAL;
5869 break;
5870 case nir_texop_texture_samples:
5871 opcode = SHADER_OPCODE_SAMPLEINFO_LOGICAL;
5872 break;
5873 case nir_texop_samples_identical: {
5874 fs_reg dst = retype(get_nir_dest(instr->dest), BRW_REGISTER_TYPE_D);
5875
5876 /* If mcs is an immediate value, it means there is no MCS. In that case
5877 * just return false.
5878 */
5879 if (srcs[TEX_LOGICAL_SRC_MCS].file == BRW_IMMEDIATE_VALUE) {
5880 bld.MOV(dst, brw_imm_ud(0u));
5881 } else if ((key_tex->msaa_16 & (1 << sampler))) {
5882 fs_reg tmp = vgrf(glsl_type::uint_type);
5883 bld.OR(tmp, srcs[TEX_LOGICAL_SRC_MCS],
5884 offset(srcs[TEX_LOGICAL_SRC_MCS], bld, 1));
5885 bld.CMP(dst, tmp, brw_imm_ud(0u), BRW_CONDITIONAL_EQ);
5886 } else {
5887 bld.CMP(dst, srcs[TEX_LOGICAL_SRC_MCS], brw_imm_ud(0u),
5888 BRW_CONDITIONAL_EQ);
5889 }
5890 return;
5891 }
5892 default:
5893 unreachable("unknown texture opcode");
5894 }
5895
5896 if (instr->op == nir_texop_tg4) {
5897 if (instr->component == 1 &&
5898 key_tex->gather_channel_quirk_mask & (1 << texture)) {
5899 /* gather4 sampler is broken for green channel on RG32F --
5900 * we must ask for blue instead.
5901 */
5902 header_bits |= 2 << 16;
5903 } else {
5904 header_bits |= instr->component << 16;
5905 }
5906 }
5907
5908 fs_reg dst = bld.vgrf(brw_type_for_nir_type(devinfo, instr->dest_type), 4);
5909 fs_inst *inst = bld.emit(opcode, dst, srcs, ARRAY_SIZE(srcs));
5910 inst->offset = header_bits;
5911
5912 const unsigned dest_size = nir_tex_instr_dest_size(instr);
5913 if (devinfo->gen >= 9 &&
5914 instr->op != nir_texop_tg4 && instr->op != nir_texop_query_levels) {
5915 unsigned write_mask = instr->dest.is_ssa ?
5916 nir_ssa_def_components_read(&instr->dest.ssa):
5917 (1 << dest_size) - 1;
5918 assert(write_mask != 0); /* dead code should have been eliminated */
5919 inst->size_written = util_last_bit(write_mask) *
5920 inst->dst.component_size(inst->exec_size);
5921 } else {
5922 inst->size_written = 4 * inst->dst.component_size(inst->exec_size);
5923 }
5924
5925 if (srcs[TEX_LOGICAL_SRC_SHADOW_C].file != BAD_FILE)
5926 inst->shadow_compare = true;
5927
5928 if (instr->op == nir_texop_tg4 && devinfo->gen == 6)
5929 emit_gen6_gather_wa(key_tex->gen6_gather_wa[texture], dst);
5930
5931 fs_reg nir_dest[4];
5932 for (unsigned i = 0; i < dest_size; i++)
5933 nir_dest[i] = offset(dst, bld, i);
5934
5935 if (instr->op == nir_texop_query_levels) {
5936 /* # levels is in .w */
5937 nir_dest[0] = offset(dst, bld, 3);
5938 } else if (instr->op == nir_texop_txs &&
5939 dest_size >= 3 && devinfo->gen < 7) {
5940 /* Gen4-6 return 0 instead of 1 for single layer surfaces. */
5941 fs_reg depth = offset(dst, bld, 2);
5942 nir_dest[2] = vgrf(glsl_type::int_type);
5943 bld.emit_minmax(nir_dest[2], depth, brw_imm_d(1), BRW_CONDITIONAL_GE);
5944 }
5945
5946 bld.LOAD_PAYLOAD(get_nir_dest(instr->dest), nir_dest, dest_size, 0);
5947 }
5948
5949 void
nir_emit_jump(const fs_builder & bld,nir_jump_instr * instr)5950 fs_visitor::nir_emit_jump(const fs_builder &bld, nir_jump_instr *instr)
5951 {
5952 switch (instr->type) {
5953 case nir_jump_break:
5954 bld.emit(BRW_OPCODE_BREAK);
5955 break;
5956 case nir_jump_continue:
5957 bld.emit(BRW_OPCODE_CONTINUE);
5958 break;
5959 case nir_jump_return:
5960 default:
5961 unreachable("unknown jump");
5962 }
5963 }
5964
5965 /*
5966 * This helper takes a source register and un/shuffles it into the destination
5967 * register.
5968 *
5969 * If source type size is smaller than destination type size the operation
5970 * needed is a component shuffle. The opposite case would be an unshuffle. If
5971 * source/destination type size is equal a shuffle is done that would be
5972 * equivalent to a simple MOV.
5973 *
5974 * For example, if source is a 16-bit type and destination is 32-bit. A 3
5975 * components .xyz 16-bit vector on SIMD8 would be.
5976 *
5977 * |x1|x2|x3|x4|x5|x6|x7|x8|y1|y2|y3|y4|y5|y6|y7|y8|
5978 * |z1|z2|z3|z4|z5|z6|z7|z8| | | | | | | | |
5979 *
5980 * This helper will return the following 2 32-bit components with the 16-bit
5981 * values shuffled:
5982 *
5983 * |x1 y1|x2 y2|x3 y3|x4 y4|x5 y5|x6 y6|x7 y7|x8 y8|
5984 * |z1 |z2 |z3 |z4 |z5 |z6 |z7 |z8 |
5985 *
5986 * For unshuffle, the example would be the opposite, a 64-bit type source
5987 * and a 32-bit destination. A 2 component .xy 64-bit vector on SIMD8
5988 * would be:
5989 *
5990 * | x1l x1h | x2l x2h | x3l x3h | x4l x4h |
5991 * | x5l x5h | x6l x6h | x7l x7h | x8l x8h |
5992 * | y1l y1h | y2l y2h | y3l y3h | y4l y4h |
5993 * | y5l y5h | y6l y6h | y7l y7h | y8l y8h |
5994 *
5995 * The returned result would be the following 4 32-bit components unshuffled:
5996 *
5997 * | x1l | x2l | x3l | x4l | x5l | x6l | x7l | x8l |
5998 * | x1h | x2h | x3h | x4h | x5h | x6h | x7h | x8h |
5999 * | y1l | y2l | y3l | y4l | y5l | y6l | y7l | y8l |
6000 * | y1h | y2h | y3h | y4h | y5h | y6h | y7h | y8h |
6001 *
6002 * - Source and destination register must not be overlapped.
6003 * - components units are measured in terms of the smaller type between
6004 * source and destination because we are un/shuffling the smaller
6005 * components from/into the bigger ones.
6006 * - first_component parameter allows skipping source components.
6007 */
6008 void
shuffle_src_to_dst(const fs_builder & bld,const fs_reg & dst,const fs_reg & src,uint32_t first_component,uint32_t components)6009 shuffle_src_to_dst(const fs_builder &bld,
6010 const fs_reg &dst,
6011 const fs_reg &src,
6012 uint32_t first_component,
6013 uint32_t components)
6014 {
6015 if (type_sz(src.type) == type_sz(dst.type)) {
6016 assert(!regions_overlap(dst,
6017 type_sz(dst.type) * bld.dispatch_width() * components,
6018 offset(src, bld, first_component),
6019 type_sz(src.type) * bld.dispatch_width() * components));
6020 for (unsigned i = 0; i < components; i++) {
6021 bld.MOV(retype(offset(dst, bld, i), src.type),
6022 offset(src, bld, i + first_component));
6023 }
6024 } else if (type_sz(src.type) < type_sz(dst.type)) {
6025 /* Source is shuffled into destination */
6026 unsigned size_ratio = type_sz(dst.type) / type_sz(src.type);
6027 assert(!regions_overlap(dst,
6028 type_sz(dst.type) * bld.dispatch_width() *
6029 DIV_ROUND_UP(components, size_ratio),
6030 offset(src, bld, first_component),
6031 type_sz(src.type) * bld.dispatch_width() * components));
6032
6033 brw_reg_type shuffle_type =
6034 brw_reg_type_from_bit_size(8 * type_sz(src.type),
6035 BRW_REGISTER_TYPE_D);
6036 for (unsigned i = 0; i < components; i++) {
6037 fs_reg shuffle_component_i =
6038 subscript(offset(dst, bld, i / size_ratio),
6039 shuffle_type, i % size_ratio);
6040 bld.MOV(shuffle_component_i,
6041 retype(offset(src, bld, i + first_component), shuffle_type));
6042 }
6043 } else {
6044 /* Source is unshuffled into destination */
6045 unsigned size_ratio = type_sz(src.type) / type_sz(dst.type);
6046 assert(!regions_overlap(dst,
6047 type_sz(dst.type) * bld.dispatch_width() * components,
6048 offset(src, bld, first_component / size_ratio),
6049 type_sz(src.type) * bld.dispatch_width() *
6050 DIV_ROUND_UP(components + (first_component % size_ratio),
6051 size_ratio)));
6052
6053 brw_reg_type shuffle_type =
6054 brw_reg_type_from_bit_size(8 * type_sz(dst.type),
6055 BRW_REGISTER_TYPE_D);
6056 for (unsigned i = 0; i < components; i++) {
6057 fs_reg shuffle_component_i =
6058 subscript(offset(src, bld, (first_component + i) / size_ratio),
6059 shuffle_type, (first_component + i) % size_ratio);
6060 bld.MOV(retype(offset(dst, bld, i), shuffle_type),
6061 shuffle_component_i);
6062 }
6063 }
6064 }
6065
6066 void
shuffle_from_32bit_read(const fs_builder & bld,const fs_reg & dst,const fs_reg & src,uint32_t first_component,uint32_t components)6067 shuffle_from_32bit_read(const fs_builder &bld,
6068 const fs_reg &dst,
6069 const fs_reg &src,
6070 uint32_t first_component,
6071 uint32_t components)
6072 {
6073 assert(type_sz(src.type) == 4);
6074
6075 /* This function takes components in units of the destination type while
6076 * shuffle_src_to_dst takes components in units of the smallest type
6077 */
6078 if (type_sz(dst.type) > 4) {
6079 assert(type_sz(dst.type) == 8);
6080 first_component *= 2;
6081 components *= 2;
6082 }
6083
6084 shuffle_src_to_dst(bld, dst, src, first_component, components);
6085 }
6086
6087 fs_reg
setup_imm_df(const fs_builder & bld,double v)6088 setup_imm_df(const fs_builder &bld, double v)
6089 {
6090 const struct gen_device_info *devinfo = bld.shader->devinfo;
6091 assert(devinfo->gen >= 7);
6092
6093 if (devinfo->gen >= 8)
6094 return brw_imm_df(v);
6095
6096 /* gen7.5 does not support DF immediates straighforward but the DIM
6097 * instruction allows to set the 64-bit immediate value.
6098 */
6099 if (devinfo->is_haswell) {
6100 const fs_builder ubld = bld.exec_all().group(1, 0);
6101 fs_reg dst = ubld.vgrf(BRW_REGISTER_TYPE_DF, 1);
6102 ubld.DIM(dst, brw_imm_df(v));
6103 return component(dst, 0);
6104 }
6105
6106 /* gen7 does not support DF immediates, so we generate a 64-bit constant by
6107 * writing the low 32-bit of the constant to suboffset 0 of a VGRF and
6108 * the high 32-bit to suboffset 4 and then applying a stride of 0.
6109 *
6110 * Alternatively, we could also produce a normal VGRF (without stride 0)
6111 * by writing to all the channels in the VGRF, however, that would hit the
6112 * gen7 bug where we have to split writes that span more than 1 register
6113 * into instructions with a width of 4 (otherwise the write to the second
6114 * register written runs into an execmask hardware bug) which isn't very
6115 * nice.
6116 */
6117 union {
6118 double d;
6119 struct {
6120 uint32_t i1;
6121 uint32_t i2;
6122 };
6123 } di;
6124
6125 di.d = v;
6126
6127 const fs_builder ubld = bld.exec_all().group(1, 0);
6128 const fs_reg tmp = ubld.vgrf(BRW_REGISTER_TYPE_UD, 2);
6129 ubld.MOV(tmp, brw_imm_ud(di.i1));
6130 ubld.MOV(horiz_offset(tmp, 1), brw_imm_ud(di.i2));
6131
6132 return component(retype(tmp, BRW_REGISTER_TYPE_DF), 0);
6133 }
6134
6135 fs_reg
setup_imm_b(const fs_builder & bld,int8_t v)6136 setup_imm_b(const fs_builder &bld, int8_t v)
6137 {
6138 const fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_B);
6139 bld.MOV(tmp, brw_imm_w(v));
6140 return tmp;
6141 }
6142
6143 fs_reg
setup_imm_ub(const fs_builder & bld,uint8_t v)6144 setup_imm_ub(const fs_builder &bld, uint8_t v)
6145 {
6146 const fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_UB);
6147 bld.MOV(tmp, brw_imm_uw(v));
6148 return tmp;
6149 }
6150