1/* 2 * Copyright © 2017 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 shall be included 12 * in all copies or substantial portions of the Software. 13 * 14 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS 15 * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 16 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL 17 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 18 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING 19 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER 20 * DEALINGS IN THE SOFTWARE. 21 */ 22 23/** 24 * @file crocus_program.c 25 * 26 * This file contains the driver interface for compiling shaders. 27 * 28 * See crocus_program_cache.c for the in-memory program cache where the 29 * compiled shaders are stored. 30 */ 31 32#include <stdio.h> 33#include <errno.h> 34#include "pipe/p_defines.h" 35#include "pipe/p_state.h" 36#include "pipe/p_context.h" 37#include "pipe/p_screen.h" 38#include "util/u_atomic.h" 39#include "util/u_upload_mgr.h" 40#include "util/debug.h" 41#include "util/u_prim.h" 42#include "compiler/nir/nir.h" 43#include "compiler/nir/nir_builder.h" 44#include "compiler/nir/nir_serialize.h" 45#include "intel/compiler/brw_compiler.h" 46#include "intel/compiler/brw_nir.h" 47#include "crocus_context.h" 48#include "nir/tgsi_to_nir.h" 49 50#define KEY_INIT_NO_ID() \ 51 .base.subgroup_size_type = BRW_SUBGROUP_SIZE_UNIFORM, \ 52 .base.tex.swizzles[0 ... MAX_SAMPLERS - 1] = 0x688, \ 53 .base.tex.compressed_multisample_layout_mask = ~0 54#define KEY_INIT() .base.program_string_id = ish->program_id, KEY_INIT_NO_ID() 55 56static void 57crocus_sanitize_tex_key(struct brw_sampler_prog_key_data *key) 58{ 59 key->gather_channel_quirk_mask = 0; 60 for (unsigned s = 0; s < MAX_SAMPLERS; s++) { 61 key->swizzles[s] = SWIZZLE_NOOP; 62 key->gfx6_gather_wa[s] = 0; 63 } 64} 65 66static uint32_t 67crocus_get_texture_swizzle(const struct crocus_context *ice, 68 const struct crocus_sampler_view *t) 69{ 70 uint32_t swiz = 0; 71 72 for (int i = 0; i < 4; i++) { 73 swiz |= t->swizzle[i] << (i * 3); 74 } 75 return swiz; 76} 77 78static inline bool can_push_ubo(const struct intel_device_info *devinfo) 79{ 80 /* push works for everyone except SNB at the moment */ 81 return devinfo->ver != 6; 82} 83 84static uint8_t 85gfx6_gather_workaround(enum pipe_format pformat) 86{ 87 switch (pformat) { 88 case PIPE_FORMAT_R8_SINT: return WA_SIGN | WA_8BIT; 89 case PIPE_FORMAT_R8_UINT: return WA_8BIT; 90 case PIPE_FORMAT_R16_SINT: return WA_SIGN | WA_16BIT; 91 case PIPE_FORMAT_R16_UINT: return WA_16BIT; 92 default: 93 /* Note that even though PIPE_FORMAT_R32_SINT and 94 * PIPE_FORMAT_R32_UINThave format overrides in 95 * the surface state, there is no shader w/a required. 96 */ 97 return 0; 98 } 99} 100 101static const unsigned crocus_gfx6_swizzle_for_offset[4] = { 102 BRW_SWIZZLE4(0, 1, 2, 3), 103 BRW_SWIZZLE4(1, 2, 3, 3), 104 BRW_SWIZZLE4(2, 3, 3, 3), 105 BRW_SWIZZLE4(3, 3, 3, 3) 106}; 107 108static void 109gfx6_gs_xfb_setup(const struct pipe_stream_output_info *so_info, 110 struct brw_gs_prog_data *gs_prog_data) 111{ 112 /* Make sure that the VUE slots won't overflow the unsigned chars in 113 * prog_data->transform_feedback_bindings[]. 114 */ 115 STATIC_ASSERT(BRW_VARYING_SLOT_COUNT <= 256); 116 117 /* Make sure that we don't need more binding table entries than we've 118 * set aside for use in transform feedback. (We shouldn't, since we 119 * set aside enough binding table entries to have one per component). 120 */ 121 assert(so_info->num_outputs <= BRW_MAX_SOL_BINDINGS); 122 123 gs_prog_data->num_transform_feedback_bindings = so_info->num_outputs; 124 for (unsigned i = 0; i < so_info->num_outputs; i++) { 125 gs_prog_data->transform_feedback_bindings[i] = 126 so_info->output[i].register_index; 127 gs_prog_data->transform_feedback_swizzles[i] = 128 crocus_gfx6_swizzle_for_offset[so_info->output[i].start_component]; 129 } 130} 131 132static void 133gfx6_ff_gs_xfb_setup(const struct pipe_stream_output_info *so_info, 134 struct brw_ff_gs_prog_key *key) 135{ 136 key->num_transform_feedback_bindings = so_info->num_outputs; 137 for (unsigned i = 0; i < so_info->num_outputs; i++) { 138 key->transform_feedback_bindings[i] = 139 so_info->output[i].register_index; 140 key->transform_feedback_swizzles[i] = 141 crocus_gfx6_swizzle_for_offset[so_info->output[i].start_component]; 142 } 143} 144 145static void 146crocus_populate_sampler_prog_key_data(struct crocus_context *ice, 147 const struct intel_device_info *devinfo, 148 gl_shader_stage stage, 149 struct crocus_uncompiled_shader *ish, 150 bool uses_texture_gather, 151 struct brw_sampler_prog_key_data *key) 152{ 153 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 154 uint32_t mask = ish->nir->info.textures_used[0]; 155 156 while (mask) { 157 const int s = u_bit_scan(&mask); 158 159 struct crocus_sampler_view *texture = ice->state.shaders[stage].textures[s]; 160 key->swizzles[s] = SWIZZLE_NOOP; 161 key->scale_factors[s] = 0.0f; 162 163 if (!texture) 164 continue; 165 if (texture->base.target == PIPE_BUFFER) 166 continue; 167 if (devinfo->verx10 < 75) { 168 key->swizzles[s] = crocus_get_texture_swizzle(ice, texture); 169 } 170 171 screen->vtbl.fill_clamp_mask(ice->state.shaders[stage].samplers[s], s, key->gl_clamp_mask); 172 173 /* gather4 for RG32* is broken in multiple ways on Gen7. */ 174 if (devinfo->ver == 7 && uses_texture_gather) { 175 switch (texture->base.format) { 176 case PIPE_FORMAT_R32G32_UINT: 177 case PIPE_FORMAT_R32G32_SINT: { 178 /* We have to override the format to R32G32_FLOAT_LD. 179 * This means that SCS_ALPHA and SCS_ONE will return 0x3f8 180 * (1.0) rather than integer 1. This needs shader hacks. 181 * 182 * On Ivybridge, we whack W (alpha) to ONE in our key's 183 * swizzle. On Haswell, we look at the original texture 184 * swizzle, and use XYZW with channels overridden to ONE, 185 * leaving normal texture swizzling to SCS. 186 */ 187 unsigned src_swizzle = key->swizzles[s]; 188 for (int i = 0; i < 4; i++) { 189 unsigned src_comp = GET_SWZ(src_swizzle, i); 190 if (src_comp == SWIZZLE_ONE || src_comp == SWIZZLE_W) { 191 key->swizzles[i] &= ~(0x7 << (3 * i)); 192 key->swizzles[i] |= SWIZZLE_ONE << (3 * i); 193 } 194 } 195 } 196 FALLTHROUGH; 197 case PIPE_FORMAT_R32G32_FLOAT: 198 /* The channel select for green doesn't work - we have to 199 * request blue. Haswell can use SCS for this, but Ivybridge 200 * needs a shader workaround. 201 */ 202 if (devinfo->verx10 < 75) 203 key->gather_channel_quirk_mask |= 1 << s; 204 break; 205 default: 206 break; 207 } 208 } 209 if (devinfo->ver == 6 && uses_texture_gather) { 210 key->gfx6_gather_wa[s] = gfx6_gather_workaround(texture->base.format); 211 } 212 } 213} 214 215static void 216crocus_lower_swizzles(struct nir_shader *nir, 217 const struct brw_sampler_prog_key_data *key_tex) 218{ 219 struct nir_lower_tex_options tex_options = { 0 }; 220 uint32_t mask = nir->info.textures_used[0]; 221 222 while (mask) { 223 const int s = u_bit_scan(&mask); 224 225 if (key_tex->swizzles[s] == SWIZZLE_NOOP) 226 continue; 227 228 tex_options.swizzle_result |= (1 << s); 229 for (unsigned c = 0; c < 4; c++) 230 tex_options.swizzles[s][c] = GET_SWZ(key_tex->swizzles[s], c); 231 } 232 if (tex_options.swizzle_result) 233 nir_lower_tex(nir, &tex_options); 234} 235 236static unsigned 237get_new_program_id(struct crocus_screen *screen) 238{ 239 return p_atomic_inc_return(&screen->program_id); 240} 241 242static nir_ssa_def * 243get_aoa_deref_offset(nir_builder *b, 244 nir_deref_instr *deref, 245 unsigned elem_size) 246{ 247 unsigned array_size = elem_size; 248 nir_ssa_def *offset = nir_imm_int(b, 0); 249 250 while (deref->deref_type != nir_deref_type_var) { 251 assert(deref->deref_type == nir_deref_type_array); 252 253 /* This level's element size is the previous level's array size */ 254 nir_ssa_def *index = nir_ssa_for_src(b, deref->arr.index, 1); 255 assert(deref->arr.index.ssa); 256 offset = nir_iadd(b, offset, 257 nir_imul(b, index, nir_imm_int(b, array_size))); 258 259 deref = nir_deref_instr_parent(deref); 260 assert(glsl_type_is_array(deref->type)); 261 array_size *= glsl_get_length(deref->type); 262 } 263 264 /* Accessing an invalid surface index with the dataport can result in a 265 * hang. According to the spec "if the index used to select an individual 266 * element is negative or greater than or equal to the size of the array, 267 * the results of the operation are undefined but may not lead to 268 * termination" -- which is one of the possible outcomes of the hang. 269 * Clamp the index to prevent access outside of the array bounds. 270 */ 271 return nir_umin(b, offset, nir_imm_int(b, array_size - elem_size)); 272} 273 274static void 275crocus_lower_storage_image_derefs(nir_shader *nir) 276{ 277 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 278 279 nir_builder b; 280 nir_builder_init(&b, impl); 281 282 nir_foreach_block(block, impl) { 283 nir_foreach_instr_safe(instr, block) { 284 if (instr->type != nir_instr_type_intrinsic) 285 continue; 286 287 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr); 288 switch (intrin->intrinsic) { 289 case nir_intrinsic_image_deref_load: 290 case nir_intrinsic_image_deref_store: 291 case nir_intrinsic_image_deref_atomic_add: 292 case nir_intrinsic_image_deref_atomic_imin: 293 case nir_intrinsic_image_deref_atomic_umin: 294 case nir_intrinsic_image_deref_atomic_imax: 295 case nir_intrinsic_image_deref_atomic_umax: 296 case nir_intrinsic_image_deref_atomic_and: 297 case nir_intrinsic_image_deref_atomic_or: 298 case nir_intrinsic_image_deref_atomic_xor: 299 case nir_intrinsic_image_deref_atomic_exchange: 300 case nir_intrinsic_image_deref_atomic_comp_swap: 301 case nir_intrinsic_image_deref_size: 302 case nir_intrinsic_image_deref_samples: 303 case nir_intrinsic_image_deref_load_raw_intel: 304 case nir_intrinsic_image_deref_store_raw_intel: { 305 nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); 306 nir_variable *var = nir_deref_instr_get_variable(deref); 307 308 b.cursor = nir_before_instr(&intrin->instr); 309 nir_ssa_def *index = 310 nir_iadd(&b, nir_imm_int(&b, var->data.driver_location), 311 get_aoa_deref_offset(&b, deref, 1)); 312 nir_rewrite_image_intrinsic(intrin, index, false); 313 break; 314 } 315 316 default: 317 break; 318 } 319 } 320 } 321} 322 323// XXX: need unify_interfaces() at link time... 324 325/** 326 * Undo nir_lower_passthrough_edgeflags but keep the inputs_read flag. 327 */ 328static bool 329crocus_fix_edge_flags(nir_shader *nir) 330{ 331 if (nir->info.stage != MESA_SHADER_VERTEX) { 332 nir_shader_preserve_all_metadata(nir); 333 return false; 334 } 335 336 nir_variable *var = nir_find_variable_with_location(nir, nir_var_shader_out, 337 VARYING_SLOT_EDGE); 338 if (!var) { 339 nir_shader_preserve_all_metadata(nir); 340 return false; 341 } 342 343 var->data.mode = nir_var_shader_temp; 344 nir->info.outputs_written &= ~VARYING_BIT_EDGE; 345 nir->info.inputs_read &= ~VERT_BIT_EDGEFLAG; 346 nir_fixup_deref_modes(nir); 347 348 nir_foreach_function(f, nir) { 349 if (f->impl) { 350 nir_metadata_preserve(f->impl, nir_metadata_block_index | 351 nir_metadata_dominance | 352 nir_metadata_live_ssa_defs | 353 nir_metadata_loop_analysis); 354 } else { 355 nir_metadata_preserve(f->impl, nir_metadata_all); 356 } 357 } 358 359 return true; 360} 361 362/** 363 * Fix an uncompiled shader's stream output info. 364 * 365 * Core Gallium stores output->register_index as a "slot" number, where 366 * slots are assigned consecutively to all outputs in info->outputs_written. 367 * This naive packing of outputs doesn't work for us - we too have slots, 368 * but the layout is defined by the VUE map, which we won't have until we 369 * compile a specific shader variant. So, we remap these and simply store 370 * VARYING_SLOT_* in our copy's output->register_index fields. 371 * 372 * We also fix up VARYING_SLOT_{LAYER,VIEWPORT,PSIZ} to select the Y/Z/W 373 * components of our VUE header. See brw_vue_map.c for the layout. 374 */ 375static void 376update_so_info(struct pipe_stream_output_info *so_info, 377 uint64_t outputs_written) 378{ 379 uint8_t reverse_map[64] = {}; 380 unsigned slot = 0; 381 while (outputs_written) { 382 reverse_map[slot++] = u_bit_scan64(&outputs_written); 383 } 384 385 for (unsigned i = 0; i < so_info->num_outputs; i++) { 386 struct pipe_stream_output *output = &so_info->output[i]; 387 388 /* Map Gallium's condensed "slots" back to real VARYING_SLOT_* enums */ 389 output->register_index = reverse_map[output->register_index]; 390 391 /* The VUE header contains three scalar fields packed together: 392 * - gl_PointSize is stored in VARYING_SLOT_PSIZ.w 393 * - gl_Layer is stored in VARYING_SLOT_PSIZ.y 394 * - gl_ViewportIndex is stored in VARYING_SLOT_PSIZ.z 395 */ 396 switch (output->register_index) { 397 case VARYING_SLOT_LAYER: 398 assert(output->num_components == 1); 399 output->register_index = VARYING_SLOT_PSIZ; 400 output->start_component = 1; 401 break; 402 case VARYING_SLOT_VIEWPORT: 403 assert(output->num_components == 1); 404 output->register_index = VARYING_SLOT_PSIZ; 405 output->start_component = 2; 406 break; 407 case VARYING_SLOT_PSIZ: 408 assert(output->num_components == 1); 409 output->start_component = 3; 410 break; 411 } 412 413 //info->outputs_written |= 1ull << output->register_index; 414 } 415} 416 417static void 418setup_vec4_image_sysval(uint32_t *sysvals, uint32_t idx, 419 unsigned offset, unsigned n) 420{ 421 assert(offset % sizeof(uint32_t) == 0); 422 423 for (unsigned i = 0; i < n; ++i) 424 sysvals[i] = BRW_PARAM_IMAGE(idx, offset / sizeof(uint32_t) + i); 425 426 for (unsigned i = n; i < 4; ++i) 427 sysvals[i] = BRW_PARAM_BUILTIN_ZERO; 428} 429 430/** 431 * Associate NIR uniform variables with the prog_data->param[] mechanism 432 * used by the backend. Also, decide which UBOs we'd like to push in an 433 * ideal situation (though the backend can reduce this). 434 */ 435static void 436crocus_setup_uniforms(const struct brw_compiler *compiler, 437 void *mem_ctx, 438 nir_shader *nir, 439 struct brw_stage_prog_data *prog_data, 440 enum brw_param_builtin **out_system_values, 441 unsigned *out_num_system_values, 442 unsigned *out_num_cbufs) 443{ 444 UNUSED const struct intel_device_info *devinfo = compiler->devinfo; 445 446 const unsigned CROCUS_MAX_SYSTEM_VALUES = 447 PIPE_MAX_SHADER_IMAGES * BRW_IMAGE_PARAM_SIZE; 448 enum brw_param_builtin *system_values = 449 rzalloc_array(mem_ctx, enum brw_param_builtin, CROCUS_MAX_SYSTEM_VALUES); 450 unsigned num_system_values = 0; 451 452 unsigned patch_vert_idx = -1; 453 unsigned ucp_idx[CROCUS_MAX_CLIP_PLANES]; 454 unsigned img_idx[PIPE_MAX_SHADER_IMAGES]; 455 unsigned variable_group_size_idx = -1; 456 memset(ucp_idx, -1, sizeof(ucp_idx)); 457 memset(img_idx, -1, sizeof(img_idx)); 458 459 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 460 461 nir_builder b; 462 nir_builder_init(&b, impl); 463 464 b.cursor = nir_before_block(nir_start_block(impl)); 465 nir_ssa_def *temp_ubo_name = nir_ssa_undef(&b, 1, 32); 466 nir_ssa_def *temp_const_ubo_name = NULL; 467 468 /* Turn system value intrinsics into uniforms */ 469 nir_foreach_block(block, impl) { 470 nir_foreach_instr_safe(instr, block) { 471 if (instr->type != nir_instr_type_intrinsic) 472 continue; 473 474 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr); 475 nir_ssa_def *offset; 476 477 switch (intrin->intrinsic) { 478 case nir_intrinsic_load_constant: { 479 /* This one is special because it reads from the shader constant 480 * data and not cbuf0 which gallium uploads for us. 481 */ 482 b.cursor = nir_before_instr(instr); 483 nir_ssa_def *offset = 484 nir_iadd_imm(&b, nir_ssa_for_src(&b, intrin->src[0], 1), 485 nir_intrinsic_base(intrin)); 486 487 if (temp_const_ubo_name == NULL) 488 temp_const_ubo_name = nir_imm_int(&b, 0); 489 490 nir_intrinsic_instr *load_ubo = 491 nir_intrinsic_instr_create(b.shader, nir_intrinsic_load_ubo); 492 load_ubo->num_components = intrin->num_components; 493 load_ubo->src[0] = nir_src_for_ssa(temp_const_ubo_name); 494 load_ubo->src[1] = nir_src_for_ssa(offset); 495 nir_intrinsic_set_align(load_ubo, 4, 0); 496 nir_intrinsic_set_range_base(load_ubo, 0); 497 nir_intrinsic_set_range(load_ubo, ~0); 498 nir_ssa_dest_init(&load_ubo->instr, &load_ubo->dest, 499 intrin->dest.ssa.num_components, 500 intrin->dest.ssa.bit_size, 501 NULL); 502 nir_builder_instr_insert(&b, &load_ubo->instr); 503 504 nir_ssa_def_rewrite_uses(&intrin->dest.ssa, 505 &load_ubo->dest.ssa); 506 nir_instr_remove(&intrin->instr); 507 continue; 508 } 509 case nir_intrinsic_load_user_clip_plane: { 510 unsigned ucp = nir_intrinsic_ucp_id(intrin); 511 512 if (ucp_idx[ucp] == -1) { 513 ucp_idx[ucp] = num_system_values; 514 num_system_values += 4; 515 } 516 517 for (int i = 0; i < 4; i++) { 518 system_values[ucp_idx[ucp] + i] = 519 BRW_PARAM_BUILTIN_CLIP_PLANE(ucp, i); 520 } 521 522 b.cursor = nir_before_instr(instr); 523 offset = nir_imm_int(&b, ucp_idx[ucp] * sizeof(uint32_t)); 524 break; 525 } 526 case nir_intrinsic_load_patch_vertices_in: 527 if (patch_vert_idx == -1) 528 patch_vert_idx = num_system_values++; 529 530 system_values[patch_vert_idx] = 531 BRW_PARAM_BUILTIN_PATCH_VERTICES_IN; 532 533 b.cursor = nir_before_instr(instr); 534 offset = nir_imm_int(&b, patch_vert_idx * sizeof(uint32_t)); 535 break; 536 case nir_intrinsic_image_deref_load_param_intel: { 537 assert(devinfo->ver < 9); 538 nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); 539 nir_variable *var = nir_deref_instr_get_variable(deref); 540 541 if (img_idx[var->data.binding] == -1) { 542 /* GL only allows arrays of arrays of images. */ 543 assert(glsl_type_is_image(glsl_without_array(var->type))); 544 unsigned num_images = MAX2(1, glsl_get_aoa_size(var->type)); 545 546 for (int i = 0; i < num_images; i++) { 547 const unsigned img = var->data.binding + i; 548 549 img_idx[img] = num_system_values; 550 num_system_values += BRW_IMAGE_PARAM_SIZE; 551 552 uint32_t *img_sv = &system_values[img_idx[img]]; 553 554 setup_vec4_image_sysval( 555 img_sv + BRW_IMAGE_PARAM_OFFSET_OFFSET, img, 556 offsetof(struct brw_image_param, offset), 2); 557 setup_vec4_image_sysval( 558 img_sv + BRW_IMAGE_PARAM_SIZE_OFFSET, img, 559 offsetof(struct brw_image_param, size), 3); 560 setup_vec4_image_sysval( 561 img_sv + BRW_IMAGE_PARAM_STRIDE_OFFSET, img, 562 offsetof(struct brw_image_param, stride), 4); 563 setup_vec4_image_sysval( 564 img_sv + BRW_IMAGE_PARAM_TILING_OFFSET, img, 565 offsetof(struct brw_image_param, tiling), 3); 566 setup_vec4_image_sysval( 567 img_sv + BRW_IMAGE_PARAM_SWIZZLING_OFFSET, img, 568 offsetof(struct brw_image_param, swizzling), 2); 569 } 570 } 571 572 b.cursor = nir_before_instr(instr); 573 offset = nir_iadd(&b, 574 get_aoa_deref_offset(&b, deref, BRW_IMAGE_PARAM_SIZE * 4), 575 nir_imm_int(&b, img_idx[var->data.binding] * 4 + 576 nir_intrinsic_base(intrin) * 16)); 577 break; 578 } 579 case nir_intrinsic_load_workgroup_size: { 580 assert(nir->info.workgroup_size_variable); 581 if (variable_group_size_idx == -1) { 582 variable_group_size_idx = num_system_values; 583 num_system_values += 3; 584 for (int i = 0; i < 3; i++) { 585 system_values[variable_group_size_idx + i] = 586 BRW_PARAM_BUILTIN_WORK_GROUP_SIZE_X + i; 587 } 588 } 589 590 b.cursor = nir_before_instr(instr); 591 offset = nir_imm_int(&b, 592 variable_group_size_idx * sizeof(uint32_t)); 593 break; 594 } 595 default: 596 continue; 597 } 598 599 unsigned comps = nir_intrinsic_dest_components(intrin); 600 601 nir_intrinsic_instr *load = 602 nir_intrinsic_instr_create(nir, nir_intrinsic_load_ubo); 603 load->num_components = comps; 604 load->src[0] = nir_src_for_ssa(temp_ubo_name); 605 load->src[1] = nir_src_for_ssa(offset); 606 nir_intrinsic_set_align(load, 4, 0); 607 nir_intrinsic_set_range_base(load, 0); 608 nir_intrinsic_set_range(load, ~0); 609 nir_ssa_dest_init(&load->instr, &load->dest, comps, 32, NULL); 610 nir_builder_instr_insert(&b, &load->instr); 611 nir_ssa_def_rewrite_uses(&intrin->dest.ssa, 612 &load->dest.ssa); 613 nir_instr_remove(instr); 614 } 615 } 616 617 nir_validate_shader(nir, "before remapping"); 618 619 /* Uniforms are stored in constant buffer 0, the 620 * user-facing UBOs are indexed by one. So if any constant buffer is 621 * needed, the constant buffer 0 will be needed, so account for it. 622 */ 623 unsigned num_cbufs = nir->info.num_ubos; 624 if (num_cbufs || nir->num_uniforms) 625 num_cbufs++; 626 627 /* Place the new params in a new cbuf. */ 628 if (num_system_values > 0) { 629 unsigned sysval_cbuf_index = num_cbufs; 630 num_cbufs++; 631 632 system_values = reralloc(mem_ctx, system_values, enum brw_param_builtin, 633 num_system_values); 634 635 nir_foreach_block(block, impl) { 636 nir_foreach_instr_safe(instr, block) { 637 if (instr->type != nir_instr_type_intrinsic) 638 continue; 639 640 nir_intrinsic_instr *load = nir_instr_as_intrinsic(instr); 641 642 if (load->intrinsic != nir_intrinsic_load_ubo) 643 continue; 644 645 b.cursor = nir_before_instr(instr); 646 647 assert(load->src[0].is_ssa); 648 649 if (load->src[0].ssa == temp_ubo_name) { 650 nir_ssa_def *imm = nir_imm_int(&b, sysval_cbuf_index); 651 nir_instr_rewrite_src(instr, &load->src[0], 652 nir_src_for_ssa(imm)); 653 } 654 } 655 } 656 657 /* We need to fold the new iadds for brw_nir_analyze_ubo_ranges */ 658 nir_opt_constant_folding(nir); 659 } else { 660 ralloc_free(system_values); 661 system_values = NULL; 662 } 663 664 assert(num_cbufs < PIPE_MAX_CONSTANT_BUFFERS); 665 nir_validate_shader(nir, "after remap"); 666 667 /* We don't use params[] but gallium leaves num_uniforms set. We use this 668 * to detect when cbuf0 exists but we don't need it anymore when we get 669 * here. Instead, zero it out so that the back-end doesn't get confused 670 * when nr_params * 4 != num_uniforms != nr_params * 4. 671 */ 672 nir->num_uniforms = 0; 673 674 /* Constant loads (if any) need to go at the end of the constant buffers so 675 * we need to know num_cbufs before we can lower to them. 676 */ 677 if (temp_const_ubo_name != NULL) { 678 nir_load_const_instr *const_ubo_index = 679 nir_instr_as_load_const(temp_const_ubo_name->parent_instr); 680 assert(const_ubo_index->def.bit_size == 32); 681 const_ubo_index->value[0].u32 = num_cbufs; 682 } 683 684 *out_system_values = system_values; 685 *out_num_system_values = num_system_values; 686 *out_num_cbufs = num_cbufs; 687} 688 689static const char *surface_group_names[] = { 690 [CROCUS_SURFACE_GROUP_RENDER_TARGET] = "render target", 691 [CROCUS_SURFACE_GROUP_RENDER_TARGET_READ] = "non-coherent render target read", 692 [CROCUS_SURFACE_GROUP_SOL] = "streamout", 693 [CROCUS_SURFACE_GROUP_CS_WORK_GROUPS] = "CS work groups", 694 [CROCUS_SURFACE_GROUP_TEXTURE] = "texture", 695 [CROCUS_SURFACE_GROUP_TEXTURE_GATHER] = "texture gather", 696 [CROCUS_SURFACE_GROUP_UBO] = "ubo", 697 [CROCUS_SURFACE_GROUP_SSBO] = "ssbo", 698 [CROCUS_SURFACE_GROUP_IMAGE] = "image", 699}; 700 701static void 702crocus_print_binding_table(FILE *fp, const char *name, 703 const struct crocus_binding_table *bt) 704{ 705 STATIC_ASSERT(ARRAY_SIZE(surface_group_names) == CROCUS_SURFACE_GROUP_COUNT); 706 707 uint32_t total = 0; 708 uint32_t compacted = 0; 709 710 for (int i = 0; i < CROCUS_SURFACE_GROUP_COUNT; i++) { 711 uint32_t size = bt->sizes[i]; 712 total += size; 713 if (size) 714 compacted += util_bitcount64(bt->used_mask[i]); 715 } 716 717 if (total == 0) { 718 fprintf(fp, "Binding table for %s is empty\n\n", name); 719 return; 720 } 721 722 if (total != compacted) { 723 fprintf(fp, "Binding table for %s " 724 "(compacted to %u entries from %u entries)\n", 725 name, compacted, total); 726 } else { 727 fprintf(fp, "Binding table for %s (%u entries)\n", name, total); 728 } 729 730 uint32_t entry = 0; 731 for (int i = 0; i < CROCUS_SURFACE_GROUP_COUNT; i++) { 732 uint64_t mask = bt->used_mask[i]; 733 while (mask) { 734 int index = u_bit_scan64(&mask); 735 fprintf(fp, " [%u] %s #%d\n", entry++, surface_group_names[i], index); 736 } 737 } 738 fprintf(fp, "\n"); 739} 740 741enum { 742 /* Max elements in a surface group. */ 743 SURFACE_GROUP_MAX_ELEMENTS = 64, 744}; 745 746static void 747rewrite_src_with_bti(nir_builder *b, struct crocus_binding_table *bt, 748 nir_instr *instr, nir_src *src, 749 enum crocus_surface_group group) 750{ 751 assert(bt->sizes[group] > 0); 752 753 b->cursor = nir_before_instr(instr); 754 nir_ssa_def *bti; 755 if (nir_src_is_const(*src)) { 756 uint32_t index = nir_src_as_uint(*src); 757 bti = nir_imm_intN_t(b, crocus_group_index_to_bti(bt, group, index), 758 src->ssa->bit_size); 759 } else { 760 /* Indirect usage makes all the surfaces of the group to be available, 761 * so we can just add the base. 762 */ 763 assert(bt->used_mask[group] == BITFIELD64_MASK(bt->sizes[group])); 764 bti = nir_iadd_imm(b, src->ssa, bt->offsets[group]); 765 } 766 nir_instr_rewrite_src(instr, src, nir_src_for_ssa(bti)); 767} 768 769static void 770mark_used_with_src(struct crocus_binding_table *bt, nir_src *src, 771 enum crocus_surface_group group) 772{ 773 assert(bt->sizes[group] > 0); 774 775 if (nir_src_is_const(*src)) { 776 uint64_t index = nir_src_as_uint(*src); 777 assert(index < bt->sizes[group]); 778 bt->used_mask[group] |= 1ull << index; 779 } else { 780 /* There's an indirect usage, we need all the surfaces. */ 781 bt->used_mask[group] = BITFIELD64_MASK(bt->sizes[group]); 782 } 783} 784 785static bool 786skip_compacting_binding_tables(void) 787{ 788 static int skip = -1; 789 if (skip < 0) 790 skip = env_var_as_boolean("INTEL_DISABLE_COMPACT_BINDING_TABLE", false); 791 return skip; 792} 793 794/** 795 * Set up the binding table indices and apply to the shader. 796 */ 797static void 798crocus_setup_binding_table(const struct intel_device_info *devinfo, 799 struct nir_shader *nir, 800 struct crocus_binding_table *bt, 801 unsigned num_render_targets, 802 unsigned num_system_values, 803 unsigned num_cbufs, 804 const struct brw_sampler_prog_key_data *key) 805{ 806 const struct shader_info *info = &nir->info; 807 808 memset(bt, 0, sizeof(*bt)); 809 810 /* Set the sizes for each surface group. For some groups, we already know 811 * upfront how many will be used, so mark them. 812 */ 813 if (info->stage == MESA_SHADER_FRAGMENT) { 814 bt->sizes[CROCUS_SURFACE_GROUP_RENDER_TARGET] = num_render_targets; 815 /* All render targets used. */ 816 bt->used_mask[CROCUS_SURFACE_GROUP_RENDER_TARGET] = 817 BITFIELD64_MASK(num_render_targets); 818 819 /* Setup render target read surface group in order to support non-coherent 820 * framebuffer fetch on Gfx7 821 */ 822 if (devinfo->ver >= 6 && info->outputs_read) { 823 bt->sizes[CROCUS_SURFACE_GROUP_RENDER_TARGET_READ] = num_render_targets; 824 bt->used_mask[CROCUS_SURFACE_GROUP_RENDER_TARGET_READ] = 825 BITFIELD64_MASK(num_render_targets); 826 } 827 } else if (info->stage == MESA_SHADER_COMPUTE) { 828 bt->sizes[CROCUS_SURFACE_GROUP_CS_WORK_GROUPS] = 1; 829 } else if (info->stage == MESA_SHADER_GEOMETRY) { 830 /* In gfx6 we reserve the first BRW_MAX_SOL_BINDINGS entries for transform 831 * feedback surfaces. 832 */ 833 if (devinfo->ver == 6) { 834 bt->sizes[CROCUS_SURFACE_GROUP_SOL] = BRW_MAX_SOL_BINDINGS; 835 bt->used_mask[CROCUS_SURFACE_GROUP_SOL] = (uint64_t)-1; 836 } 837 } 838 839 bt->sizes[CROCUS_SURFACE_GROUP_TEXTURE] = BITSET_LAST_BIT(info->textures_used); 840 bt->used_mask[CROCUS_SURFACE_GROUP_TEXTURE] = info->textures_used[0]; 841 842 if (info->uses_texture_gather && devinfo->ver < 8) { 843 bt->sizes[CROCUS_SURFACE_GROUP_TEXTURE_GATHER] = BITSET_LAST_BIT(info->textures_used); 844 bt->used_mask[CROCUS_SURFACE_GROUP_TEXTURE_GATHER] = info->textures_used[0]; 845 } 846 847 bt->sizes[CROCUS_SURFACE_GROUP_IMAGE] = info->num_images; 848 849 /* Allocate an extra slot in the UBO section for NIR constants. 850 * Binding table compaction will remove it if unnecessary. 851 * 852 * We don't include them in crocus_compiled_shader::num_cbufs because 853 * they are uploaded separately from shs->constbufs[], but from a shader 854 * point of view, they're another UBO (at the end of the section). 855 */ 856 bt->sizes[CROCUS_SURFACE_GROUP_UBO] = num_cbufs + 1; 857 858 bt->sizes[CROCUS_SURFACE_GROUP_SSBO] = info->num_ssbos; 859 860 for (int i = 0; i < CROCUS_SURFACE_GROUP_COUNT; i++) 861 assert(bt->sizes[i] <= SURFACE_GROUP_MAX_ELEMENTS); 862 863 /* Mark surfaces used for the cases we don't have the information available 864 * upfront. 865 */ 866 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 867 nir_foreach_block (block, impl) { 868 nir_foreach_instr (instr, block) { 869 if (instr->type != nir_instr_type_intrinsic) 870 continue; 871 872 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr); 873 switch (intrin->intrinsic) { 874 case nir_intrinsic_load_num_workgroups: 875 bt->used_mask[CROCUS_SURFACE_GROUP_CS_WORK_GROUPS] = 1; 876 break; 877 878 case nir_intrinsic_load_output: 879 if (devinfo->ver >= 6) { 880 mark_used_with_src(bt, &intrin->src[0], 881 CROCUS_SURFACE_GROUP_RENDER_TARGET_READ); 882 } 883 break; 884 885 case nir_intrinsic_image_size: 886 case nir_intrinsic_image_load: 887 case nir_intrinsic_image_store: 888 case nir_intrinsic_image_atomic_add: 889 case nir_intrinsic_image_atomic_imin: 890 case nir_intrinsic_image_atomic_umin: 891 case nir_intrinsic_image_atomic_imax: 892 case nir_intrinsic_image_atomic_umax: 893 case nir_intrinsic_image_atomic_and: 894 case nir_intrinsic_image_atomic_or: 895 case nir_intrinsic_image_atomic_xor: 896 case nir_intrinsic_image_atomic_exchange: 897 case nir_intrinsic_image_atomic_comp_swap: 898 case nir_intrinsic_image_load_raw_intel: 899 case nir_intrinsic_image_store_raw_intel: 900 mark_used_with_src(bt, &intrin->src[0], CROCUS_SURFACE_GROUP_IMAGE); 901 break; 902 903 case nir_intrinsic_load_ubo: 904 mark_used_with_src(bt, &intrin->src[0], CROCUS_SURFACE_GROUP_UBO); 905 break; 906 907 case nir_intrinsic_store_ssbo: 908 mark_used_with_src(bt, &intrin->src[1], CROCUS_SURFACE_GROUP_SSBO); 909 break; 910 911 case nir_intrinsic_get_ssbo_size: 912 case nir_intrinsic_ssbo_atomic_add: 913 case nir_intrinsic_ssbo_atomic_imin: 914 case nir_intrinsic_ssbo_atomic_umin: 915 case nir_intrinsic_ssbo_atomic_imax: 916 case nir_intrinsic_ssbo_atomic_umax: 917 case nir_intrinsic_ssbo_atomic_and: 918 case nir_intrinsic_ssbo_atomic_or: 919 case nir_intrinsic_ssbo_atomic_xor: 920 case nir_intrinsic_ssbo_atomic_exchange: 921 case nir_intrinsic_ssbo_atomic_comp_swap: 922 case nir_intrinsic_ssbo_atomic_fmin: 923 case nir_intrinsic_ssbo_atomic_fmax: 924 case nir_intrinsic_ssbo_atomic_fcomp_swap: 925 case nir_intrinsic_load_ssbo: 926 mark_used_with_src(bt, &intrin->src[0], CROCUS_SURFACE_GROUP_SSBO); 927 break; 928 929 default: 930 break; 931 } 932 } 933 } 934 935 /* When disable we just mark everything as used. */ 936 if (unlikely(skip_compacting_binding_tables())) { 937 for (int i = 0; i < CROCUS_SURFACE_GROUP_COUNT; i++) 938 bt->used_mask[i] = BITFIELD64_MASK(bt->sizes[i]); 939 } 940 941 /* Calculate the offsets and the binding table size based on the used 942 * surfaces. After this point, the functions to go between "group indices" 943 * and binding table indices can be used. 944 */ 945 uint32_t next = 0; 946 for (int i = 0; i < CROCUS_SURFACE_GROUP_COUNT; i++) { 947 if (bt->used_mask[i] != 0) { 948 bt->offsets[i] = next; 949 next += util_bitcount64(bt->used_mask[i]); 950 } 951 } 952 bt->size_bytes = next * 4; 953 954 if (INTEL_DEBUG(DEBUG_BT)) { 955 crocus_print_binding_table(stderr, gl_shader_stage_name(info->stage), bt); 956 } 957 958 /* Apply the binding table indices. The backend compiler is not expected 959 * to change those, as we haven't set any of the *_start entries in brw 960 * binding_table. 961 */ 962 nir_builder b; 963 nir_builder_init(&b, impl); 964 965 nir_foreach_block (block, impl) { 966 nir_foreach_instr (instr, block) { 967 if (instr->type == nir_instr_type_tex) { 968 nir_tex_instr *tex = nir_instr_as_tex(instr); 969 bool is_gather = devinfo->ver < 8 && tex->op == nir_texop_tg4; 970 971 /* rewrite the tg4 component from green to blue before replacing the 972 texture index */ 973 if (devinfo->verx10 == 70) { 974 if (tex->component == 1) 975 if (key->gather_channel_quirk_mask & (1 << tex->texture_index)) 976 tex->component = 2; 977 } 978 979 if (is_gather && devinfo->ver == 6 && key->gfx6_gather_wa[tex->texture_index]) { 980 b.cursor = nir_after_instr(instr); 981 enum gfx6_gather_sampler_wa wa = key->gfx6_gather_wa[tex->texture_index]; 982 int width = (wa & WA_8BIT) ? 8 : 16; 983 984 nir_ssa_def *val = nir_fmul_imm(&b, &tex->dest.ssa, (1 << width) - 1); 985 val = nir_f2u32(&b, val); 986 if (wa & WA_SIGN) { 987 val = nir_ishl(&b, val, nir_imm_int(&b, 32 - width)); 988 val = nir_ishr(&b, val, nir_imm_int(&b, 32 - width)); 989 } 990 nir_ssa_def_rewrite_uses_after(&tex->dest.ssa, val, val->parent_instr); 991 } 992 993 tex->texture_index = 994 crocus_group_index_to_bti(bt, is_gather ? CROCUS_SURFACE_GROUP_TEXTURE_GATHER : CROCUS_SURFACE_GROUP_TEXTURE, 995 tex->texture_index); 996 continue; 997 } 998 999 if (instr->type != nir_instr_type_intrinsic) 1000 continue; 1001 1002 nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr); 1003 switch (intrin->intrinsic) { 1004 case nir_intrinsic_image_size: 1005 case nir_intrinsic_image_load: 1006 case nir_intrinsic_image_store: 1007 case nir_intrinsic_image_atomic_add: 1008 case nir_intrinsic_image_atomic_imin: 1009 case nir_intrinsic_image_atomic_umin: 1010 case nir_intrinsic_image_atomic_imax: 1011 case nir_intrinsic_image_atomic_umax: 1012 case nir_intrinsic_image_atomic_and: 1013 case nir_intrinsic_image_atomic_or: 1014 case nir_intrinsic_image_atomic_xor: 1015 case nir_intrinsic_image_atomic_exchange: 1016 case nir_intrinsic_image_atomic_comp_swap: 1017 case nir_intrinsic_image_load_raw_intel: 1018 case nir_intrinsic_image_store_raw_intel: 1019 rewrite_src_with_bti(&b, bt, instr, &intrin->src[0], 1020 CROCUS_SURFACE_GROUP_IMAGE); 1021 break; 1022 1023 case nir_intrinsic_load_ubo: 1024 rewrite_src_with_bti(&b, bt, instr, &intrin->src[0], 1025 CROCUS_SURFACE_GROUP_UBO); 1026 break; 1027 1028 case nir_intrinsic_store_ssbo: 1029 rewrite_src_with_bti(&b, bt, instr, &intrin->src[1], 1030 CROCUS_SURFACE_GROUP_SSBO); 1031 break; 1032 1033 case nir_intrinsic_load_output: 1034 if (devinfo->ver >= 6) { 1035 rewrite_src_with_bti(&b, bt, instr, &intrin->src[0], 1036 CROCUS_SURFACE_GROUP_RENDER_TARGET_READ); 1037 } 1038 break; 1039 1040 case nir_intrinsic_get_ssbo_size: 1041 case nir_intrinsic_ssbo_atomic_add: 1042 case nir_intrinsic_ssbo_atomic_imin: 1043 case nir_intrinsic_ssbo_atomic_umin: 1044 case nir_intrinsic_ssbo_atomic_imax: 1045 case nir_intrinsic_ssbo_atomic_umax: 1046 case nir_intrinsic_ssbo_atomic_and: 1047 case nir_intrinsic_ssbo_atomic_or: 1048 case nir_intrinsic_ssbo_atomic_xor: 1049 case nir_intrinsic_ssbo_atomic_exchange: 1050 case nir_intrinsic_ssbo_atomic_comp_swap: 1051 case nir_intrinsic_ssbo_atomic_fmin: 1052 case nir_intrinsic_ssbo_atomic_fmax: 1053 case nir_intrinsic_ssbo_atomic_fcomp_swap: 1054 case nir_intrinsic_load_ssbo: 1055 rewrite_src_with_bti(&b, bt, instr, &intrin->src[0], 1056 CROCUS_SURFACE_GROUP_SSBO); 1057 break; 1058 1059 default: 1060 break; 1061 } 1062 } 1063 } 1064} 1065 1066static void 1067crocus_debug_recompile(struct crocus_context *ice, 1068 struct shader_info *info, 1069 const struct brw_base_prog_key *key) 1070{ 1071 struct crocus_screen *screen = (struct crocus_screen *) ice->ctx.screen; 1072 const struct brw_compiler *c = screen->compiler; 1073 1074 if (!info) 1075 return; 1076 1077 brw_shader_perf_log(c, &ice->dbg, "Recompiling %s shader for program %s: %s\n", 1078 _mesa_shader_stage_to_string(info->stage), 1079 info->name ? info->name : "(no identifier)", 1080 info->label ? info->label : ""); 1081 1082 const void *old_key = 1083 crocus_find_previous_compile(ice, info->stage, key->program_string_id); 1084 1085 brw_debug_key_recompile(c, &ice->dbg, info->stage, old_key, key); 1086} 1087 1088/** 1089 * Get the shader for the last enabled geometry stage. 1090 * 1091 * This stage is the one which will feed stream output and the rasterizer. 1092 */ 1093static gl_shader_stage 1094last_vue_stage(struct crocus_context *ice) 1095{ 1096 if (ice->shaders.uncompiled[MESA_SHADER_GEOMETRY]) 1097 return MESA_SHADER_GEOMETRY; 1098 1099 if (ice->shaders.uncompiled[MESA_SHADER_TESS_EVAL]) 1100 return MESA_SHADER_TESS_EVAL; 1101 1102 return MESA_SHADER_VERTEX; 1103} 1104 1105static GLbitfield64 1106crocus_vs_outputs_written(struct crocus_context *ice, 1107 const struct brw_vs_prog_key *key, 1108 GLbitfield64 user_varyings) 1109{ 1110 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1111 const struct intel_device_info *devinfo = &screen->devinfo; 1112 GLbitfield64 outputs_written = user_varyings; 1113 1114 if (devinfo->ver < 6) { 1115 1116 if (key->copy_edgeflag) 1117 outputs_written |= BITFIELD64_BIT(VARYING_SLOT_EDGE); 1118 1119 /* Put dummy slots into the VUE for the SF to put the replaced 1120 * point sprite coords in. We shouldn't need these dummy slots, 1121 * which take up precious URB space, but it would mean that the SF 1122 * doesn't get nice aligned pairs of input coords into output 1123 * coords, which would be a pain to handle. 1124 */ 1125 for (unsigned i = 0; i < 8; i++) { 1126 if (key->point_coord_replace & (1 << i)) 1127 outputs_written |= BITFIELD64_BIT(VARYING_SLOT_TEX0 + i); 1128 } 1129 1130 /* if back colors are written, allocate slots for front colors too */ 1131 if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_BFC0)) 1132 outputs_written |= BITFIELD64_BIT(VARYING_SLOT_COL0); 1133 if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_BFC1)) 1134 outputs_written |= BITFIELD64_BIT(VARYING_SLOT_COL1); 1135 } 1136 1137 /* In order for legacy clipping to work, we need to populate the clip 1138 * distance varying slots whenever clipping is enabled, even if the vertex 1139 * shader doesn't write to gl_ClipDistance. 1140 */ 1141 if (key->nr_userclip_plane_consts > 0) { 1142 outputs_written |= BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST0); 1143 outputs_written |= BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST1); 1144 } 1145 1146 return outputs_written; 1147} 1148 1149/* 1150 * If no edgeflags come from the user, gen4/5 1151 * require giving the clip shader a default edgeflag. 1152 * 1153 * This will always be 1.0. 1154 */ 1155static void 1156crocus_lower_default_edgeflags(struct nir_shader *nir) 1157{ 1158 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 1159 1160 nir_builder b; 1161 nir_builder_init(&b, impl); 1162 1163 b.cursor = nir_after_cf_list(&b.impl->body); 1164 nir_variable *var = nir_variable_create(nir, nir_var_shader_out, 1165 glsl_float_type(), 1166 "edgeflag"); 1167 var->data.location = VARYING_SLOT_EDGE; 1168 nir_store_var(&b, var, nir_imm_float(&b, 1.0), 0x1); 1169} 1170 1171/** 1172 * Compile a vertex shader, and upload the assembly. 1173 */ 1174static struct crocus_compiled_shader * 1175crocus_compile_vs(struct crocus_context *ice, 1176 struct crocus_uncompiled_shader *ish, 1177 const struct brw_vs_prog_key *key) 1178{ 1179 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1180 const struct brw_compiler *compiler = screen->compiler; 1181 const struct intel_device_info *devinfo = &screen->devinfo; 1182 void *mem_ctx = ralloc_context(NULL); 1183 struct brw_vs_prog_data *vs_prog_data = 1184 rzalloc(mem_ctx, struct brw_vs_prog_data); 1185 struct brw_vue_prog_data *vue_prog_data = &vs_prog_data->base; 1186 struct brw_stage_prog_data *prog_data = &vue_prog_data->base; 1187 enum brw_param_builtin *system_values; 1188 unsigned num_system_values; 1189 unsigned num_cbufs; 1190 1191 nir_shader *nir = nir_shader_clone(mem_ctx, ish->nir); 1192 1193 if (key->nr_userclip_plane_consts) { 1194 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 1195 nir_lower_clip_vs(nir, (1 << key->nr_userclip_plane_consts) - 1, true, 1196 false, NULL); 1197 nir_lower_io_to_temporaries(nir, impl, true, false); 1198 nir_lower_global_vars_to_local(nir); 1199 nir_lower_vars_to_ssa(nir); 1200 nir_shader_gather_info(nir, impl); 1201 } 1202 1203 prog_data->use_alt_mode = nir->info.is_arb_asm; 1204 1205 crocus_setup_uniforms(compiler, mem_ctx, nir, prog_data, &system_values, 1206 &num_system_values, &num_cbufs); 1207 1208 crocus_lower_swizzles(nir, &key->base.tex); 1209 1210 if (devinfo->ver <= 5 && 1211 !(nir->info.inputs_read & BITFIELD64_BIT(VERT_ATTRIB_EDGEFLAG))) 1212 crocus_lower_default_edgeflags(nir); 1213 1214 struct crocus_binding_table bt; 1215 crocus_setup_binding_table(devinfo, nir, &bt, /* num_render_targets */ 0, 1216 num_system_values, num_cbufs, &key->base.tex); 1217 1218 if (can_push_ubo(devinfo)) 1219 brw_nir_analyze_ubo_ranges(compiler, nir, NULL, prog_data->ubo_ranges); 1220 1221 uint64_t outputs_written = 1222 crocus_vs_outputs_written(ice, key, nir->info.outputs_written); 1223 brw_compute_vue_map(devinfo, 1224 &vue_prog_data->vue_map, outputs_written, 1225 nir->info.separate_shader, /* pos slots */ 1); 1226 1227 /* Don't tell the backend about our clip plane constants, we've already 1228 * lowered them in NIR and we don't want it doing it again. 1229 */ 1230 struct brw_vs_prog_key key_no_ucp = *key; 1231 key_no_ucp.nr_userclip_plane_consts = 0; 1232 key_no_ucp.copy_edgeflag = false; 1233 crocus_sanitize_tex_key(&key_no_ucp.base.tex); 1234 1235 struct brw_compile_vs_params params = { 1236 .nir = nir, 1237 .key = &key_no_ucp, 1238 .prog_data = vs_prog_data, 1239 .edgeflag_is_last = devinfo->ver < 6, 1240 .log_data = &ice->dbg, 1241 }; 1242 const unsigned *program = 1243 brw_compile_vs(compiler, mem_ctx, ¶ms); 1244 if (program == NULL) { 1245 dbg_printf("Failed to compile vertex shader: %s\n", params.error_str); 1246 ralloc_free(mem_ctx); 1247 return false; 1248 } 1249 1250 if (ish->compiled_once) { 1251 crocus_debug_recompile(ice, &nir->info, &key->base); 1252 } else { 1253 ish->compiled_once = true; 1254 } 1255 1256 uint32_t *so_decls = NULL; 1257 if (devinfo->ver > 6) 1258 so_decls = screen->vtbl.create_so_decl_list(&ish->stream_output, 1259 &vue_prog_data->vue_map); 1260 1261 struct crocus_compiled_shader *shader = 1262 crocus_upload_shader(ice, CROCUS_CACHE_VS, sizeof(*key), key, program, 1263 prog_data->program_size, 1264 prog_data, sizeof(*vs_prog_data), so_decls, 1265 system_values, num_system_values, 1266 num_cbufs, &bt); 1267 1268 crocus_disk_cache_store(screen->disk_cache, ish, shader, 1269 ice->shaders.cache_bo_map, 1270 key, sizeof(*key)); 1271 1272 ralloc_free(mem_ctx); 1273 return shader; 1274} 1275 1276/** 1277 * Update the current vertex shader variant. 1278 * 1279 * Fill out the key, look in the cache, compile and bind if needed. 1280 */ 1281static void 1282crocus_update_compiled_vs(struct crocus_context *ice) 1283{ 1284 struct crocus_shader_state *shs = &ice->state.shaders[MESA_SHADER_VERTEX]; 1285 struct crocus_uncompiled_shader *ish = 1286 ice->shaders.uncompiled[MESA_SHADER_VERTEX]; 1287 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1288 const struct intel_device_info *devinfo = &screen->devinfo; 1289 struct brw_vs_prog_key key = { KEY_INIT() }; 1290 1291 if (ish->nos & (1ull << CROCUS_NOS_TEXTURES)) 1292 crocus_populate_sampler_prog_key_data(ice, devinfo, MESA_SHADER_VERTEX, ish, 1293 ish->nir->info.uses_texture_gather, &key.base.tex); 1294 screen->vtbl.populate_vs_key(ice, &ish->nir->info, last_vue_stage(ice), &key); 1295 1296 struct crocus_compiled_shader *old = ice->shaders.prog[CROCUS_CACHE_VS]; 1297 struct crocus_compiled_shader *shader = 1298 crocus_find_cached_shader(ice, CROCUS_CACHE_VS, sizeof(key), &key); 1299 1300 if (!shader) 1301 shader = crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key)); 1302 1303 if (!shader) 1304 shader = crocus_compile_vs(ice, ish, &key); 1305 1306 if (old != shader) { 1307 ice->shaders.prog[CROCUS_CACHE_VS] = shader; 1308 if (devinfo->ver == 8) 1309 ice->state.dirty |= CROCUS_DIRTY_GEN8_VF_SGVS; 1310 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_VS | 1311 CROCUS_STAGE_DIRTY_BINDINGS_VS | 1312 CROCUS_STAGE_DIRTY_CONSTANTS_VS; 1313 shs->sysvals_need_upload = true; 1314 1315 const struct brw_vs_prog_data *vs_prog_data = 1316 (void *) shader->prog_data; 1317 const bool uses_draw_params = vs_prog_data->uses_firstvertex || 1318 vs_prog_data->uses_baseinstance; 1319 const bool uses_derived_draw_params = vs_prog_data->uses_drawid || 1320 vs_prog_data->uses_is_indexed_draw; 1321 const bool needs_sgvs_element = uses_draw_params || 1322 vs_prog_data->uses_instanceid || 1323 vs_prog_data->uses_vertexid; 1324 1325 if (ice->state.vs_uses_draw_params != uses_draw_params || 1326 ice->state.vs_uses_derived_draw_params != uses_derived_draw_params || 1327 ice->state.vs_needs_edge_flag != ish->needs_edge_flag || 1328 ice->state.vs_uses_vertexid != vs_prog_data->uses_vertexid || 1329 ice->state.vs_uses_instanceid != vs_prog_data->uses_instanceid) { 1330 ice->state.dirty |= CROCUS_DIRTY_VERTEX_BUFFERS | 1331 CROCUS_DIRTY_VERTEX_ELEMENTS; 1332 } 1333 ice->state.vs_uses_draw_params = uses_draw_params; 1334 ice->state.vs_uses_derived_draw_params = uses_derived_draw_params; 1335 ice->state.vs_needs_sgvs_element = needs_sgvs_element; 1336 ice->state.vs_needs_edge_flag = ish->needs_edge_flag; 1337 ice->state.vs_uses_vertexid = vs_prog_data->uses_vertexid; 1338 ice->state.vs_uses_instanceid = vs_prog_data->uses_instanceid; 1339 } 1340} 1341 1342/** 1343 * Get the shader_info for a given stage, or NULL if the stage is disabled. 1344 */ 1345const struct shader_info * 1346crocus_get_shader_info(const struct crocus_context *ice, gl_shader_stage stage) 1347{ 1348 const struct crocus_uncompiled_shader *ish = ice->shaders.uncompiled[stage]; 1349 1350 if (!ish) 1351 return NULL; 1352 1353 const nir_shader *nir = ish->nir; 1354 return &nir->info; 1355} 1356 1357/** 1358 * Get the union of TCS output and TES input slots. 1359 * 1360 * TCS and TES need to agree on a common URB entry layout. In particular, 1361 * the data for all patch vertices is stored in a single URB entry (unlike 1362 * GS which has one entry per input vertex). This means that per-vertex 1363 * array indexing needs a stride. 1364 * 1365 * SSO requires locations to match, but doesn't require the number of 1366 * outputs/inputs to match (in fact, the TCS often has extra outputs). 1367 * So, we need to take the extra step of unifying these on the fly. 1368 */ 1369static void 1370get_unified_tess_slots(const struct crocus_context *ice, 1371 uint64_t *per_vertex_slots, 1372 uint32_t *per_patch_slots) 1373{ 1374 const struct shader_info *tcs = 1375 crocus_get_shader_info(ice, MESA_SHADER_TESS_CTRL); 1376 const struct shader_info *tes = 1377 crocus_get_shader_info(ice, MESA_SHADER_TESS_EVAL); 1378 1379 *per_vertex_slots = tes->inputs_read; 1380 *per_patch_slots = tes->patch_inputs_read; 1381 1382 if (tcs) { 1383 *per_vertex_slots |= tcs->outputs_written; 1384 *per_patch_slots |= tcs->patch_outputs_written; 1385 } 1386} 1387 1388/** 1389 * Compile a tessellation control shader, and upload the assembly. 1390 */ 1391static struct crocus_compiled_shader * 1392crocus_compile_tcs(struct crocus_context *ice, 1393 struct crocus_uncompiled_shader *ish, 1394 const struct brw_tcs_prog_key *key) 1395{ 1396 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1397 const struct brw_compiler *compiler = screen->compiler; 1398 const struct nir_shader_compiler_options *options = 1399 compiler->glsl_compiler_options[MESA_SHADER_TESS_CTRL].NirOptions; 1400 void *mem_ctx = ralloc_context(NULL); 1401 struct brw_tcs_prog_data *tcs_prog_data = 1402 rzalloc(mem_ctx, struct brw_tcs_prog_data); 1403 struct brw_vue_prog_data *vue_prog_data = &tcs_prog_data->base; 1404 struct brw_stage_prog_data *prog_data = &vue_prog_data->base; 1405 const struct intel_device_info *devinfo = &screen->devinfo; 1406 enum brw_param_builtin *system_values = NULL; 1407 unsigned num_system_values = 0; 1408 unsigned num_cbufs = 0; 1409 1410 nir_shader *nir; 1411 1412 struct crocus_binding_table bt; 1413 1414 if (ish) { 1415 nir = nir_shader_clone(mem_ctx, ish->nir); 1416 1417 crocus_setup_uniforms(compiler, mem_ctx, nir, prog_data, &system_values, 1418 &num_system_values, &num_cbufs); 1419 1420 crocus_lower_swizzles(nir, &key->base.tex); 1421 crocus_setup_binding_table(devinfo, nir, &bt, /* num_render_targets */ 0, 1422 num_system_values, num_cbufs, &key->base.tex); 1423 if (can_push_ubo(devinfo)) 1424 brw_nir_analyze_ubo_ranges(compiler, nir, NULL, prog_data->ubo_ranges); 1425 } else { 1426 nir = brw_nir_create_passthrough_tcs(mem_ctx, compiler, options, key); 1427 1428 /* Reserve space for passing the default tess levels as constants. */ 1429 num_cbufs = 1; 1430 num_system_values = 8; 1431 system_values = 1432 rzalloc_array(mem_ctx, enum brw_param_builtin, num_system_values); 1433 prog_data->param = rzalloc_array(mem_ctx, uint32_t, num_system_values); 1434 prog_data->nr_params = num_system_values; 1435 1436 if (key->tes_primitive_mode == GL_QUADS) { 1437 for (int i = 0; i < 4; i++) 1438 system_values[7 - i] = BRW_PARAM_BUILTIN_TESS_LEVEL_OUTER_X + i; 1439 1440 system_values[3] = BRW_PARAM_BUILTIN_TESS_LEVEL_INNER_X; 1441 system_values[2] = BRW_PARAM_BUILTIN_TESS_LEVEL_INNER_Y; 1442 } else if (key->tes_primitive_mode == GL_TRIANGLES) { 1443 for (int i = 0; i < 3; i++) 1444 system_values[7 - i] = BRW_PARAM_BUILTIN_TESS_LEVEL_OUTER_X + i; 1445 1446 system_values[4] = BRW_PARAM_BUILTIN_TESS_LEVEL_INNER_X; 1447 } else { 1448 assert(key->tes_primitive_mode == GL_ISOLINES); 1449 system_values[7] = BRW_PARAM_BUILTIN_TESS_LEVEL_OUTER_Y; 1450 system_values[6] = BRW_PARAM_BUILTIN_TESS_LEVEL_OUTER_X; 1451 } 1452 1453 /* Manually setup the TCS binding table. */ 1454 memset(&bt, 0, sizeof(bt)); 1455 bt.sizes[CROCUS_SURFACE_GROUP_UBO] = 1; 1456 bt.used_mask[CROCUS_SURFACE_GROUP_UBO] = 1; 1457 bt.size_bytes = 4; 1458 1459 prog_data->ubo_ranges[0].length = 1; 1460 } 1461 1462 struct brw_tcs_prog_key key_clean = *key; 1463 crocus_sanitize_tex_key(&key_clean.base.tex); 1464 char *error_str = NULL; 1465 const unsigned *program = 1466 brw_compile_tcs(compiler, &ice->dbg, mem_ctx, &key_clean, tcs_prog_data, nir, 1467 -1, NULL, &error_str); 1468 if (program == NULL) { 1469 dbg_printf("Failed to compile control shader: %s\n", error_str); 1470 ralloc_free(mem_ctx); 1471 return false; 1472 } 1473 1474 if (ish) { 1475 if (ish->compiled_once) { 1476 crocus_debug_recompile(ice, &nir->info, &key->base); 1477 } else { 1478 ish->compiled_once = true; 1479 } 1480 } 1481 1482 struct crocus_compiled_shader *shader = 1483 crocus_upload_shader(ice, CROCUS_CACHE_TCS, sizeof(*key), key, program, 1484 prog_data->program_size, 1485 prog_data, sizeof(*tcs_prog_data), NULL, 1486 system_values, num_system_values, 1487 num_cbufs, &bt); 1488 1489 if (ish) 1490 crocus_disk_cache_store(screen->disk_cache, ish, shader, 1491 ice->shaders.cache_bo_map, 1492 key, sizeof(*key)); 1493 1494 ralloc_free(mem_ctx); 1495 return shader; 1496} 1497 1498/** 1499 * Update the current tessellation control shader variant. 1500 * 1501 * Fill out the key, look in the cache, compile and bind if needed. 1502 */ 1503static void 1504crocus_update_compiled_tcs(struct crocus_context *ice) 1505{ 1506 struct crocus_shader_state *shs = &ice->state.shaders[MESA_SHADER_TESS_CTRL]; 1507 struct crocus_uncompiled_shader *tcs = 1508 ice->shaders.uncompiled[MESA_SHADER_TESS_CTRL]; 1509 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1510 const struct intel_device_info *devinfo = &screen->devinfo; 1511 1512 const struct shader_info *tes_info = 1513 crocus_get_shader_info(ice, MESA_SHADER_TESS_EVAL); 1514 struct brw_tcs_prog_key key = { 1515 KEY_INIT_NO_ID(), 1516 .base.program_string_id = tcs ? tcs->program_id : 0, 1517 .tes_primitive_mode = tes_info->tess.primitive_mode, 1518 .input_vertices = ice->state.vertices_per_patch, 1519 .quads_workaround = tes_info->tess.primitive_mode == GL_QUADS && 1520 tes_info->tess.spacing == TESS_SPACING_EQUAL, 1521 }; 1522 1523 if (tcs && tcs->nos & (1ull << CROCUS_NOS_TEXTURES)) 1524 crocus_populate_sampler_prog_key_data(ice, devinfo, MESA_SHADER_TESS_CTRL, tcs, 1525 tcs->nir->info.uses_texture_gather, &key.base.tex); 1526 get_unified_tess_slots(ice, &key.outputs_written, 1527 &key.patch_outputs_written); 1528 screen->vtbl.populate_tcs_key(ice, &key); 1529 1530 struct crocus_compiled_shader *old = ice->shaders.prog[CROCUS_CACHE_TCS]; 1531 struct crocus_compiled_shader *shader = 1532 crocus_find_cached_shader(ice, CROCUS_CACHE_TCS, sizeof(key), &key); 1533 1534 if (tcs && !shader) 1535 shader = crocus_disk_cache_retrieve(ice, tcs, &key, sizeof(key)); 1536 1537 if (!shader) 1538 shader = crocus_compile_tcs(ice, tcs, &key); 1539 1540 if (old != shader) { 1541 ice->shaders.prog[CROCUS_CACHE_TCS] = shader; 1542 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_TCS | 1543 CROCUS_STAGE_DIRTY_BINDINGS_TCS | 1544 CROCUS_STAGE_DIRTY_CONSTANTS_TCS; 1545 shs->sysvals_need_upload = true; 1546 } 1547} 1548 1549/** 1550 * Compile a tessellation evaluation shader, and upload the assembly. 1551 */ 1552static struct crocus_compiled_shader * 1553crocus_compile_tes(struct crocus_context *ice, 1554 struct crocus_uncompiled_shader *ish, 1555 const struct brw_tes_prog_key *key) 1556{ 1557 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1558 const struct brw_compiler *compiler = screen->compiler; 1559 void *mem_ctx = ralloc_context(NULL); 1560 struct brw_tes_prog_data *tes_prog_data = 1561 rzalloc(mem_ctx, struct brw_tes_prog_data); 1562 struct brw_vue_prog_data *vue_prog_data = &tes_prog_data->base; 1563 struct brw_stage_prog_data *prog_data = &vue_prog_data->base; 1564 enum brw_param_builtin *system_values; 1565 const struct intel_device_info *devinfo = &screen->devinfo; 1566 unsigned num_system_values; 1567 unsigned num_cbufs; 1568 1569 nir_shader *nir = nir_shader_clone(mem_ctx, ish->nir); 1570 1571 if (key->nr_userclip_plane_consts) { 1572 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 1573 nir_lower_clip_vs(nir, (1 << key->nr_userclip_plane_consts) - 1, true, 1574 false, NULL); 1575 nir_lower_io_to_temporaries(nir, impl, true, false); 1576 nir_lower_global_vars_to_local(nir); 1577 nir_lower_vars_to_ssa(nir); 1578 nir_shader_gather_info(nir, impl); 1579 } 1580 1581 crocus_setup_uniforms(compiler, mem_ctx, nir, prog_data, &system_values, 1582 &num_system_values, &num_cbufs); 1583 crocus_lower_swizzles(nir, &key->base.tex); 1584 struct crocus_binding_table bt; 1585 crocus_setup_binding_table(devinfo, nir, &bt, /* num_render_targets */ 0, 1586 num_system_values, num_cbufs, &key->base.tex); 1587 1588 if (can_push_ubo(devinfo)) 1589 brw_nir_analyze_ubo_ranges(compiler, nir, NULL, prog_data->ubo_ranges); 1590 1591 struct brw_vue_map input_vue_map; 1592 brw_compute_tess_vue_map(&input_vue_map, key->inputs_read, 1593 key->patch_inputs_read); 1594 1595 struct brw_tes_prog_key key_clean = *key; 1596 crocus_sanitize_tex_key(&key_clean.base.tex); 1597 char *error_str = NULL; 1598 const unsigned *program = 1599 brw_compile_tes(compiler, &ice->dbg, mem_ctx, &key_clean, &input_vue_map, 1600 tes_prog_data, nir, -1, NULL, &error_str); 1601 if (program == NULL) { 1602 dbg_printf("Failed to compile evaluation shader: %s\n", error_str); 1603 ralloc_free(mem_ctx); 1604 return false; 1605 } 1606 1607 if (ish->compiled_once) { 1608 crocus_debug_recompile(ice, &nir->info, &key->base); 1609 } else { 1610 ish->compiled_once = true; 1611 } 1612 1613 uint32_t *so_decls = NULL; 1614 if (devinfo->ver > 6) 1615 so_decls = screen->vtbl.create_so_decl_list(&ish->stream_output, 1616 &vue_prog_data->vue_map); 1617 1618 struct crocus_compiled_shader *shader = 1619 crocus_upload_shader(ice, CROCUS_CACHE_TES, sizeof(*key), key, program, 1620 prog_data->program_size, 1621 prog_data, sizeof(*tes_prog_data), so_decls, 1622 system_values, num_system_values, 1623 num_cbufs, &bt); 1624 1625 crocus_disk_cache_store(screen->disk_cache, ish, shader, 1626 ice->shaders.cache_bo_map, 1627 key, sizeof(*key)); 1628 1629 ralloc_free(mem_ctx); 1630 return shader; 1631} 1632 1633/** 1634 * Update the current tessellation evaluation shader variant. 1635 * 1636 * Fill out the key, look in the cache, compile and bind if needed. 1637 */ 1638static void 1639crocus_update_compiled_tes(struct crocus_context *ice) 1640{ 1641 struct crocus_shader_state *shs = &ice->state.shaders[MESA_SHADER_TESS_EVAL]; 1642 struct crocus_uncompiled_shader *ish = 1643 ice->shaders.uncompiled[MESA_SHADER_TESS_EVAL]; 1644 struct brw_tes_prog_key key = { KEY_INIT() }; 1645 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1646 const struct intel_device_info *devinfo = &screen->devinfo; 1647 1648 if (ish->nos & (1ull << CROCUS_NOS_TEXTURES)) 1649 crocus_populate_sampler_prog_key_data(ice, devinfo, MESA_SHADER_TESS_EVAL, ish, 1650 ish->nir->info.uses_texture_gather, &key.base.tex); 1651 get_unified_tess_slots(ice, &key.inputs_read, &key.patch_inputs_read); 1652 screen->vtbl.populate_tes_key(ice, &ish->nir->info, last_vue_stage(ice), &key); 1653 1654 struct crocus_compiled_shader *old = ice->shaders.prog[CROCUS_CACHE_TES]; 1655 struct crocus_compiled_shader *shader = 1656 crocus_find_cached_shader(ice, CROCUS_CACHE_TES, sizeof(key), &key); 1657 1658 if (!shader) 1659 shader = crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key)); 1660 1661 if (!shader) 1662 shader = crocus_compile_tes(ice, ish, &key); 1663 1664 if (old != shader) { 1665 ice->shaders.prog[CROCUS_CACHE_TES] = shader; 1666 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_TES | 1667 CROCUS_STAGE_DIRTY_BINDINGS_TES | 1668 CROCUS_STAGE_DIRTY_CONSTANTS_TES; 1669 shs->sysvals_need_upload = true; 1670 } 1671 1672 /* TODO: Could compare and avoid flagging this. */ 1673 const struct shader_info *tes_info = &ish->nir->info; 1674 if (BITSET_TEST(tes_info->system_values_read, SYSTEM_VALUE_VERTICES_IN)) { 1675 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_CONSTANTS_TES; 1676 ice->state.shaders[MESA_SHADER_TESS_EVAL].sysvals_need_upload = true; 1677 } 1678} 1679 1680/** 1681 * Compile a geometry shader, and upload the assembly. 1682 */ 1683static struct crocus_compiled_shader * 1684crocus_compile_gs(struct crocus_context *ice, 1685 struct crocus_uncompiled_shader *ish, 1686 const struct brw_gs_prog_key *key) 1687{ 1688 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1689 const struct brw_compiler *compiler = screen->compiler; 1690 const struct intel_device_info *devinfo = &screen->devinfo; 1691 void *mem_ctx = ralloc_context(NULL); 1692 struct brw_gs_prog_data *gs_prog_data = 1693 rzalloc(mem_ctx, struct brw_gs_prog_data); 1694 struct brw_vue_prog_data *vue_prog_data = &gs_prog_data->base; 1695 struct brw_stage_prog_data *prog_data = &vue_prog_data->base; 1696 enum brw_param_builtin *system_values; 1697 unsigned num_system_values; 1698 unsigned num_cbufs; 1699 1700 nir_shader *nir = nir_shader_clone(mem_ctx, ish->nir); 1701 1702 if (key->nr_userclip_plane_consts) { 1703 nir_function_impl *impl = nir_shader_get_entrypoint(nir); 1704 nir_lower_clip_gs(nir, (1 << key->nr_userclip_plane_consts) - 1, false, 1705 NULL); 1706 nir_lower_io_to_temporaries(nir, impl, true, false); 1707 nir_lower_global_vars_to_local(nir); 1708 nir_lower_vars_to_ssa(nir); 1709 nir_shader_gather_info(nir, impl); 1710 } 1711 1712 crocus_setup_uniforms(compiler, mem_ctx, nir, prog_data, &system_values, 1713 &num_system_values, &num_cbufs); 1714 crocus_lower_swizzles(nir, &key->base.tex); 1715 struct crocus_binding_table bt; 1716 crocus_setup_binding_table(devinfo, nir, &bt, /* num_render_targets */ 0, 1717 num_system_values, num_cbufs, &key->base.tex); 1718 1719 if (can_push_ubo(devinfo)) 1720 brw_nir_analyze_ubo_ranges(compiler, nir, NULL, prog_data->ubo_ranges); 1721 1722 brw_compute_vue_map(devinfo, 1723 &vue_prog_data->vue_map, nir->info.outputs_written, 1724 nir->info.separate_shader, /* pos slots */ 1); 1725 1726 if (devinfo->ver == 6) 1727 gfx6_gs_xfb_setup(&ish->stream_output, gs_prog_data); 1728 struct brw_gs_prog_key key_clean = *key; 1729 crocus_sanitize_tex_key(&key_clean.base.tex); 1730 1731 char *error_str = NULL; 1732 const unsigned *program = 1733 brw_compile_gs(compiler, &ice->dbg, mem_ctx, &key_clean, gs_prog_data, nir, 1734 -1, NULL, &error_str); 1735 if (program == NULL) { 1736 dbg_printf("Failed to compile geometry shader: %s\n", error_str); 1737 ralloc_free(mem_ctx); 1738 return false; 1739 } 1740 1741 if (ish->compiled_once) { 1742 crocus_debug_recompile(ice, &nir->info, &key->base); 1743 } else { 1744 ish->compiled_once = true; 1745 } 1746 1747 uint32_t *so_decls = NULL; 1748 if (devinfo->ver > 6) 1749 so_decls = screen->vtbl.create_so_decl_list(&ish->stream_output, 1750 &vue_prog_data->vue_map); 1751 1752 struct crocus_compiled_shader *shader = 1753 crocus_upload_shader(ice, CROCUS_CACHE_GS, sizeof(*key), key, program, 1754 prog_data->program_size, 1755 prog_data, sizeof(*gs_prog_data), so_decls, 1756 system_values, num_system_values, 1757 num_cbufs, &bt); 1758 1759 crocus_disk_cache_store(screen->disk_cache, ish, shader, 1760 ice->shaders.cache_bo_map, 1761 key, sizeof(*key)); 1762 1763 ralloc_free(mem_ctx); 1764 return shader; 1765} 1766 1767/** 1768 * Update the current geometry shader variant. 1769 * 1770 * Fill out the key, look in the cache, compile and bind if needed. 1771 */ 1772static void 1773crocus_update_compiled_gs(struct crocus_context *ice) 1774{ 1775 struct crocus_shader_state *shs = &ice->state.shaders[MESA_SHADER_GEOMETRY]; 1776 struct crocus_uncompiled_shader *ish = 1777 ice->shaders.uncompiled[MESA_SHADER_GEOMETRY]; 1778 struct crocus_compiled_shader *old = ice->shaders.prog[CROCUS_CACHE_GS]; 1779 struct crocus_compiled_shader *shader = NULL; 1780 1781 if (ish) { 1782 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1783 const struct intel_device_info *devinfo = &screen->devinfo; 1784 struct brw_gs_prog_key key = { KEY_INIT() }; 1785 1786 if (ish->nos & (1ull << CROCUS_NOS_TEXTURES)) 1787 crocus_populate_sampler_prog_key_data(ice, devinfo, MESA_SHADER_GEOMETRY, ish, 1788 ish->nir->info.uses_texture_gather, &key.base.tex); 1789 screen->vtbl.populate_gs_key(ice, &ish->nir->info, last_vue_stage(ice), &key); 1790 1791 shader = 1792 crocus_find_cached_shader(ice, CROCUS_CACHE_GS, sizeof(key), &key); 1793 1794 if (!shader) 1795 shader = crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key)); 1796 1797 if (!shader) 1798 shader = crocus_compile_gs(ice, ish, &key); 1799 } 1800 1801 if (old != shader) { 1802 ice->shaders.prog[CROCUS_CACHE_GS] = shader; 1803 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_GS | 1804 CROCUS_STAGE_DIRTY_BINDINGS_GS | 1805 CROCUS_STAGE_DIRTY_CONSTANTS_GS; 1806 shs->sysvals_need_upload = true; 1807 } 1808} 1809 1810/** 1811 * Compile a fragment (pixel) shader, and upload the assembly. 1812 */ 1813static struct crocus_compiled_shader * 1814crocus_compile_fs(struct crocus_context *ice, 1815 struct crocus_uncompiled_shader *ish, 1816 const struct brw_wm_prog_key *key, 1817 struct brw_vue_map *vue_map) 1818{ 1819 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1820 const struct brw_compiler *compiler = screen->compiler; 1821 void *mem_ctx = ralloc_context(NULL); 1822 struct brw_wm_prog_data *fs_prog_data = 1823 rzalloc(mem_ctx, struct brw_wm_prog_data); 1824 struct brw_stage_prog_data *prog_data = &fs_prog_data->base; 1825 enum brw_param_builtin *system_values; 1826 const struct intel_device_info *devinfo = &screen->devinfo; 1827 unsigned num_system_values; 1828 unsigned num_cbufs; 1829 1830 nir_shader *nir = nir_shader_clone(mem_ctx, ish->nir); 1831 1832 prog_data->use_alt_mode = nir->info.is_arb_asm; 1833 1834 crocus_setup_uniforms(compiler, mem_ctx, nir, prog_data, &system_values, 1835 &num_system_values, &num_cbufs); 1836 1837 /* Lower output variables to load_output intrinsics before setting up 1838 * binding tables, so crocus_setup_binding_table can map any load_output 1839 * intrinsics to CROCUS_SURFACE_GROUP_RENDER_TARGET_READ on Gen8 for 1840 * non-coherent framebuffer fetches. 1841 */ 1842 brw_nir_lower_fs_outputs(nir); 1843 1844 /* lower swizzles before binding table */ 1845 crocus_lower_swizzles(nir, &key->base.tex); 1846 int null_rts = 1; 1847 1848 struct crocus_binding_table bt; 1849 crocus_setup_binding_table(devinfo, nir, &bt, 1850 MAX2(key->nr_color_regions, null_rts), 1851 num_system_values, num_cbufs, 1852 &key->base.tex); 1853 1854 if (can_push_ubo(devinfo)) 1855 brw_nir_analyze_ubo_ranges(compiler, nir, NULL, prog_data->ubo_ranges); 1856 1857 struct brw_wm_prog_key key_clean = *key; 1858 crocus_sanitize_tex_key(&key_clean.base.tex); 1859 1860 struct brw_compile_fs_params params = { 1861 .nir = nir, 1862 .key = &key_clean, 1863 .prog_data = fs_prog_data, 1864 1865 .allow_spilling = true, 1866 .vue_map = vue_map, 1867 1868 .log_data = &ice->dbg, 1869 }; 1870 const unsigned *program = 1871 brw_compile_fs(compiler, mem_ctx, ¶ms); 1872 if (program == NULL) { 1873 dbg_printf("Failed to compile fragment shader: %s\n", params.error_str); 1874 ralloc_free(mem_ctx); 1875 return false; 1876 } 1877 1878 if (ish->compiled_once) { 1879 crocus_debug_recompile(ice, &nir->info, &key->base); 1880 } else { 1881 ish->compiled_once = true; 1882 } 1883 1884 struct crocus_compiled_shader *shader = 1885 crocus_upload_shader(ice, CROCUS_CACHE_FS, sizeof(*key), key, program, 1886 prog_data->program_size, 1887 prog_data, sizeof(*fs_prog_data), NULL, 1888 system_values, num_system_values, 1889 num_cbufs, &bt); 1890 1891 crocus_disk_cache_store(screen->disk_cache, ish, shader, 1892 ice->shaders.cache_bo_map, 1893 key, sizeof(*key)); 1894 1895 ralloc_free(mem_ctx); 1896 return shader; 1897} 1898 1899/** 1900 * Update the current fragment shader variant. 1901 * 1902 * Fill out the key, look in the cache, compile and bind if needed. 1903 */ 1904static void 1905crocus_update_compiled_fs(struct crocus_context *ice) 1906{ 1907 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1908 const struct intel_device_info *devinfo = &screen->devinfo; 1909 struct crocus_shader_state *shs = &ice->state.shaders[MESA_SHADER_FRAGMENT]; 1910 struct crocus_uncompiled_shader *ish = 1911 ice->shaders.uncompiled[MESA_SHADER_FRAGMENT]; 1912 struct brw_wm_prog_key key = { KEY_INIT() }; 1913 1914 if (ish->nos & (1ull << CROCUS_NOS_TEXTURES)) 1915 crocus_populate_sampler_prog_key_data(ice, devinfo, MESA_SHADER_FRAGMENT, ish, 1916 ish->nir->info.uses_texture_gather, &key.base.tex); 1917 screen->vtbl.populate_fs_key(ice, &ish->nir->info, &key); 1918 1919 if (ish->nos & (1ull << CROCUS_NOS_LAST_VUE_MAP)) 1920 key.input_slots_valid = ice->shaders.last_vue_map->slots_valid; 1921 1922 struct crocus_compiled_shader *old = ice->shaders.prog[CROCUS_CACHE_FS]; 1923 struct crocus_compiled_shader *shader = 1924 crocus_find_cached_shader(ice, CROCUS_CACHE_FS, sizeof(key), &key); 1925 1926 if (!shader) 1927 shader = crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key)); 1928 1929 if (!shader) 1930 shader = crocus_compile_fs(ice, ish, &key, ice->shaders.last_vue_map); 1931 1932 if (old != shader) { 1933 // XXX: only need to flag CLIP if barycentric has NONPERSPECTIVE 1934 // toggles. might be able to avoid flagging SBE too. 1935 ice->shaders.prog[CROCUS_CACHE_FS] = shader; 1936 ice->state.dirty |= CROCUS_DIRTY_WM; 1937 /* gen4 clip/sf rely on fs prog_data */ 1938 if (devinfo->ver < 6) 1939 ice->state.dirty |= CROCUS_DIRTY_GEN4_CLIP_PROG | CROCUS_DIRTY_GEN4_SF_PROG; 1940 else 1941 ice->state.dirty |= CROCUS_DIRTY_CLIP | CROCUS_DIRTY_GEN6_BLEND_STATE; 1942 if (devinfo->ver == 6) 1943 ice->state.dirty |= CROCUS_DIRTY_RASTER; 1944 if (devinfo->ver >= 7) 1945 ice->state.dirty |= CROCUS_DIRTY_GEN7_SBE; 1946 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_FS | 1947 CROCUS_STAGE_DIRTY_BINDINGS_FS | 1948 CROCUS_STAGE_DIRTY_CONSTANTS_FS; 1949 shs->sysvals_need_upload = true; 1950 } 1951} 1952 1953/** 1954 * Update the last enabled stage's VUE map. 1955 * 1956 * When the shader feeding the rasterizer's output interface changes, we 1957 * need to re-emit various packets. 1958 */ 1959static void 1960update_last_vue_map(struct crocus_context *ice, 1961 struct brw_stage_prog_data *prog_data) 1962{ 1963 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 1964 const struct intel_device_info *devinfo = &screen->devinfo; 1965 struct brw_vue_prog_data *vue_prog_data = (void *) prog_data; 1966 struct brw_vue_map *vue_map = &vue_prog_data->vue_map; 1967 struct brw_vue_map *old_map = ice->shaders.last_vue_map; 1968 const uint64_t changed_slots = 1969 (old_map ? old_map->slots_valid : 0ull) ^ vue_map->slots_valid; 1970 1971 if (changed_slots & VARYING_BIT_VIEWPORT) { 1972 ice->state.num_viewports = 1973 (vue_map->slots_valid & VARYING_BIT_VIEWPORT) ? CROCUS_MAX_VIEWPORTS : 1; 1974 ice->state.dirty |= CROCUS_DIRTY_SF_CL_VIEWPORT | 1975 CROCUS_DIRTY_CC_VIEWPORT; 1976 if (devinfo->ver < 6) 1977 ice->state.dirty |= CROCUS_DIRTY_GEN4_CLIP_PROG | CROCUS_DIRTY_GEN4_SF_PROG; 1978 1979 if (devinfo->ver <= 6) 1980 ice->state.dirty |= CROCUS_DIRTY_GEN4_FF_GS_PROG; 1981 1982 if (devinfo->ver >= 6) 1983 ice->state.dirty |= CROCUS_DIRTY_CLIP | 1984 CROCUS_DIRTY_GEN6_SCISSOR_RECT;; 1985 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_UNCOMPILED_FS | 1986 ice->state.stage_dirty_for_nos[CROCUS_NOS_LAST_VUE_MAP]; 1987 } 1988 1989 if (changed_slots || (old_map && old_map->separate != vue_map->separate)) { 1990 ice->state.dirty |= CROCUS_DIRTY_GEN7_SBE; 1991 if (devinfo->ver < 6) 1992 ice->state.dirty |= CROCUS_DIRTY_GEN4_FF_GS_PROG; 1993 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_UNCOMPILED_FS; 1994 } 1995 1996 ice->shaders.last_vue_map = &vue_prog_data->vue_map; 1997} 1998 1999static void 2000crocus_update_pull_constant_descriptors(struct crocus_context *ice, 2001 gl_shader_stage stage) 2002{ 2003 struct crocus_compiled_shader *shader = ice->shaders.prog[stage]; 2004 2005 if (!shader || !shader->prog_data->has_ubo_pull) 2006 return; 2007 2008 struct crocus_shader_state *shs = &ice->state.shaders[stage]; 2009 bool any_new_descriptors = 2010 shader->num_system_values > 0 && shs->sysvals_need_upload; 2011 2012 unsigned bound_cbufs = shs->bound_cbufs; 2013 2014 while (bound_cbufs) { 2015 const int i = u_bit_scan(&bound_cbufs); 2016 struct pipe_constant_buffer *cbuf = &shs->constbufs[i]; 2017 if (cbuf->buffer) { 2018 any_new_descriptors = true; 2019 } 2020 } 2021 2022 if (any_new_descriptors) 2023 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_BINDINGS_VS << stage; 2024} 2025 2026/** 2027 * Get the prog_data for a given stage, or NULL if the stage is disabled. 2028 */ 2029static struct brw_vue_prog_data * 2030get_vue_prog_data(struct crocus_context *ice, gl_shader_stage stage) 2031{ 2032 if (!ice->shaders.prog[stage]) 2033 return NULL; 2034 2035 return (void *) ice->shaders.prog[stage]->prog_data; 2036} 2037 2038static struct crocus_compiled_shader * 2039crocus_compile_clip(struct crocus_context *ice, struct brw_clip_prog_key *key) 2040{ 2041 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2042 const struct brw_compiler *compiler = screen->compiler; 2043 void *mem_ctx; 2044 unsigned program_size; 2045 mem_ctx = ralloc_context(NULL); 2046 2047 struct brw_clip_prog_data *clip_prog_data = 2048 rzalloc(mem_ctx, struct brw_clip_prog_data); 2049 2050 const unsigned *program = brw_compile_clip(compiler, mem_ctx, key, clip_prog_data, 2051 ice->shaders.last_vue_map, &program_size); 2052 2053 if (program == NULL) { 2054 dbg_printf("failed to compile clip shader\n"); 2055 ralloc_free(mem_ctx); 2056 return false; 2057 } 2058 struct crocus_binding_table bt; 2059 memset(&bt, 0, sizeof(bt)); 2060 2061 struct crocus_compiled_shader *shader = 2062 crocus_upload_shader(ice, CROCUS_CACHE_CLIP, sizeof(*key), key, program, 2063 program_size, 2064 (struct brw_stage_prog_data *)clip_prog_data, sizeof(*clip_prog_data), 2065 NULL, NULL, 0, 0, &bt); 2066 ralloc_free(mem_ctx); 2067 return shader; 2068} 2069static void 2070crocus_update_compiled_clip(struct crocus_context *ice) 2071{ 2072 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2073 struct brw_clip_prog_key key; 2074 struct crocus_compiled_shader *old = ice->shaders.clip_prog; 2075 memset(&key, 0, sizeof(key)); 2076 2077 const struct brw_wm_prog_data *wm_prog_data = brw_wm_prog_data(ice->shaders.prog[MESA_SHADER_FRAGMENT]->prog_data); 2078 if (wm_prog_data) { 2079 key.contains_flat_varying = wm_prog_data->contains_flat_varying; 2080 key.contains_noperspective_varying = 2081 wm_prog_data->contains_noperspective_varying; 2082 memcpy(key.interp_mode, wm_prog_data->interp_mode, sizeof(key.interp_mode)); 2083 } 2084 2085 key.primitive = ice->state.reduced_prim_mode; 2086 key.attrs = ice->shaders.last_vue_map->slots_valid; 2087 2088 struct pipe_rasterizer_state *rs_state = crocus_get_rast_state(ice); 2089 key.pv_first = rs_state->flatshade_first; 2090 2091 if (rs_state->clip_plane_enable) 2092 key.nr_userclip = util_logbase2(rs_state->clip_plane_enable) + 1; 2093 2094 if (screen->devinfo.ver == 5) 2095 key.clip_mode = BRW_CLIP_MODE_KERNEL_CLIP; 2096 else 2097 key.clip_mode = BRW_CLIP_MODE_NORMAL; 2098 2099 if (key.primitive == PIPE_PRIM_TRIANGLES) { 2100 if (rs_state->cull_face == PIPE_FACE_FRONT_AND_BACK) 2101 key.clip_mode = BRW_CLIP_MODE_REJECT_ALL; 2102 else { 2103 uint32_t fill_front = BRW_CLIP_FILL_MODE_CULL; 2104 uint32_t fill_back = BRW_CLIP_FILL_MODE_CULL; 2105 uint32_t offset_front = 0; 2106 uint32_t offset_back = 0; 2107 2108 if (!(rs_state->cull_face & PIPE_FACE_FRONT)) { 2109 switch (rs_state->fill_front) { 2110 case PIPE_POLYGON_MODE_FILL: 2111 fill_front = BRW_CLIP_FILL_MODE_FILL; 2112 offset_front = 0; 2113 break; 2114 case PIPE_POLYGON_MODE_LINE: 2115 fill_front = BRW_CLIP_FILL_MODE_LINE; 2116 offset_front = rs_state->offset_line; 2117 break; 2118 case PIPE_POLYGON_MODE_POINT: 2119 fill_front = BRW_CLIP_FILL_MODE_POINT; 2120 offset_front = rs_state->offset_point; 2121 break; 2122 } 2123 } 2124 2125 if (!(rs_state->cull_face & PIPE_FACE_BACK)) { 2126 switch (rs_state->fill_back) { 2127 case PIPE_POLYGON_MODE_FILL: 2128 fill_back = BRW_CLIP_FILL_MODE_FILL; 2129 offset_back = 0; 2130 break; 2131 case PIPE_POLYGON_MODE_LINE: 2132 fill_back = BRW_CLIP_FILL_MODE_LINE; 2133 offset_back = rs_state->offset_line; 2134 break; 2135 case PIPE_POLYGON_MODE_POINT: 2136 fill_back = BRW_CLIP_FILL_MODE_POINT; 2137 offset_back = rs_state->offset_point; 2138 break; 2139 } 2140 } 2141 2142 if (rs_state->fill_back != PIPE_POLYGON_MODE_FILL || 2143 rs_state->fill_front != PIPE_POLYGON_MODE_FILL) { 2144 key.do_unfilled = 1; 2145 2146 /* Most cases the fixed function units will handle. Cases where 2147 * one or more polygon faces are unfilled will require help: 2148 */ 2149 key.clip_mode = BRW_CLIP_MODE_CLIP_NON_REJECTED; 2150 2151 if (offset_back || offset_front) { 2152 double mrd = 0.0; 2153 if (ice->state.framebuffer.zsbuf) 2154 mrd = util_get_depth_format_mrd(util_format_description(ice->state.framebuffer.zsbuf->format)); 2155 key.offset_units = rs_state->offset_units * mrd * 2; 2156 key.offset_factor = rs_state->offset_scale * mrd; 2157 key.offset_clamp = rs_state->offset_clamp * mrd; 2158 } 2159 2160 if (!(rs_state->front_ccw ^ rs_state->bottom_edge_rule)) { 2161 key.fill_ccw = fill_front; 2162 key.fill_cw = fill_back; 2163 key.offset_ccw = offset_front; 2164 key.offset_cw = offset_back; 2165 if (rs_state->light_twoside && 2166 key.fill_cw != BRW_CLIP_FILL_MODE_CULL) 2167 key.copy_bfc_cw = 1; 2168 } else { 2169 key.fill_cw = fill_front; 2170 key.fill_ccw = fill_back; 2171 key.offset_cw = offset_front; 2172 key.offset_ccw = offset_back; 2173 if (rs_state->light_twoside && 2174 key.fill_ccw != BRW_CLIP_FILL_MODE_CULL) 2175 key.copy_bfc_ccw = 1; 2176 } 2177 } 2178 } 2179 } 2180 struct crocus_compiled_shader *shader = 2181 crocus_find_cached_shader(ice, CROCUS_CACHE_CLIP, sizeof(key), &key); 2182 2183 if (!shader) 2184 shader = crocus_compile_clip(ice, &key); 2185 2186 if (old != shader) { 2187 ice->state.dirty |= CROCUS_DIRTY_CLIP; 2188 ice->shaders.clip_prog = shader; 2189 } 2190} 2191 2192static struct crocus_compiled_shader * 2193crocus_compile_sf(struct crocus_context *ice, struct brw_sf_prog_key *key) 2194{ 2195 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2196 const struct brw_compiler *compiler = screen->compiler; 2197 void *mem_ctx; 2198 unsigned program_size; 2199 mem_ctx = ralloc_context(NULL); 2200 2201 struct brw_sf_prog_data *sf_prog_data = 2202 rzalloc(mem_ctx, struct brw_sf_prog_data); 2203 2204 const unsigned *program = brw_compile_sf(compiler, mem_ctx, key, sf_prog_data, 2205 ice->shaders.last_vue_map, &program_size); 2206 2207 if (program == NULL) { 2208 dbg_printf("failed to compile sf shader\n"); 2209 ralloc_free(mem_ctx); 2210 return false; 2211 } 2212 2213 struct crocus_binding_table bt; 2214 memset(&bt, 0, sizeof(bt)); 2215 struct crocus_compiled_shader *shader = 2216 crocus_upload_shader(ice, CROCUS_CACHE_SF, sizeof(*key), key, program, 2217 program_size, 2218 (struct brw_stage_prog_data *)sf_prog_data, sizeof(*sf_prog_data), 2219 NULL, NULL, 0, 0, &bt); 2220 ralloc_free(mem_ctx); 2221 return shader; 2222} 2223 2224static void 2225crocus_update_compiled_sf(struct crocus_context *ice) 2226{ 2227 struct brw_sf_prog_key key; 2228 struct crocus_compiled_shader *old = ice->shaders.sf_prog; 2229 memset(&key, 0, sizeof(key)); 2230 2231 key.attrs = ice->shaders.last_vue_map->slots_valid; 2232 2233 switch (ice->state.reduced_prim_mode) { 2234 case GL_TRIANGLES: 2235 default: 2236 if (key.attrs & BITFIELD64_BIT(VARYING_SLOT_EDGE)) 2237 key.primitive = BRW_SF_PRIM_UNFILLED_TRIS; 2238 else 2239 key.primitive = BRW_SF_PRIM_TRIANGLES; 2240 break; 2241 case GL_LINES: 2242 key.primitive = BRW_SF_PRIM_LINES; 2243 break; 2244 case GL_POINTS: 2245 key.primitive = BRW_SF_PRIM_POINTS; 2246 break; 2247 } 2248 2249 struct pipe_rasterizer_state *rs_state = crocus_get_rast_state(ice); 2250 key.userclip_active = rs_state->clip_plane_enable != 0; 2251 const struct brw_wm_prog_data *wm_prog_data = brw_wm_prog_data(ice->shaders.prog[MESA_SHADER_FRAGMENT]->prog_data); 2252 if (wm_prog_data) { 2253 key.contains_flat_varying = wm_prog_data->contains_flat_varying; 2254 memcpy(key.interp_mode, wm_prog_data->interp_mode, sizeof(key.interp_mode)); 2255 } 2256 2257 key.do_twoside_color = rs_state->light_twoside; 2258 2259 key.do_point_sprite = rs_state->point_quad_rasterization; 2260 if (key.do_point_sprite) { 2261 key.point_sprite_coord_replace = rs_state->sprite_coord_enable & 0xff; 2262 if (rs_state->sprite_coord_enable & (1 << 8)) 2263 key.do_point_coord = 1; 2264 if (wm_prog_data && wm_prog_data->urb_setup[VARYING_SLOT_PNTC] != -1) 2265 key.do_point_coord = 1; 2266 } 2267 2268 key.sprite_origin_lower_left = rs_state->sprite_coord_mode == PIPE_SPRITE_COORD_LOWER_LEFT; 2269 2270 if (key.do_twoside_color) { 2271 key.frontface_ccw = rs_state->front_ccw; 2272 } 2273 struct crocus_compiled_shader *shader = 2274 crocus_find_cached_shader(ice, CROCUS_CACHE_SF, sizeof(key), &key); 2275 2276 if (!shader) 2277 shader = crocus_compile_sf(ice, &key); 2278 2279 if (old != shader) { 2280 ice->state.dirty |= CROCUS_DIRTY_RASTER; 2281 ice->shaders.sf_prog = shader; 2282 } 2283} 2284 2285static struct crocus_compiled_shader * 2286crocus_compile_ff_gs(struct crocus_context *ice, struct brw_ff_gs_prog_key *key) 2287{ 2288 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2289 struct brw_compiler *compiler = screen->compiler; 2290 void *mem_ctx; 2291 unsigned program_size; 2292 mem_ctx = ralloc_context(NULL); 2293 2294 struct brw_ff_gs_prog_data *ff_gs_prog_data = 2295 rzalloc(mem_ctx, struct brw_ff_gs_prog_data); 2296 2297 const unsigned *program = brw_compile_ff_gs_prog(compiler, mem_ctx, key, ff_gs_prog_data, 2298 ice->shaders.last_vue_map, &program_size); 2299 2300 if (program == NULL) { 2301 dbg_printf("failed to compile sf shader\n"); 2302 ralloc_free(mem_ctx); 2303 return false; 2304 } 2305 2306 struct crocus_binding_table bt; 2307 memset(&bt, 0, sizeof(bt)); 2308 2309 if (screen->devinfo.ver == 6) { 2310 bt.sizes[CROCUS_SURFACE_GROUP_SOL] = BRW_MAX_SOL_BINDINGS; 2311 bt.used_mask[CROCUS_SURFACE_GROUP_SOL] = (uint64_t)-1; 2312 2313 bt.size_bytes = BRW_MAX_SOL_BINDINGS * 4; 2314 } 2315 2316 struct crocus_compiled_shader *shader = 2317 crocus_upload_shader(ice, CROCUS_CACHE_FF_GS, sizeof(*key), key, program, 2318 program_size, 2319 (struct brw_stage_prog_data *)ff_gs_prog_data, sizeof(*ff_gs_prog_data), 2320 NULL, NULL, 0, 0, &bt); 2321 ralloc_free(mem_ctx); 2322 return shader; 2323} 2324 2325static void 2326crocus_update_compiled_ff_gs(struct crocus_context *ice) 2327{ 2328 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2329 const struct intel_device_info *devinfo = &screen->devinfo; 2330 struct brw_ff_gs_prog_key key; 2331 struct crocus_compiled_shader *old = ice->shaders.ff_gs_prog; 2332 memset(&key, 0, sizeof(key)); 2333 2334 assert(devinfo->ver < 7); 2335 2336 key.attrs = ice->shaders.last_vue_map->slots_valid; 2337 2338 key.primitive = screen->vtbl.translate_prim_type(ice->state.prim_mode, 0); 2339 2340 struct pipe_rasterizer_state *rs_state = crocus_get_rast_state(ice); 2341 key.pv_first = rs_state->flatshade_first; 2342 2343 if (key.primitive == _3DPRIM_QUADLIST && !rs_state->flatshade) { 2344 /* Provide consistenbbbbbt primitive order with brw_set_prim's 2345 * optimization of single quads to trifans. 2346 */ 2347 key.pv_first = true; 2348 } 2349 2350 if (devinfo->ver >= 6) { 2351 key.need_gs_prog = ice->state.streamout_active; 2352 if (key.need_gs_prog) { 2353 struct crocus_uncompiled_shader *vs = 2354 ice->shaders.uncompiled[MESA_SHADER_VERTEX]; 2355 gfx6_ff_gs_xfb_setup(&vs->stream_output, 2356 &key); 2357 } 2358 } else { 2359 key.need_gs_prog = (key.primitive == _3DPRIM_QUADLIST || 2360 key.primitive == _3DPRIM_QUADSTRIP || 2361 key.primitive == _3DPRIM_LINELOOP); 2362 } 2363 2364 struct crocus_compiled_shader *shader = NULL; 2365 if (key.need_gs_prog) { 2366 shader = crocus_find_cached_shader(ice, CROCUS_CACHE_FF_GS, 2367 sizeof(key), &key); 2368 if (!shader) 2369 shader = crocus_compile_ff_gs(ice, &key); 2370 } 2371 if (old != shader) { 2372 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_GS; 2373 if (!!old != !!shader) 2374 ice->state.dirty |= CROCUS_DIRTY_GEN6_URB; 2375 ice->shaders.ff_gs_prog = shader; 2376 if (shader) { 2377 const struct brw_ff_gs_prog_data *gs_prog_data = (struct brw_ff_gs_prog_data *)ice->shaders.ff_gs_prog->prog_data; 2378 ice->state.last_xfb_verts_per_prim = gs_prog_data->svbi_postincrement_value; 2379 } 2380 } 2381} 2382 2383// XXX: crocus_compiled_shaders are space-leaking :( 2384// XXX: do remember to unbind them if deleting them. 2385 2386/** 2387 * Update the current shader variants for the given state. 2388 * 2389 * This should be called on every draw call to ensure that the correct 2390 * shaders are bound. It will also flag any dirty state triggered by 2391 * swapping out those shaders. 2392 */ 2393bool 2394crocus_update_compiled_shaders(struct crocus_context *ice) 2395{ 2396 struct crocus_screen *screen = (void *) ice->ctx.screen; 2397 const uint64_t stage_dirty = ice->state.stage_dirty; 2398 2399 struct brw_vue_prog_data *old_prog_datas[4]; 2400 if (!(ice->state.dirty & CROCUS_DIRTY_GEN6_URB)) { 2401 for (int i = MESA_SHADER_VERTEX; i <= MESA_SHADER_GEOMETRY; i++) 2402 old_prog_datas[i] = get_vue_prog_data(ice, i); 2403 } 2404 2405 if (stage_dirty & (CROCUS_STAGE_DIRTY_UNCOMPILED_TCS | 2406 CROCUS_STAGE_DIRTY_UNCOMPILED_TES)) { 2407 struct crocus_uncompiled_shader *tes = 2408 ice->shaders.uncompiled[MESA_SHADER_TESS_EVAL]; 2409 if (tes) { 2410 crocus_update_compiled_tcs(ice); 2411 crocus_update_compiled_tes(ice); 2412 } else { 2413 ice->shaders.prog[CROCUS_CACHE_TCS] = NULL; 2414 ice->shaders.prog[CROCUS_CACHE_TES] = NULL; 2415 ice->state.stage_dirty |= 2416 CROCUS_STAGE_DIRTY_TCS | CROCUS_STAGE_DIRTY_TES | 2417 CROCUS_STAGE_DIRTY_BINDINGS_TCS | CROCUS_STAGE_DIRTY_BINDINGS_TES | 2418 CROCUS_STAGE_DIRTY_CONSTANTS_TCS | CROCUS_STAGE_DIRTY_CONSTANTS_TES; 2419 } 2420 } 2421 2422 if (stage_dirty & CROCUS_STAGE_DIRTY_UNCOMPILED_VS) 2423 crocus_update_compiled_vs(ice); 2424 if (stage_dirty & CROCUS_STAGE_DIRTY_UNCOMPILED_GS) 2425 crocus_update_compiled_gs(ice); 2426 2427 if (stage_dirty & (CROCUS_STAGE_DIRTY_UNCOMPILED_GS | 2428 CROCUS_STAGE_DIRTY_UNCOMPILED_TES)) { 2429 const struct crocus_compiled_shader *gs = 2430 ice->shaders.prog[MESA_SHADER_GEOMETRY]; 2431 const struct crocus_compiled_shader *tes = 2432 ice->shaders.prog[MESA_SHADER_TESS_EVAL]; 2433 2434 bool points_or_lines = false; 2435 2436 if (gs) { 2437 const struct brw_gs_prog_data *gs_prog_data = (void *) gs->prog_data; 2438 points_or_lines = 2439 gs_prog_data->output_topology == _3DPRIM_POINTLIST || 2440 gs_prog_data->output_topology == _3DPRIM_LINESTRIP; 2441 } else if (tes) { 2442 const struct brw_tes_prog_data *tes_data = (void *) tes->prog_data; 2443 points_or_lines = 2444 tes_data->output_topology == BRW_TESS_OUTPUT_TOPOLOGY_LINE || 2445 tes_data->output_topology == BRW_TESS_OUTPUT_TOPOLOGY_POINT; 2446 } 2447 2448 if (ice->shaders.output_topology_is_points_or_lines != points_or_lines) { 2449 /* Outbound to XY Clip enables */ 2450 ice->shaders.output_topology_is_points_or_lines = points_or_lines; 2451 ice->state.dirty |= CROCUS_DIRTY_CLIP; 2452 } 2453 } 2454 2455 if (!ice->shaders.prog[MESA_SHADER_VERTEX]) 2456 return false; 2457 2458 gl_shader_stage last_stage = last_vue_stage(ice); 2459 struct crocus_compiled_shader *shader = ice->shaders.prog[last_stage]; 2460 struct crocus_uncompiled_shader *ish = ice->shaders.uncompiled[last_stage]; 2461 update_last_vue_map(ice, shader->prog_data); 2462 if (ice->state.streamout != shader->streamout) { 2463 ice->state.streamout = shader->streamout; 2464 ice->state.dirty |= CROCUS_DIRTY_SO_DECL_LIST | CROCUS_DIRTY_STREAMOUT; 2465 } 2466 2467 if (ice->state.streamout_active) { 2468 screen->vtbl.update_so_strides(ice, ish->stream_output.stride); 2469 } 2470 2471 /* use ice->state version as last_vue_map can dirty this bit */ 2472 if (ice->state.stage_dirty & CROCUS_STAGE_DIRTY_UNCOMPILED_FS) 2473 crocus_update_compiled_fs(ice); 2474 2475 if (screen->devinfo.ver <= 6) { 2476 if (ice->state.dirty & CROCUS_DIRTY_GEN4_FF_GS_PROG && 2477 !ice->shaders.prog[MESA_SHADER_GEOMETRY]) 2478 crocus_update_compiled_ff_gs(ice); 2479 } 2480 2481 if (screen->devinfo.ver < 6) { 2482 if (ice->state.dirty & CROCUS_DIRTY_GEN4_CLIP_PROG) 2483 crocus_update_compiled_clip(ice); 2484 if (ice->state.dirty & CROCUS_DIRTY_GEN4_SF_PROG) 2485 crocus_update_compiled_sf(ice); 2486 } 2487 2488 2489 /* Changing shader interfaces may require a URB configuration. */ 2490 if (!(ice->state.dirty & CROCUS_DIRTY_GEN6_URB)) { 2491 for (int i = MESA_SHADER_VERTEX; i <= MESA_SHADER_GEOMETRY; i++) { 2492 struct brw_vue_prog_data *old = old_prog_datas[i]; 2493 struct brw_vue_prog_data *new = get_vue_prog_data(ice, i); 2494 if (!!old != !!new || 2495 (new && new->urb_entry_size != old->urb_entry_size)) { 2496 ice->state.dirty |= CROCUS_DIRTY_GEN6_URB; 2497 break; 2498 } 2499 } 2500 } 2501 2502 if (ice->state.stage_dirty & CROCUS_RENDER_STAGE_DIRTY_CONSTANTS) { 2503 for (int i = MESA_SHADER_VERTEX; i <= MESA_SHADER_FRAGMENT; i++) { 2504 if (ice->state.stage_dirty & (CROCUS_STAGE_DIRTY_CONSTANTS_VS << i)) 2505 crocus_update_pull_constant_descriptors(ice, i); 2506 } 2507 } 2508 return true; 2509} 2510 2511static struct crocus_compiled_shader * 2512crocus_compile_cs(struct crocus_context *ice, 2513 struct crocus_uncompiled_shader *ish, 2514 const struct brw_cs_prog_key *key) 2515{ 2516 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2517 const struct brw_compiler *compiler = screen->compiler; 2518 void *mem_ctx = ralloc_context(NULL); 2519 struct brw_cs_prog_data *cs_prog_data = 2520 rzalloc(mem_ctx, struct brw_cs_prog_data); 2521 struct brw_stage_prog_data *prog_data = &cs_prog_data->base; 2522 enum brw_param_builtin *system_values; 2523 const struct intel_device_info *devinfo = &screen->devinfo; 2524 unsigned num_system_values; 2525 unsigned num_cbufs; 2526 2527 nir_shader *nir = nir_shader_clone(mem_ctx, ish->nir); 2528 2529 NIR_PASS_V(nir, brw_nir_lower_cs_intrinsics); 2530 2531 crocus_setup_uniforms(compiler, mem_ctx, nir, prog_data, &system_values, 2532 &num_system_values, &num_cbufs); 2533 crocus_lower_swizzles(nir, &key->base.tex); 2534 struct crocus_binding_table bt; 2535 crocus_setup_binding_table(devinfo, nir, &bt, /* num_render_targets */ 0, 2536 num_system_values, num_cbufs, &key->base.tex); 2537 2538 struct brw_compile_cs_params params = { 2539 .nir = nir, 2540 .key = key, 2541 .prog_data = cs_prog_data, 2542 .log_data = &ice->dbg, 2543 }; 2544 2545 const unsigned *program = 2546 brw_compile_cs(compiler, mem_ctx, ¶ms); 2547 if (program == NULL) { 2548 dbg_printf("Failed to compile compute shader: %s\n", params.error_str); 2549 ralloc_free(mem_ctx); 2550 return false; 2551 } 2552 2553 if (ish->compiled_once) { 2554 crocus_debug_recompile(ice, &nir->info, &key->base); 2555 } else { 2556 ish->compiled_once = true; 2557 } 2558 2559 struct crocus_compiled_shader *shader = 2560 crocus_upload_shader(ice, CROCUS_CACHE_CS, sizeof(*key), key, program, 2561 prog_data->program_size, 2562 prog_data, sizeof(*cs_prog_data), NULL, 2563 system_values, num_system_values, 2564 num_cbufs, &bt); 2565 2566 crocus_disk_cache_store(screen->disk_cache, ish, shader, 2567 ice->shaders.cache_bo_map, 2568 key, sizeof(*key)); 2569 2570 ralloc_free(mem_ctx); 2571 return shader; 2572} 2573 2574static void 2575crocus_update_compiled_cs(struct crocus_context *ice) 2576{ 2577 struct crocus_shader_state *shs = &ice->state.shaders[MESA_SHADER_COMPUTE]; 2578 struct crocus_uncompiled_shader *ish = 2579 ice->shaders.uncompiled[MESA_SHADER_COMPUTE]; 2580 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2581 const struct intel_device_info *devinfo = &screen->devinfo; 2582 struct brw_cs_prog_key key = { KEY_INIT() }; 2583 2584 if (ish->nos & (1ull << CROCUS_NOS_TEXTURES)) 2585 crocus_populate_sampler_prog_key_data(ice, devinfo, MESA_SHADER_COMPUTE, ish, 2586 ish->nir->info.uses_texture_gather, &key.base.tex); 2587 screen->vtbl.populate_cs_key(ice, &key); 2588 2589 struct crocus_compiled_shader *old = ice->shaders.prog[CROCUS_CACHE_CS]; 2590 struct crocus_compiled_shader *shader = 2591 crocus_find_cached_shader(ice, CROCUS_CACHE_CS, sizeof(key), &key); 2592 2593 if (!shader) 2594 shader = crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key)); 2595 2596 if (!shader) 2597 shader = crocus_compile_cs(ice, ish, &key); 2598 2599 if (old != shader) { 2600 ice->shaders.prog[CROCUS_CACHE_CS] = shader; 2601 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_CS | 2602 CROCUS_STAGE_DIRTY_BINDINGS_CS | 2603 CROCUS_STAGE_DIRTY_CONSTANTS_CS; 2604 shs->sysvals_need_upload = true; 2605 } 2606} 2607 2608void 2609crocus_update_compiled_compute_shader(struct crocus_context *ice) 2610{ 2611 if (ice->state.stage_dirty & CROCUS_STAGE_DIRTY_UNCOMPILED_CS) 2612 crocus_update_compiled_cs(ice); 2613 2614 if (ice->state.stage_dirty & CROCUS_STAGE_DIRTY_CONSTANTS_CS) 2615 crocus_update_pull_constant_descriptors(ice, MESA_SHADER_COMPUTE); 2616} 2617 2618void 2619crocus_fill_cs_push_const_buffer(struct brw_cs_prog_data *cs_prog_data, 2620 unsigned threads, 2621 uint32_t *dst) 2622{ 2623 assert(brw_cs_push_const_total_size(cs_prog_data, threads) > 0); 2624 assert(cs_prog_data->push.cross_thread.size == 0); 2625 assert(cs_prog_data->push.per_thread.dwords == 1); 2626 assert(cs_prog_data->base.param[0] == BRW_PARAM_BUILTIN_SUBGROUP_ID); 2627 for (unsigned t = 0; t < threads; t++) 2628 dst[8 * t] = t; 2629} 2630 2631/** 2632 * Allocate scratch BOs as needed for the given per-thread size and stage. 2633 */ 2634struct crocus_bo * 2635crocus_get_scratch_space(struct crocus_context *ice, 2636 unsigned per_thread_scratch, 2637 gl_shader_stage stage) 2638{ 2639 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 2640 struct crocus_bufmgr *bufmgr = screen->bufmgr; 2641 const struct intel_device_info *devinfo = &screen->devinfo; 2642 2643 unsigned encoded_size = ffs(per_thread_scratch) - 11; 2644 assert(encoded_size < (1 << 16)); 2645 2646 struct crocus_bo **bop = &ice->shaders.scratch_bos[encoded_size][stage]; 2647 2648 if (!*bop) { 2649 assert(stage < ARRAY_SIZE(devinfo->max_scratch_ids)); 2650 uint32_t size = per_thread_scratch * devinfo->max_scratch_ids[stage]; 2651 *bop = crocus_bo_alloc(bufmgr, "scratch", size); 2652 } 2653 2654 return *bop; 2655} 2656 2657/* ------------------------------------------------------------------- */ 2658 2659/** 2660 * The pipe->create_[stage]_state() driver hooks. 2661 * 2662 * Performs basic NIR preprocessing, records any state dependencies, and 2663 * returns an crocus_uncompiled_shader as the Gallium CSO. 2664 * 2665 * Actual shader compilation to assembly happens later, at first use. 2666 */ 2667static void * 2668crocus_create_uncompiled_shader(struct pipe_context *ctx, 2669 nir_shader *nir, 2670 const struct pipe_stream_output_info *so_info) 2671{ 2672 struct crocus_screen *screen = (struct crocus_screen *)ctx->screen; 2673 const struct intel_device_info *devinfo = &screen->devinfo; 2674 struct crocus_uncompiled_shader *ish = 2675 calloc(1, sizeof(struct crocus_uncompiled_shader)); 2676 if (!ish) 2677 return NULL; 2678 2679 if (devinfo->ver >= 6) 2680 NIR_PASS(ish->needs_edge_flag, nir, crocus_fix_edge_flags); 2681 else 2682 ish->needs_edge_flag = false; 2683 2684 brw_preprocess_nir(screen->compiler, nir, NULL); 2685 2686 NIR_PASS_V(nir, brw_nir_lower_storage_image, devinfo); 2687 NIR_PASS_V(nir, crocus_lower_storage_image_derefs); 2688 2689 nir_sweep(nir); 2690 2691 ish->program_id = get_new_program_id(screen); 2692 ish->nir = nir; 2693 if (so_info) { 2694 memcpy(&ish->stream_output, so_info, sizeof(*so_info)); 2695 update_so_info(&ish->stream_output, nir->info.outputs_written); 2696 } 2697 2698 if (screen->disk_cache) { 2699 /* Serialize the NIR to a binary blob that we can hash for the disk 2700 * cache. Drop unnecessary information (like variable names) 2701 * so the serialized NIR is smaller, and also to let us detect more 2702 * isomorphic shaders when hashing, increasing cache hits. 2703 */ 2704 struct blob blob; 2705 blob_init(&blob); 2706 nir_serialize(&blob, nir, true); 2707 _mesa_sha1_compute(blob.data, blob.size, ish->nir_sha1); 2708 blob_finish(&blob); 2709 } 2710 2711 return ish; 2712} 2713 2714static struct crocus_uncompiled_shader * 2715crocus_create_shader_state(struct pipe_context *ctx, 2716 const struct pipe_shader_state *state) 2717{ 2718 struct nir_shader *nir; 2719 2720 if (state->type == PIPE_SHADER_IR_TGSI) 2721 nir = tgsi_to_nir(state->tokens, ctx->screen, false); 2722 else 2723 nir = state->ir.nir; 2724 2725 return crocus_create_uncompiled_shader(ctx, nir, &state->stream_output); 2726} 2727 2728static void * 2729crocus_create_vs_state(struct pipe_context *ctx, 2730 const struct pipe_shader_state *state) 2731{ 2732 struct crocus_context *ice = (void *) ctx; 2733 struct crocus_screen *screen = (void *) ctx->screen; 2734 struct crocus_uncompiled_shader *ish = crocus_create_shader_state(ctx, state); 2735 2736 ish->nos |= (1ull << CROCUS_NOS_TEXTURES); 2737 /* User clip planes or gen5 sprite coord enable */ 2738 if (ish->nir->info.clip_distance_array_size == 0 || 2739 screen->devinfo.ver <= 5) 2740 ish->nos |= (1ull << CROCUS_NOS_RASTERIZER); 2741 2742 if (screen->devinfo.verx10 < 75) 2743 ish->nos |= (1ull << CROCUS_NOS_VERTEX_ELEMENTS); 2744 2745 if (screen->precompile) { 2746 struct brw_vs_prog_key key = { KEY_INIT() }; 2747 2748 if (!crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key))) 2749 crocus_compile_vs(ice, ish, &key); 2750 } 2751 2752 return ish; 2753} 2754 2755static void * 2756crocus_create_tcs_state(struct pipe_context *ctx, 2757 const struct pipe_shader_state *state) 2758{ 2759 struct crocus_context *ice = (void *) ctx; 2760 struct crocus_screen *screen = (void *) ctx->screen; 2761 struct crocus_uncompiled_shader *ish = crocus_create_shader_state(ctx, state); 2762 struct shader_info *info = &ish->nir->info; 2763 2764 ish->nos |= (1ull << CROCUS_NOS_TEXTURES); 2765 if (screen->precompile) { 2766 const unsigned _GL_TRIANGLES = 0x0004; 2767 struct brw_tcs_prog_key key = { 2768 KEY_INIT(), 2769 // XXX: make sure the linker fills this out from the TES... 2770 .tes_primitive_mode = 2771 info->tess.primitive_mode ? info->tess.primitive_mode 2772 : _GL_TRIANGLES, 2773 .outputs_written = info->outputs_written, 2774 .patch_outputs_written = info->patch_outputs_written, 2775 }; 2776 2777 key.input_vertices = info->tess.tcs_vertices_out; 2778 2779 if (!crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key))) 2780 crocus_compile_tcs(ice, ish, &key); 2781 } 2782 2783 return ish; 2784} 2785 2786static void * 2787crocus_create_tes_state(struct pipe_context *ctx, 2788 const struct pipe_shader_state *state) 2789{ 2790 struct crocus_context *ice = (void *) ctx; 2791 struct crocus_screen *screen = (void *) ctx->screen; 2792 struct crocus_uncompiled_shader *ish = crocus_create_shader_state(ctx, state); 2793 struct shader_info *info = &ish->nir->info; 2794 2795 ish->nos |= (1ull << CROCUS_NOS_TEXTURES); 2796 /* User clip planes */ 2797 if (ish->nir->info.clip_distance_array_size == 0) 2798 ish->nos |= (1ull << CROCUS_NOS_RASTERIZER); 2799 2800 if (screen->precompile) { 2801 struct brw_tes_prog_key key = { 2802 KEY_INIT(), 2803 // XXX: not ideal, need TCS output/TES input unification 2804 .inputs_read = info->inputs_read, 2805 .patch_inputs_read = info->patch_inputs_read, 2806 }; 2807 2808 if (!crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key))) 2809 crocus_compile_tes(ice, ish, &key); 2810 } 2811 2812 return ish; 2813} 2814 2815static void * 2816crocus_create_gs_state(struct pipe_context *ctx, 2817 const struct pipe_shader_state *state) 2818{ 2819 struct crocus_context *ice = (void *) ctx; 2820 struct crocus_screen *screen = (void *) ctx->screen; 2821 struct crocus_uncompiled_shader *ish = crocus_create_shader_state(ctx, state); 2822 2823 ish->nos |= (1ull << CROCUS_NOS_TEXTURES); 2824 /* User clip planes */ 2825 if (ish->nir->info.clip_distance_array_size == 0) 2826 ish->nos |= (1ull << CROCUS_NOS_RASTERIZER); 2827 2828 if (screen->precompile) { 2829 struct brw_gs_prog_key key = { KEY_INIT() }; 2830 2831 if (!crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key))) 2832 crocus_compile_gs(ice, ish, &key); 2833 } 2834 2835 return ish; 2836} 2837 2838static void * 2839crocus_create_fs_state(struct pipe_context *ctx, 2840 const struct pipe_shader_state *state) 2841{ 2842 struct crocus_context *ice = (void *) ctx; 2843 struct crocus_screen *screen = (void *) ctx->screen; 2844 struct crocus_uncompiled_shader *ish = crocus_create_shader_state(ctx, state); 2845 struct shader_info *info = &ish->nir->info; 2846 2847 ish->nos |= (1ull << CROCUS_NOS_FRAMEBUFFER) | 2848 (1ull << CROCUS_NOS_DEPTH_STENCIL_ALPHA) | 2849 (1ull << CROCUS_NOS_RASTERIZER) | 2850 (1ull << CROCUS_NOS_TEXTURES) | 2851 (1ull << CROCUS_NOS_BLEND); 2852 2853 /* The program key needs the VUE map if there are > 16 inputs or gen4/5 */ 2854 if (screen->devinfo.ver < 6 || util_bitcount64(ish->nir->info.inputs_read & 2855 BRW_FS_VARYING_INPUT_MASK) > 16) { 2856 ish->nos |= (1ull << CROCUS_NOS_LAST_VUE_MAP); 2857 } 2858 2859 if (screen->precompile) { 2860 const uint64_t color_outputs = info->outputs_written & 2861 ~(BITFIELD64_BIT(FRAG_RESULT_DEPTH) | 2862 BITFIELD64_BIT(FRAG_RESULT_STENCIL) | 2863 BITFIELD64_BIT(FRAG_RESULT_SAMPLE_MASK)); 2864 2865 bool can_rearrange_varyings = 2866 screen->devinfo.ver > 6 && util_bitcount64(info->inputs_read & BRW_FS_VARYING_INPUT_MASK) <= 16; 2867 2868 const struct intel_device_info *devinfo = &screen->devinfo; 2869 struct brw_wm_prog_key key = { 2870 KEY_INIT(), 2871 .nr_color_regions = util_bitcount(color_outputs), 2872 .coherent_fb_fetch = false, 2873 .input_slots_valid = 2874 can_rearrange_varyings ? 0 : info->inputs_read | VARYING_BIT_POS, 2875 }; 2876 2877 struct brw_vue_map vue_map; 2878 if (devinfo->ver < 6) { 2879 brw_compute_vue_map(devinfo, &vue_map, 2880 info->inputs_read | VARYING_BIT_POS, 2881 false, /* pos slots */ 1); 2882 } 2883 if (!crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key))) 2884 crocus_compile_fs(ice, ish, &key, &vue_map); 2885 } 2886 2887 return ish; 2888} 2889 2890static void * 2891crocus_create_compute_state(struct pipe_context *ctx, 2892 const struct pipe_compute_state *state) 2893{ 2894 assert(state->ir_type == PIPE_SHADER_IR_NIR); 2895 2896 struct crocus_context *ice = (void *) ctx; 2897 struct crocus_screen *screen = (void *) ctx->screen; 2898 struct crocus_uncompiled_shader *ish = 2899 crocus_create_uncompiled_shader(ctx, (void *) state->prog, NULL); 2900 2901 ish->nos |= (1ull << CROCUS_NOS_TEXTURES); 2902 // XXX: disallow more than 64KB of shared variables 2903 2904 if (screen->precompile) { 2905 struct brw_cs_prog_key key = { KEY_INIT() }; 2906 2907 if (!crocus_disk_cache_retrieve(ice, ish, &key, sizeof(key))) 2908 crocus_compile_cs(ice, ish, &key); 2909 } 2910 2911 return ish; 2912} 2913 2914/** 2915 * The pipe->delete_[stage]_state() driver hooks. 2916 * 2917 * Frees the crocus_uncompiled_shader. 2918 */ 2919static void 2920crocus_delete_shader_state(struct pipe_context *ctx, void *state, gl_shader_stage stage) 2921{ 2922 struct crocus_uncompiled_shader *ish = state; 2923 struct crocus_context *ice = (void *) ctx; 2924 2925 if (ice->shaders.uncompiled[stage] == ish) { 2926 ice->shaders.uncompiled[stage] = NULL; 2927 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_UNCOMPILED_VS << stage; 2928 } 2929 2930 if (ish->const_data) { 2931 pipe_resource_reference(&ish->const_data, NULL); 2932 pipe_resource_reference(&ish->const_data_state.res, NULL); 2933 } 2934 2935 ralloc_free(ish->nir); 2936 free(ish); 2937} 2938 2939static void 2940crocus_delete_vs_state(struct pipe_context *ctx, void *state) 2941{ 2942 crocus_delete_shader_state(ctx, state, MESA_SHADER_VERTEX); 2943} 2944 2945static void 2946crocus_delete_tcs_state(struct pipe_context *ctx, void *state) 2947{ 2948 crocus_delete_shader_state(ctx, state, MESA_SHADER_TESS_CTRL); 2949} 2950 2951static void 2952crocus_delete_tes_state(struct pipe_context *ctx, void *state) 2953{ 2954 crocus_delete_shader_state(ctx, state, MESA_SHADER_TESS_EVAL); 2955} 2956 2957static void 2958crocus_delete_gs_state(struct pipe_context *ctx, void *state) 2959{ 2960 crocus_delete_shader_state(ctx, state, MESA_SHADER_GEOMETRY); 2961} 2962 2963static void 2964crocus_delete_fs_state(struct pipe_context *ctx, void *state) 2965{ 2966 crocus_delete_shader_state(ctx, state, MESA_SHADER_FRAGMENT); 2967} 2968 2969static void 2970crocus_delete_cs_state(struct pipe_context *ctx, void *state) 2971{ 2972 crocus_delete_shader_state(ctx, state, MESA_SHADER_COMPUTE); 2973} 2974 2975/** 2976 * The pipe->bind_[stage]_state() driver hook. 2977 * 2978 * Binds an uncompiled shader as the current one for a particular stage. 2979 * Updates dirty tracking to account for the shader's NOS. 2980 */ 2981static void 2982bind_shader_state(struct crocus_context *ice, 2983 struct crocus_uncompiled_shader *ish, 2984 gl_shader_stage stage) 2985{ 2986 uint64_t dirty_bit = CROCUS_STAGE_DIRTY_UNCOMPILED_VS << stage; 2987 const uint64_t nos = ish ? ish->nos : 0; 2988 2989 const struct shader_info *old_info = crocus_get_shader_info(ice, stage); 2990 const struct shader_info *new_info = ish ? &ish->nir->info : NULL; 2991 2992 if ((old_info ? BITSET_LAST_BIT(old_info->textures_used) : 0) != 2993 (new_info ? BITSET_LAST_BIT(new_info->textures_used) : 0)) { 2994 ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_SAMPLER_STATES_VS << stage; 2995 } 2996 2997 ice->shaders.uncompiled[stage] = ish; 2998 ice->state.stage_dirty |= dirty_bit; 2999 3000 /* Record that CSOs need to mark CROCUS_DIRTY_UNCOMPILED_XS when they change 3001 * (or that they no longer need to do so). 3002 */ 3003 for (int i = 0; i < CROCUS_NOS_COUNT; i++) { 3004 if (nos & (1 << i)) 3005 ice->state.stage_dirty_for_nos[i] |= dirty_bit; 3006 else 3007 ice->state.stage_dirty_for_nos[i] &= ~dirty_bit; 3008 } 3009} 3010 3011static void 3012crocus_bind_vs_state(struct pipe_context *ctx, void *state) 3013{ 3014 struct crocus_context *ice = (struct crocus_context *)ctx; 3015 struct crocus_uncompiled_shader *new_ish = state; 3016 struct crocus_screen *screen = (struct crocus_screen *)ice->ctx.screen; 3017 const struct intel_device_info *devinfo = &screen->devinfo; 3018 3019 if (new_ish && 3020 ice->state.window_space_position != 3021 new_ish->nir->info.vs.window_space_position) { 3022 ice->state.window_space_position = 3023 new_ish->nir->info.vs.window_space_position; 3024 3025 ice->state.dirty |= CROCUS_DIRTY_CLIP | 3026 CROCUS_DIRTY_RASTER | 3027 CROCUS_DIRTY_CC_VIEWPORT; 3028 } 3029 3030 if (devinfo->ver == 6) { 3031 ice->state.stage_dirty |= CROCUS_DIRTY_GEN4_FF_GS_PROG; 3032 } 3033 3034 bind_shader_state((void *) ctx, state, MESA_SHADER_VERTEX); 3035} 3036 3037static void 3038crocus_bind_tcs_state(struct pipe_context *ctx, void *state) 3039{ 3040 bind_shader_state((void *) ctx, state, MESA_SHADER_TESS_CTRL); 3041} 3042 3043static void 3044crocus_bind_tes_state(struct pipe_context *ctx, void *state) 3045{ 3046 struct crocus_context *ice = (struct crocus_context *)ctx; 3047 3048 /* Enabling/disabling optional stages requires a URB reconfiguration. */ 3049 if (!!state != !!ice->shaders.uncompiled[MESA_SHADER_TESS_EVAL]) 3050 ice->state.dirty |= CROCUS_DIRTY_GEN6_URB; 3051 3052 bind_shader_state((void *) ctx, state, MESA_SHADER_TESS_EVAL); 3053} 3054 3055static void 3056crocus_bind_gs_state(struct pipe_context *ctx, void *state) 3057{ 3058 struct crocus_context *ice = (struct crocus_context *)ctx; 3059 3060 /* Enabling/disabling optional stages requires a URB reconfiguration. */ 3061 if (!!state != !!ice->shaders.uncompiled[MESA_SHADER_GEOMETRY]) 3062 ice->state.dirty |= CROCUS_DIRTY_GEN6_URB; 3063 3064 bind_shader_state((void *) ctx, state, MESA_SHADER_GEOMETRY); 3065} 3066 3067static void 3068crocus_bind_fs_state(struct pipe_context *ctx, void *state) 3069{ 3070 struct crocus_context *ice = (struct crocus_context *) ctx; 3071 struct crocus_screen *screen = (struct crocus_screen *) ctx->screen; 3072 const struct intel_device_info *devinfo = &screen->devinfo; 3073 struct crocus_uncompiled_shader *old_ish = 3074 ice->shaders.uncompiled[MESA_SHADER_FRAGMENT]; 3075 struct crocus_uncompiled_shader *new_ish = state; 3076 3077 const unsigned color_bits = 3078 BITFIELD64_BIT(FRAG_RESULT_COLOR) | 3079 BITFIELD64_RANGE(FRAG_RESULT_DATA0, BRW_MAX_DRAW_BUFFERS); 3080 3081 /* Fragment shader outputs influence HasWriteableRT */ 3082 if (!old_ish || !new_ish || 3083 (old_ish->nir->info.outputs_written & color_bits) != 3084 (new_ish->nir->info.outputs_written & color_bits)) { 3085 if (devinfo->ver == 8) 3086 ice->state.dirty |= CROCUS_DIRTY_GEN8_PS_BLEND; 3087 else 3088 ice->state.dirty |= CROCUS_DIRTY_WM; 3089 } 3090 3091 if (devinfo->ver == 8) 3092 ice->state.dirty |= CROCUS_DIRTY_GEN8_PMA_FIX; 3093 bind_shader_state((void *) ctx, state, MESA_SHADER_FRAGMENT); 3094} 3095 3096static void 3097crocus_bind_cs_state(struct pipe_context *ctx, void *state) 3098{ 3099 bind_shader_state((void *) ctx, state, MESA_SHADER_COMPUTE); 3100} 3101 3102void 3103crocus_init_program_functions(struct pipe_context *ctx) 3104{ 3105 ctx->create_vs_state = crocus_create_vs_state; 3106 ctx->create_tcs_state = crocus_create_tcs_state; 3107 ctx->create_tes_state = crocus_create_tes_state; 3108 ctx->create_gs_state = crocus_create_gs_state; 3109 ctx->create_fs_state = crocus_create_fs_state; 3110 ctx->create_compute_state = crocus_create_compute_state; 3111 3112 ctx->delete_vs_state = crocus_delete_vs_state; 3113 ctx->delete_tcs_state = crocus_delete_tcs_state; 3114 ctx->delete_tes_state = crocus_delete_tes_state; 3115 ctx->delete_gs_state = crocus_delete_gs_state; 3116 ctx->delete_fs_state = crocus_delete_fs_state; 3117 ctx->delete_compute_state = crocus_delete_cs_state; 3118 3119 ctx->bind_vs_state = crocus_bind_vs_state; 3120 ctx->bind_tcs_state = crocus_bind_tcs_state; 3121 ctx->bind_tes_state = crocus_bind_tes_state; 3122 ctx->bind_gs_state = crocus_bind_gs_state; 3123 ctx->bind_fs_state = crocus_bind_fs_state; 3124 ctx->bind_compute_state = crocus_bind_cs_state; 3125} 3126