/* * Copyright © 2022 Collabora Ltd. and Red Hat Inc. * SPDX-License-Identifier: MIT */ #include "nvk_cmd_buffer.h" #include "nvk_descriptor_set_layout.h" #include "nvk_descriptor_types.h" #include "nvk_shader.h" #include "vk_pipeline.h" #include "nir_builder.h" #include "nir_deref.h" #include "clc397.h" #include "clc597.h" struct lower_desc_cbuf { struct nvk_cbuf key; uint32_t use_count; uint64_t start; uint64_t end; }; DERIVE_HASH_TABLE(nvk_cbuf); static int compar_cbufs(const void *_a, const void *_b) { const struct lower_desc_cbuf *a = _a; const struct lower_desc_cbuf *b = _b; #define COMPAR(field, pos) \ if (a->field < b->field) return -(pos); \ if (a->field > b->field) return (pos); /* Sort by most used first */ COMPAR(use_count, -1) /* Keep the list stable by then sorting by key fields. */ COMPAR(key.type, 1) COMPAR(key.desc_set, 1) COMPAR(key.dynamic_idx, 1) COMPAR(key.desc_offset, 1) #undef COMPAR return 0; } struct lower_descriptors_ctx { const struct nv_device_info *dev_info; const struct nvk_descriptor_set_layout *set_layouts[NVK_MAX_SETS]; bool use_bindless_cbuf; bool use_edb_buffer_views; bool clamp_desc_array_bounds; nir_address_format ubo_addr_format; nir_address_format ssbo_addr_format; struct hash_table *cbufs; struct nvk_cbuf_map *cbuf_map; }; static bool descriptor_type_is_ubo(VkDescriptorType desc_type) { switch (desc_type) { case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: case VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK: return true; default: return false; } } static bool descriptor_type_is_ssbo(VkDescriptorType desc_type) { switch (desc_type) { case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: return true; default: return false; } } static void record_cbuf_use(const struct nvk_cbuf *key, uint64_t start, uint64_t end, struct lower_descriptors_ctx *ctx) { struct hash_entry *entry = _mesa_hash_table_search(ctx->cbufs, key); if (entry != NULL) { struct lower_desc_cbuf *cbuf = entry->data; cbuf->use_count++; cbuf->start = MIN2(cbuf->start, start); cbuf->end = MAX2(cbuf->end, end); } else { struct lower_desc_cbuf *cbuf = ralloc(ctx->cbufs, struct lower_desc_cbuf); *cbuf = (struct lower_desc_cbuf) { .key = *key, .use_count = 1, .start = start, .end = end, }; _mesa_hash_table_insert(ctx->cbufs, &cbuf->key, cbuf); } } static const struct nvk_descriptor_set_binding_layout * get_binding_layout(uint32_t set, uint32_t binding, const struct lower_descriptors_ctx *ctx) { assert(set < NVK_MAX_SETS); assert(ctx->set_layouts[set] != NULL); const struct nvk_descriptor_set_layout *set_layout = ctx->set_layouts[set]; assert(binding < set_layout->binding_count); return &set_layout->binding[binding]; } static void record_descriptor_cbuf_use(uint32_t set, uint32_t binding, nir_src *index, struct lower_descriptors_ctx *ctx) { const struct nvk_descriptor_set_binding_layout *binding_layout = get_binding_layout(set, binding, ctx); const struct nvk_cbuf key = { .type = NVK_CBUF_TYPE_DESC_SET, .desc_set = set, }; uint64_t start, end; if (index == NULL) { /* When we don't have an index, assume 0 */ start = binding_layout->offset; end = start + binding_layout->stride; } else if (nir_src_is_const(*index)) { start = binding_layout->offset + nir_src_as_uint(*index) * binding_layout->stride; end = start + binding_layout->stride; } else { start = binding_layout->offset; end = start + binding_layout->array_size * binding_layout->stride; } record_cbuf_use(&key, start, end, ctx); } static void record_vulkan_resource_cbuf_use(nir_intrinsic_instr *intrin, struct lower_descriptors_ctx *ctx) { assert(intrin->intrinsic == nir_intrinsic_vulkan_resource_index); /* These we'll handle later */ if (descriptor_type_is_ubo(nir_intrinsic_desc_type(intrin))) return; record_descriptor_cbuf_use(nir_intrinsic_desc_set(intrin), nir_intrinsic_binding(intrin), &intrin->src[0], ctx); } static void record_deref_descriptor_cbuf_use(nir_deref_instr *deref, struct lower_descriptors_ctx *ctx) { nir_src *index_src = NULL; if (deref->deref_type == nir_deref_type_array) { index_src = &deref->arr.index; deref = nir_deref_instr_parent(deref); } assert(deref->deref_type == nir_deref_type_var); nir_variable *var = deref->var; record_descriptor_cbuf_use(var->data.descriptor_set, var->data.binding, index_src, ctx); } static void record_tex_descriptor_cbuf_use(nir_tex_instr *tex, struct lower_descriptors_ctx *ctx) { const int texture_src_idx = nir_tex_instr_src_index(tex, nir_tex_src_texture_deref); const int sampler_src_idx = nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref); if (texture_src_idx >= 0) { nir_deref_instr *deref = nir_src_as_deref(tex->src[texture_src_idx].src); record_deref_descriptor_cbuf_use(deref, ctx); } if (sampler_src_idx >= 0) { nir_deref_instr *deref = nir_src_as_deref(tex->src[sampler_src_idx].src); record_deref_descriptor_cbuf_use(deref, ctx); } } static struct nvk_cbuf ubo_deref_to_cbuf(nir_deref_instr *deref, nir_intrinsic_instr **resource_index_out, uint64_t *offset_out, uint64_t *start_out, uint64_t *end_out, const struct lower_descriptors_ctx *ctx) { assert(nir_deref_mode_is(deref, nir_var_mem_ubo)); /* In case we early return */ *offset_out = 0; *start_out = 0; *end_out = UINT64_MAX; *resource_index_out = NULL; const struct nvk_cbuf invalid = { .type = NVK_CBUF_TYPE_INVALID, }; uint64_t offset = 0; uint64_t range = glsl_get_explicit_size(deref->type, false); bool offset_valid = true; while (deref->deref_type != nir_deref_type_cast) { nir_deref_instr *parent = nir_deref_instr_parent(deref); switch (deref->deref_type) { case nir_deref_type_var: unreachable("Buffers don't use variables in Vulkan"); case nir_deref_type_array: case nir_deref_type_array_wildcard: { uint32_t stride = nir_deref_instr_array_stride(deref); if (deref->deref_type == nir_deref_type_array && nir_src_is_const(deref->arr.index)) { offset += nir_src_as_uint(deref->arr.index) * stride; } else { range = glsl_get_length(parent->type) * stride; } break; } case nir_deref_type_ptr_as_array: /* All bets are off. We shouldn't see these most of the time * anyway, even with variable pointers. */ offset_valid = false; unreachable("Variable pointers aren't allowed on UBOs"); break; case nir_deref_type_struct: { offset += glsl_get_struct_field_offset(parent->type, deref->strct.index); break; } default: unreachable("Unknown deref type"); } deref = parent; } nir_intrinsic_instr *load_desc = nir_src_as_intrinsic(deref->parent); if (load_desc == NULL || load_desc->intrinsic != nir_intrinsic_load_vulkan_descriptor) return invalid; nir_intrinsic_instr *res_index = nir_src_as_intrinsic(load_desc->src[0]); if (res_index == NULL || res_index->intrinsic != nir_intrinsic_vulkan_resource_index) return invalid; /* We try to early return as little as possible prior to this point so we * can return the resource index intrinsic in as many cases as possible. * After this point, though, early returns are fair game. */ *resource_index_out = res_index; if (!offset_valid || !nir_src_is_const(res_index->src[0])) return invalid; uint32_t set = nir_intrinsic_desc_set(res_index); uint32_t binding = nir_intrinsic_binding(res_index); uint32_t index = nir_src_as_uint(res_index->src[0]); const struct nvk_descriptor_set_binding_layout *binding_layout = get_binding_layout(set, binding, ctx); switch (binding_layout->type) { case VK_DESCRIPTOR_TYPE_MUTABLE_EXT: case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER: { *offset_out = 0; *start_out = offset; *end_out = offset + range; return (struct nvk_cbuf) { .type = NVK_CBUF_TYPE_UBO_DESC, .desc_set = set, .desc_offset = binding_layout->offset + index * binding_layout->stride, }; } case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: { *offset_out = 0; *start_out = offset; *end_out = offset + range; return (struct nvk_cbuf) { .type = NVK_CBUF_TYPE_DYNAMIC_UBO, .desc_set = set, .dynamic_idx = binding_layout->dynamic_buffer_index + index, }; } case VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK: { *offset_out = binding_layout->offset; *start_out = binding_layout->offset + offset; *end_out = *start_out + range; return (struct nvk_cbuf) { .type = NVK_CBUF_TYPE_DESC_SET, .desc_set = set, }; } default: return invalid; } } static void record_load_ubo_cbuf_uses(nir_deref_instr *deref, struct lower_descriptors_ctx *ctx) { assert(nir_deref_mode_is(deref, nir_var_mem_ubo)); UNUSED uint64_t offset; uint64_t start, end; nir_intrinsic_instr *res_index; struct nvk_cbuf cbuf = ubo_deref_to_cbuf(deref, &res_index, &offset, &start, &end, ctx); if (cbuf.type != NVK_CBUF_TYPE_INVALID) { record_cbuf_use(&cbuf, start, end, ctx); } else if (res_index != NULL) { record_vulkan_resource_cbuf_use(res_index, ctx); } } static bool record_cbuf_uses_instr(UNUSED nir_builder *b, nir_instr *instr, void *_ctx) { struct lower_descriptors_ctx *ctx = _ctx; switch (instr->type) { case nir_instr_type_intrinsic: { nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr); switch (intrin->intrinsic) { case nir_intrinsic_vulkan_resource_index: record_vulkan_resource_cbuf_use(intrin, ctx); return false; case nir_intrinsic_load_deref: { nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); if (nir_deref_mode_is(deref, nir_var_mem_ubo)) record_load_ubo_cbuf_uses(deref, ctx); return false; } case nir_intrinsic_image_deref_load: case nir_intrinsic_image_deref_store: case nir_intrinsic_image_deref_atomic: case nir_intrinsic_image_deref_atomic_swap: case nir_intrinsic_image_deref_size: case nir_intrinsic_image_deref_samples: { nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); record_deref_descriptor_cbuf_use(deref, ctx); return false; } default: return false; } unreachable("All cases return false"); } case nir_instr_type_tex: record_tex_descriptor_cbuf_use(nir_instr_as_tex(instr), ctx); return false; default: return false; } } static void build_cbuf_map(nir_shader *nir, struct lower_descriptors_ctx *ctx) { ctx->cbuf_map->cbuf_count = 0; /* Root descriptors always go in cbuf 0 */ ctx->cbuf_map->cbufs[ctx->cbuf_map->cbuf_count++] = (struct nvk_cbuf) { .type = NVK_CBUF_TYPE_ROOT_DESC, }; /* If we have constant data, put it at cbuf 1 */ if (nir->constant_data_size > 0) { ctx->cbuf_map->cbufs[ctx->cbuf_map->cbuf_count++] = (struct nvk_cbuf) { .type = NVK_CBUF_TYPE_SHADER_DATA, }; } ctx->cbufs = nvk_cbuf_table_create(NULL); nir_shader_instructions_pass(nir, record_cbuf_uses_instr, nir_metadata_all, (void *)ctx); struct lower_desc_cbuf *cbufs = ralloc_array(ctx->cbufs, struct lower_desc_cbuf, _mesa_hash_table_num_entries(ctx->cbufs)); uint32_t num_cbufs = 0; hash_table_foreach(ctx->cbufs, entry) { struct lower_desc_cbuf *cbuf = entry->data; /* We currently only start cbufs at the beginning so if it starts after * the max cbuf size, there's no point in including it in the list. */ if (cbuf->start > NVK_MAX_CBUF_SIZE) continue; cbufs[num_cbufs++] = *cbuf; } qsort(cbufs, num_cbufs, sizeof(*cbufs), compar_cbufs); uint8_t max_cbuf_bindings; if (nir->info.stage == MESA_SHADER_COMPUTE || nir->info.stage == MESA_SHADER_KERNEL) { max_cbuf_bindings = 8; } else { max_cbuf_bindings = 16; } for (uint32_t i = 0; i < num_cbufs; i++) { if (ctx->cbuf_map->cbuf_count >= max_cbuf_bindings) break; /* We can't support indirect cbufs in compute yet */ if ((nir->info.stage == MESA_SHADER_COMPUTE || nir->info.stage == MESA_SHADER_KERNEL) && cbufs[i].key.type == NVK_CBUF_TYPE_UBO_DESC) continue; /* Prior to Turing, indirect cbufs require splitting the pushbuf and * pushing bits of the descriptor set. Doing this every draw call is * probably more overhead than it's worth. */ if (ctx->dev_info->cls_eng3d < TURING_A && cbufs[i].key.type == NVK_CBUF_TYPE_UBO_DESC) continue; ctx->cbuf_map->cbufs[ctx->cbuf_map->cbuf_count++] = cbufs[i].key; } ralloc_free(ctx->cbufs); ctx->cbufs = NULL; } static int get_mapped_cbuf_idx(const struct nvk_cbuf *key, const struct lower_descriptors_ctx *ctx) { if (ctx->cbuf_map == NULL) return -1; for (uint32_t c = 0; c < ctx->cbuf_map->cbuf_count; c++) { if (nvk_cbuf_equal(&ctx->cbuf_map->cbufs[c], key)) { return c; } } return -1; } static bool lower_load_ubo_intrin(nir_builder *b, nir_intrinsic_instr *load, void *_ctx) { const struct lower_descriptors_ctx *ctx = _ctx; if (load->intrinsic != nir_intrinsic_load_deref) return false; nir_deref_instr *deref = nir_src_as_deref(load->src[0]); if (!nir_deref_mode_is(deref, nir_var_mem_ubo)) return false; uint64_t offset, end; UNUSED uint64_t start; UNUSED nir_intrinsic_instr *res_index; struct nvk_cbuf cbuf_key = ubo_deref_to_cbuf(deref, &res_index, &offset, &start, &end, ctx); if (cbuf_key.type == NVK_CBUF_TYPE_INVALID) return false; if (end > NVK_MAX_CBUF_SIZE) return false; int cbuf_idx = get_mapped_cbuf_idx(&cbuf_key, ctx); if (cbuf_idx < 0) return false; b->cursor = nir_before_instr(&load->instr); nir_deref_path path; nir_deref_path_init(&path, deref, NULL); nir_def *addr = nir_imm_ivec2(b, cbuf_idx, offset); nir_address_format addr_format = nir_address_format_32bit_index_offset; for (nir_deref_instr **p = &path.path[1]; *p != NULL; p++) addr = nir_explicit_io_address_from_deref(b, *p, addr, addr_format); nir_deref_path_finish(&path); nir_lower_explicit_io_instr(b, load, addr, addr_format); return true; } static bool lower_load_constant(nir_builder *b, nir_intrinsic_instr *load, const struct lower_descriptors_ctx *ctx) { assert(load->intrinsic == nir_intrinsic_load_constant); const struct nvk_cbuf cbuf_key = { .type = NVK_CBUF_TYPE_SHADER_DATA, }; int cbuf_idx = get_mapped_cbuf_idx(&cbuf_key, ctx); assert(cbuf_idx >= 0); uint32_t base = nir_intrinsic_base(load); b->cursor = nir_before_instr(&load->instr); nir_def *offset = nir_iadd_imm(b, load->src[0].ssa, base); nir_def *data = nir_ldc_nv(b, load->def.num_components, load->def.bit_size, nir_imm_int(b, cbuf_idx), offset, .align_mul = nir_intrinsic_align_mul(load), .align_offset = nir_intrinsic_align_offset(load)); nir_def_rewrite_uses(&load->def, data); return true; } static nir_def * load_descriptor_set_addr(nir_builder *b, uint32_t set, UNUSED const struct lower_descriptors_ctx *ctx) { uint32_t set_addr_offset = nvk_root_descriptor_offset(sets) + set * sizeof(struct nvk_buffer_address); return nir_ldc_nv(b, 1, 64, nir_imm_int(b, 0), nir_imm_int(b, set_addr_offset), .align_mul = 8, .align_offset = 0); } static nir_def * load_dynamic_buffer_start(nir_builder *b, uint32_t set, const struct lower_descriptors_ctx *ctx) { int dynamic_buffer_start_imm = 0; for (uint32_t s = 0; s < set; s++) { if (ctx->set_layouts[s] == NULL) { dynamic_buffer_start_imm = -1; break; } dynamic_buffer_start_imm += ctx->set_layouts[s]->dynamic_buffer_count; } if (dynamic_buffer_start_imm >= 0) { return nir_imm_int(b, dynamic_buffer_start_imm); } else { uint32_t root_offset = nvk_root_descriptor_offset(set_dynamic_buffer_start) + set; return nir_u2u32(b, nir_ldc_nv(b, 1, 8, nir_imm_int(b, 0), nir_imm_int(b, root_offset), .align_mul = 1, .align_offset = 0)); } } static nir_def * load_descriptor(nir_builder *b, unsigned num_components, unsigned bit_size, uint32_t set, uint32_t binding, nir_def *index, unsigned offset_B, const struct lower_descriptors_ctx *ctx) { const struct nvk_descriptor_set_binding_layout *binding_layout = get_binding_layout(set, binding, ctx); if (ctx->clamp_desc_array_bounds) index = nir_umin(b, index, nir_imm_int(b, binding_layout->array_size - 1)); switch (binding_layout->type) { case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC: case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: { /* Get the index in the root descriptor table dynamic_buffers array. */ nir_def *dynamic_buffer_start = load_dynamic_buffer_start(b, set, ctx); index = nir_iadd(b, index, nir_iadd_imm(b, dynamic_buffer_start, binding_layout->dynamic_buffer_index)); uint32_t desc_size = sizeof(union nvk_buffer_descriptor); nir_def *root_desc_offset = nir_iadd_imm(b, nir_imul_imm(b, index, desc_size), nvk_root_descriptor_offset(dynamic_buffers)); assert(num_components * bit_size <= desc_size * 8); return nir_ldc_nv(b, num_components, bit_size, nir_imm_int(b, 0), root_desc_offset, .align_mul = 16, .align_offset = 0); } case VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK: { nir_def *base_addr = nir_iadd_imm(b, load_descriptor_set_addr(b, set, ctx), binding_layout->offset); assert(binding_layout->stride == 1); const uint32_t binding_size = binding_layout->array_size; if (ctx->use_bindless_cbuf) { assert(num_components == 1 && bit_size == 64); const uint32_t size = align(binding_size, 16); return nir_ior_imm(b, nir_ishr_imm(b, base_addr, 4), ((uint64_t)size >> 4) << 45); } else { /* Convert it to nir_address_format_64bit_bounded_global */ assert(num_components == 4 && bit_size == 32); return nir_vec4(b, nir_unpack_64_2x32_split_x(b, base_addr), nir_unpack_64_2x32_split_y(b, base_addr), nir_imm_int(b, binding_size), nir_imm_int(b, 0)); } } default: { assert(binding_layout->stride > 0); nir_def *desc_ubo_offset = nir_iadd_imm(b, nir_imul_imm(b, index, binding_layout->stride), binding_layout->offset + offset_B); uint64_t max_desc_ubo_offset = binding_layout->offset + binding_layout->array_size * binding_layout->stride; unsigned desc_align_mul = (1 << (ffs(binding_layout->stride) - 1)); desc_align_mul = MIN2(desc_align_mul, 16); unsigned desc_align_offset = binding_layout->offset + offset_B; desc_align_offset %= desc_align_mul; const struct nvk_cbuf cbuf_key = { .type = NVK_CBUF_TYPE_DESC_SET, .desc_set = set, }; int cbuf_idx = get_mapped_cbuf_idx(&cbuf_key, ctx); if (cbuf_idx >= 0 && max_desc_ubo_offset <= NVK_MAX_CBUF_SIZE) { return nir_ldc_nv(b, num_components, bit_size, nir_imm_int(b, cbuf_idx), desc_ubo_offset, .align_mul = desc_align_mul, .align_offset = desc_align_offset); } else { nir_def *set_addr = load_descriptor_set_addr(b, set, ctx); return nir_load_global_constant_offset(b, num_components, bit_size, set_addr, desc_ubo_offset, .align_mul = desc_align_mul, .align_offset = desc_align_offset); } } } } static bool is_idx_intrin(nir_intrinsic_instr *intrin) { while (intrin->intrinsic == nir_intrinsic_vulkan_resource_reindex) { intrin = nir_src_as_intrinsic(intrin->src[0]); if (intrin == NULL) return false; } return intrin->intrinsic == nir_intrinsic_vulkan_resource_index; } static nir_def * buffer_address_to_ldcx_handle(nir_builder *b, nir_def *addr) { nir_def *base_addr = nir_pack_64_2x32(b, nir_channels(b, addr, 0x3)); nir_def *size = nir_channel(b, addr, 2); nir_def *offset = nir_channel(b, addr, 3); nir_def *addr16 = nir_ushr_imm(b, base_addr, 4); nir_def *addr16_lo = nir_unpack_64_2x32_split_x(b, addr16); nir_def *addr16_hi = nir_unpack_64_2x32_split_y(b, addr16); /* If we assume the top bis of the address are 0 as well as the bottom two * bits of the size. (We can trust it since it's a descriptor) then * * ((size >> 4) << 13) | addr * * is just an imad. */ nir_def *handle_hi = nir_imad(b, size, nir_imm_int(b, 1 << 9), addr16_hi); return nir_vec3(b, addr16_lo, handle_hi, offset); } static nir_def * load_descriptor_for_idx_intrin(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { nir_def *index = nir_imm_int(b, 0); while (intrin->intrinsic == nir_intrinsic_vulkan_resource_reindex) { index = nir_iadd(b, index, intrin->src[1].ssa); intrin = nir_src_as_intrinsic(intrin->src[0]); } assert(intrin->intrinsic == nir_intrinsic_vulkan_resource_index); uint32_t set = nir_intrinsic_desc_set(intrin); uint32_t binding = nir_intrinsic_binding(intrin); index = nir_iadd(b, index, intrin->src[0].ssa); const VkDescriptorType desc_type = nir_intrinsic_desc_type(intrin); if (descriptor_type_is_ubo(desc_type) && ctx->use_bindless_cbuf) { nir_def *desc = load_descriptor(b, 1, 64, set, binding, index, 0, ctx); /* The descriptor is just the handle. NIR also needs an offset. */ return nir_vec3(b, nir_unpack_64_2x32_split_x(b, desc), nir_unpack_64_2x32_split_y(b, desc), nir_imm_int(b, 0)); } else { nir_def *desc = load_descriptor(b, 4, 32, set, binding, index, 0, ctx); /* We know a priori that the the .w compnent (offset) is zero */ return nir_vec4(b, nir_channel(b, desc, 0), nir_channel(b, desc, 1), nir_channel(b, desc, 2), nir_imm_int(b, 0)); } } static bool try_lower_load_vulkan_descriptor(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { ASSERTED const VkDescriptorType desc_type = nir_intrinsic_desc_type(intrin); b->cursor = nir_before_instr(&intrin->instr); nir_intrinsic_instr *idx_intrin = nir_src_as_intrinsic(intrin->src[0]); if (idx_intrin == NULL || !is_idx_intrin(idx_intrin)) { assert(descriptor_type_is_ssbo(desc_type)); return false; } nir_def *desc = load_descriptor_for_idx_intrin(b, idx_intrin, ctx); nir_def_rewrite_uses(&intrin->def, desc); return true; } static bool _lower_sysval_to_root_table(nir_builder *b, nir_intrinsic_instr *intrin, uint32_t root_table_offset, const struct lower_descriptors_ctx *ctx) { b->cursor = nir_instr_remove(&intrin->instr); nir_def *val = nir_ldc_nv(b, intrin->def.num_components, intrin->def.bit_size, nir_imm_int(b, 0), /* Root table */ nir_imm_int(b, root_table_offset), .align_mul = 4, .align_offset = 0); nir_def_rewrite_uses(&intrin->def, val); return true; } #define lower_sysval_to_root_table(b, intrin, member, ctx) \ _lower_sysval_to_root_table(b, intrin, \ nvk_root_descriptor_offset(member), \ ctx) static bool lower_load_push_constant(nir_builder *b, nir_intrinsic_instr *load, const struct lower_descriptors_ctx *ctx) { const uint32_t push_region_offset = nvk_root_descriptor_offset(push); const uint32_t base = nir_intrinsic_base(load); b->cursor = nir_before_instr(&load->instr); nir_def *offset = nir_iadd_imm(b, load->src[0].ssa, push_region_offset + base); nir_def *val = nir_ldc_nv(b, load->def.num_components, load->def.bit_size, nir_imm_int(b, 0), offset, .align_mul = load->def.bit_size / 8, .align_offset = 0); nir_def_rewrite_uses(&load->def, val); return true; } static void get_resource_deref_binding(nir_builder *b, nir_deref_instr *deref, uint32_t *set, uint32_t *binding, nir_def **index) { if (deref->deref_type == nir_deref_type_array) { *index = deref->arr.index.ssa; deref = nir_deref_instr_parent(deref); } else { *index = nir_imm_int(b, 0); } assert(deref->deref_type == nir_deref_type_var); nir_variable *var = deref->var; *set = var->data.descriptor_set; *binding = var->data.binding; } static nir_def * load_resource_deref_desc(nir_builder *b, unsigned num_components, unsigned bit_size, nir_deref_instr *deref, unsigned offset_B, const struct lower_descriptors_ctx *ctx) { uint32_t set, binding; nir_def *index; get_resource_deref_binding(b, deref, &set, &binding, &index); return load_descriptor(b, num_components, bit_size, set, binding, index, offset_B, ctx); } static void lower_msaa_image_intrin(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { assert(nir_intrinsic_image_dim(intrin) == GLSL_SAMPLER_DIM_MS); b->cursor = nir_before_instr(&intrin->instr); nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); nir_def *desc = load_resource_deref_desc(b, 1, 32, deref, 0, ctx); nir_def *img_index = nir_ubitfield_extract_imm(b, desc, 0, 20); nir_rewrite_image_intrinsic(intrin, img_index, true); nir_def *sw_log2 = nir_ubitfield_extract_imm(b, desc, 20, 2); nir_def *sh_log2 = nir_ubitfield_extract_imm(b, desc, 22, 2); nir_def *sw = nir_ishl(b, nir_imm_int(b, 1), sw_log2); nir_def *sh = nir_ishl(b, nir_imm_int(b, 1), sh_log2); nir_def *num_samples = nir_imul(b, sw, sh); switch (intrin->intrinsic) { case nir_intrinsic_bindless_image_load: case nir_intrinsic_bindless_image_store: case nir_intrinsic_bindless_image_atomic: case nir_intrinsic_bindless_image_atomic_swap: { nir_def *x = nir_channel(b, intrin->src[1].ssa, 0); nir_def *y = nir_channel(b, intrin->src[1].ssa, 1); nir_def *z = nir_channel(b, intrin->src[1].ssa, 2); nir_def *w = nir_channel(b, intrin->src[1].ssa, 3); nir_def *s = intrin->src[2].ssa; nir_def *sw_mask = nir_iadd_imm(b, sw, -1); nir_def *sx = nir_iand(b, s, sw_mask); nir_def *sy = nir_ishr(b, s, sw_log2); x = nir_imad(b, x, sw, sx); y = nir_imad(b, y, sh, sy); /* Make OOB sample indices OOB X/Y indices */ x = nir_bcsel(b, nir_ult(b, s, num_samples), x, nir_imm_int(b, -1)); nir_src_rewrite(&intrin->src[1], nir_vec4(b, x, y, z, w)); nir_src_rewrite(&intrin->src[2], nir_undef(b, 1, 32)); break; } case nir_intrinsic_bindless_image_size: { b->cursor = nir_after_instr(&intrin->instr); nir_def *size = &intrin->def; nir_def *w = nir_channel(b, size, 0); nir_def *h = nir_channel(b, size, 1); w = nir_ushr(b, w, sw_log2); h = nir_ushr(b, h, sh_log2); size = nir_vector_insert_imm(b, size, w, 0); size = nir_vector_insert_imm(b, size, h, 1); nir_def_rewrite_uses_after(&intrin->def, size, size->parent_instr); break; } case nir_intrinsic_bindless_image_samples: { /* We need to handle NULL descriptors explicitly */ nir_def *samples = nir_bcsel(b, nir_ieq(b, desc, nir_imm_int(b, 0)), nir_imm_int(b, 0), num_samples); nir_def_rewrite_uses(&intrin->def, samples); break; } default: unreachable("Unknown image intrinsic"); } nir_intrinsic_set_image_dim(intrin, GLSL_SAMPLER_DIM_2D); } static bool is_edb_buffer_view(nir_deref_instr *deref, const struct lower_descriptors_ctx *ctx) { if (glsl_get_sampler_dim(deref->type) != GLSL_SAMPLER_DIM_BUF) return false; if (ctx->use_edb_buffer_views) return true; nir_variable *var = nir_deref_instr_get_variable(deref); uint8_t set = var->data.descriptor_set; return ctx->set_layouts[set]->flags & VK_DESCRIPTOR_SET_LAYOUT_CREATE_DESCRIPTOR_BUFFER_BIT_EXT; } static nir_def * edb_buffer_view_is_null(nir_builder *b, nir_def *desc) { assert(desc->num_components == 4); nir_def *index = nir_channel(b, desc, 0); return nir_ieq_imm(b, index, 0); } static nir_def * edb_buffer_view_offset_el(nir_builder *b, nir_def *desc) { assert(desc->num_components == 4); return nir_channel(b, desc, 1); } static nir_def * edb_buffer_view_size_el(nir_builder *b, nir_def *desc) { assert(desc->num_components == 4); return nir_channel(b, desc, 2); } static nir_def * edb_buffer_view_oob_alpha(nir_builder *b, nir_def *desc) { assert(desc->num_components == 4); return nir_channel(b, desc, 3); } static nir_def * edb_buffer_view_coord_is_in_bounds(nir_builder *b, nir_def *desc, nir_def *coord) { assert(desc->num_components == 4); return nir_ult(b, coord, edb_buffer_view_size_el(b, desc)); } static nir_def * edb_buffer_view_index(nir_builder *b, nir_def *desc, nir_def *in_bounds) { assert(desc->num_components == 4); nir_def *index = nir_channel(b, desc, 0); /* Use the NULL descriptor for OOB access */ return nir_bcsel(b, in_bounds, index, nir_imm_int(b, 0)); } static nir_def * adjust_edb_buffer_view_coord(nir_builder *b, nir_def *desc, nir_def *coord) { return nir_iadd(b, coord, edb_buffer_view_offset_el(b, desc)); } static nir_def * fixup_edb_buffer_view_result(nir_builder *b, nir_def *desc, nir_def *in_bounds, nir_def *res, nir_alu_type dest_type) { if (res->num_components < 4) return res; nir_def *is_null = edb_buffer_view_is_null(b, desc); nir_def *oob_alpha = edb_buffer_view_oob_alpha(b, desc); nir_def *a = nir_channel(b, res, 3); a = nir_bcsel(b, nir_ior(b, in_bounds, is_null), a, oob_alpha); return nir_vector_insert_imm(b, res, a, 3); } static void lower_edb_buffer_image_intrin(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { assert(nir_intrinsic_image_dim(intrin) == GLSL_SAMPLER_DIM_BUF); b->cursor = nir_before_instr(&intrin->instr); nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); nir_def *desc = load_resource_deref_desc(b, 4, 32, deref, 0, ctx); switch (intrin->intrinsic) { case nir_intrinsic_image_deref_load: case nir_intrinsic_image_deref_store: case nir_intrinsic_image_deref_atomic: case nir_intrinsic_image_deref_atomic_swap: { nir_def *pos = intrin->src[1].ssa; nir_def *x = nir_channel(b, pos, 0); nir_def *in_bounds = edb_buffer_view_coord_is_in_bounds(b, desc, x); nir_def *index = edb_buffer_view_index(b, desc, in_bounds); nir_def *new_x = adjust_edb_buffer_view_coord(b, desc, x); pos = nir_vector_insert_imm(b, pos, new_x, 0); nir_src_rewrite(&intrin->src[1], pos); if (intrin->intrinsic == nir_intrinsic_image_deref_load) { b->cursor = nir_after_instr(&intrin->instr); nir_def *res = &intrin->def; res = fixup_edb_buffer_view_result(b, desc, in_bounds, res, nir_intrinsic_dest_type(intrin)); nir_def_rewrite_uses_after(&intrin->def, res, res->parent_instr); } nir_rewrite_image_intrinsic(intrin, index, true); break; } case nir_intrinsic_image_deref_size: { assert(intrin->def.num_components == 1); nir_def *size_el = nir_channel(b, desc, 2); nir_def_rewrite_uses(&intrin->def, size_el); break; } default: unreachable("Unknown image intrinsic"); } } static bool lower_image_intrin(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]); if (glsl_get_sampler_dim(deref->type) == GLSL_SAMPLER_DIM_MS) { lower_msaa_image_intrin(b, intrin, ctx); } else if (is_edb_buffer_view(deref, ctx)) { lower_edb_buffer_image_intrin(b, intrin, ctx); } else { b->cursor = nir_before_instr(&intrin->instr); nir_def *desc = load_resource_deref_desc(b, 1, 32, deref, 0, ctx); nir_rewrite_image_intrinsic(intrin, desc, true); } return true; } static bool lower_interp_at_sample(nir_builder *b, nir_intrinsic_instr *interp, const struct lower_descriptors_ctx *ctx) { const uint32_t root_table_offset = nvk_root_descriptor_offset(draw.sample_locations); nir_def *sample = interp->src[1].ssa; b->cursor = nir_before_instr(&interp->instr); nir_def *loc = nir_ldc_nv(b, 1, 64, nir_imm_int(b, 0), /* Root table */ nir_imm_int(b, root_table_offset), .align_mul = 8, .align_offset = 0); /* Yay little endian */ loc = nir_ushr(b, loc, nir_imul_imm(b, sample, 8)); nir_def *loc_x_u4 = nir_iand_imm(b, loc, 0xf); nir_def *loc_y_u4 = nir_iand_imm(b, nir_ushr_imm(b, loc, 4), 0xf); nir_def *loc_u4 = nir_vec2(b, loc_x_u4, loc_y_u4); nir_def *loc_f = nir_fmul_imm(b, nir_i2f32(b, loc_u4), 1.0 / 16.0); nir_def *offset = nir_fadd_imm(b, loc_f, -0.5); assert(interp->intrinsic == nir_intrinsic_interp_deref_at_sample); interp->intrinsic = nir_intrinsic_interp_deref_at_offset; nir_src_rewrite(&interp->src[1], offset); return true; } static bool try_lower_intrin(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { switch (intrin->intrinsic) { case nir_intrinsic_load_constant: return lower_load_constant(b, intrin, ctx); case nir_intrinsic_load_vulkan_descriptor: return try_lower_load_vulkan_descriptor(b, intrin, ctx); case nir_intrinsic_load_workgroup_size: unreachable("Should have been lowered by nir_lower_cs_intrinsics()"); case nir_intrinsic_load_num_workgroups: return lower_sysval_to_root_table(b, intrin, cs.group_count, ctx); case nir_intrinsic_load_base_workgroup_id: return lower_sysval_to_root_table(b, intrin, cs.base_group, ctx); case nir_intrinsic_load_push_constant: return lower_load_push_constant(b, intrin, ctx); case nir_intrinsic_load_base_vertex: case nir_intrinsic_load_first_vertex: return lower_sysval_to_root_table(b, intrin, draw.base_vertex, ctx); case nir_intrinsic_load_base_instance: return lower_sysval_to_root_table(b, intrin, draw.base_instance, ctx); case nir_intrinsic_load_draw_id: return lower_sysval_to_root_table(b, intrin, draw.draw_index, ctx); case nir_intrinsic_load_view_index: return lower_sysval_to_root_table(b, intrin, draw.view_index, ctx); case nir_intrinsic_image_deref_load: case nir_intrinsic_image_deref_sparse_load: case nir_intrinsic_image_deref_store: case nir_intrinsic_image_deref_atomic: case nir_intrinsic_image_deref_atomic_swap: case nir_intrinsic_image_deref_size: case nir_intrinsic_image_deref_samples: return lower_image_intrin(b, intrin, ctx); case nir_intrinsic_interp_deref_at_sample: return lower_interp_at_sample(b, intrin, ctx); default: return false; } } static void lower_edb_buffer_tex_instr(nir_builder *b, nir_tex_instr *tex, const struct lower_descriptors_ctx *ctx) { assert(tex->sampler_dim == GLSL_SAMPLER_DIM_BUF); b->cursor = nir_before_instr(&tex->instr); const int texture_src_idx = nir_tex_instr_src_index(tex, nir_tex_src_texture_deref); nir_deref_instr *texture = nir_src_as_deref(tex->src[texture_src_idx].src); nir_def *plane_ssa = nir_steal_tex_src(tex, nir_tex_src_plane); ASSERTED const uint32_t plane = plane_ssa ? nir_src_as_uint(nir_src_for_ssa(plane_ssa)) : 0; assert(plane == 0); nir_def *desc = load_resource_deref_desc(b, 4, 32, texture, 0, ctx); switch (tex->op) { case nir_texop_txf: { const int coord_src_idx = nir_tex_instr_src_index(tex, nir_tex_src_coord); assert(coord_src_idx >= 0); nir_def *coord = tex->src[coord_src_idx].src.ssa; nir_def *in_bounds = edb_buffer_view_coord_is_in_bounds(b, desc, coord); nir_def *index = edb_buffer_view_index(b, desc, in_bounds); nir_src_rewrite(&tex->src[texture_src_idx].src, index); tex->src[texture_src_idx].src_type = nir_tex_src_texture_handle; nir_def *new_coord = adjust_edb_buffer_view_coord(b, desc, coord); nir_src_rewrite(&tex->src[coord_src_idx].src, new_coord); b->cursor = nir_after_instr(&tex->instr); nir_def *res = &tex->def; res = fixup_edb_buffer_view_result(b, desc, in_bounds, res, tex->dest_type); nir_def_rewrite_uses_after(&tex->def, res, res->parent_instr); break; } case nir_texop_txs: { assert(tex->def.num_components == 1); nir_def *size_el = edb_buffer_view_size_el(b, desc); nir_def_rewrite_uses(&tex->def, size_el); break; } default: unreachable("Invalid buffer texture op"); } } static bool lower_tex(nir_builder *b, nir_tex_instr *tex, const struct lower_descriptors_ctx *ctx) { const int texture_src_idx = nir_tex_instr_src_index(tex, nir_tex_src_texture_deref); const int sampler_src_idx = nir_tex_instr_src_index(tex, nir_tex_src_sampler_deref); if (texture_src_idx < 0) { assert(sampler_src_idx < 0); return false; } nir_deref_instr *texture = nir_src_as_deref(tex->src[texture_src_idx].src); nir_deref_instr *sampler = sampler_src_idx < 0 ? NULL : nir_src_as_deref(tex->src[sampler_src_idx].src); assert(texture); if (is_edb_buffer_view(texture, ctx)) { lower_edb_buffer_tex_instr(b, tex, ctx); return true; } b->cursor = nir_before_instr(&tex->instr); nir_def *plane_ssa = nir_steal_tex_src(tex, nir_tex_src_plane); const uint32_t plane = plane_ssa ? nir_src_as_uint(nir_src_for_ssa(plane_ssa)) : 0; const uint64_t plane_offset_B = plane * sizeof(struct nvk_sampled_image_descriptor); nir_def *texture_desc = load_resource_deref_desc(b, 1, 32, texture, plane_offset_B, ctx); nir_def *combined_handle; if (texture == sampler) { combined_handle = texture_desc; } else { combined_handle = nir_iand_imm(b, texture_desc, NVK_IMAGE_DESCRIPTOR_IMAGE_INDEX_MASK); if (sampler != NULL) { nir_def *sampler_desc = load_resource_deref_desc(b, 1, 32, sampler, plane_offset_B, ctx); nir_def *sampler_index = nir_iand_imm(b, sampler_desc, NVK_IMAGE_DESCRIPTOR_SAMPLER_INDEX_MASK); combined_handle = nir_ior(b, combined_handle, sampler_index); } } /* TODO: The nv50 back-end assumes it's 64-bit because of GL */ combined_handle = nir_u2u64(b, combined_handle); /* TODO: The nv50 back-end assumes it gets handles both places, even for * texelFetch. */ nir_src_rewrite(&tex->src[texture_src_idx].src, combined_handle); tex->src[texture_src_idx].src_type = nir_tex_src_texture_handle; if (sampler_src_idx < 0) { nir_tex_instr_add_src(tex, nir_tex_src_sampler_handle, combined_handle); } else { nir_src_rewrite(&tex->src[sampler_src_idx].src, combined_handle); tex->src[sampler_src_idx].src_type = nir_tex_src_sampler_handle; } /* On pre-Volta hardware, we don't have real null descriptors. Null * descriptors work well enough for sampling but they may not return the * correct query results. */ if (ctx->dev_info->cls_eng3d < VOLTA_A && nir_tex_instr_is_query(tex)) { b->cursor = nir_after_instr(&tex->instr); /* This should get CSE'd with the earlier load */ nir_def *texture_handle = nir_iand_imm(b, texture_desc, NVK_IMAGE_DESCRIPTOR_IMAGE_INDEX_MASK); nir_def *is_null = nir_ieq_imm(b, texture_handle, 0); nir_def *zero = nir_imm_zero(b, tex->def.num_components, tex->def.bit_size); nir_def *res = nir_bcsel(b, is_null, zero, &tex->def); nir_def_rewrite_uses_after(&tex->def, res, res->parent_instr); } return true; } static bool try_lower_descriptors_instr(nir_builder *b, nir_instr *instr, void *_data) { const struct lower_descriptors_ctx *ctx = _data; switch (instr->type) { case nir_instr_type_tex: return lower_tex(b, nir_instr_as_tex(instr), ctx); case nir_instr_type_intrinsic: return try_lower_intrin(b, nir_instr_as_intrinsic(instr), ctx); default: return false; } } #define ROOT_DESC_BASE_ADDR_HI 0x0057de3c static bool lower_ssbo_resource_index(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { if (!descriptor_type_is_ssbo(nir_intrinsic_desc_type(intrin))) return false; b->cursor = nir_instr_remove(&intrin->instr); uint32_t set = nir_intrinsic_desc_set(intrin); uint32_t binding = nir_intrinsic_binding(intrin); nir_def *index = intrin->src[0].ssa; const struct nvk_descriptor_set_binding_layout *binding_layout = get_binding_layout(set, binding, ctx); nir_def *binding_addr; uint8_t binding_stride; switch (binding_layout->type) { case VK_DESCRIPTOR_TYPE_MUTABLE_EXT: case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER: { nir_def *set_addr = load_descriptor_set_addr(b, set, ctx); binding_addr = nir_iadd_imm(b, set_addr, binding_layout->offset); binding_stride = binding_layout->stride; break; } case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC: { nir_def *dynamic_buffer_start = nir_iadd_imm(b, load_dynamic_buffer_start(b, set, ctx), binding_layout->dynamic_buffer_index); nir_def *dynamic_binding_offset = nir_iadd_imm(b, nir_imul_imm(b, dynamic_buffer_start, sizeof(struct nvk_buffer_address)), nvk_root_descriptor_offset(dynamic_buffers)); binding_addr = nir_pack_64_2x32_split(b, dynamic_binding_offset, nir_imm_int(b, ROOT_DESC_BASE_ADDR_HI)); binding_stride = sizeof(struct nvk_buffer_address); break; } default: unreachable("Not an SSBO descriptor"); } /* Tuck the stride in the top 8 bits of the binding address */ binding_addr = nir_ior_imm(b, binding_addr, (uint64_t)binding_stride << 56); const uint32_t binding_size = binding_layout->array_size * binding_stride; nir_def *offset_in_binding = nir_imul_imm(b, index, binding_stride); /* We depend on this when we load descrptors */ assert(binding_layout->array_size >= 1); nir_def *addr; switch (ctx->ssbo_addr_format) { case nir_address_format_64bit_global_32bit_offset: case nir_address_format_64bit_bounded_global: addr = nir_vec4(b, nir_unpack_64_2x32_split_x(b, binding_addr), nir_unpack_64_2x32_split_y(b, binding_addr), nir_imm_int(b, binding_size), offset_in_binding); break; default: unreachable("Unknown address mode"); } nir_def_rewrite_uses(&intrin->def, addr); return true; } static bool lower_ssbo_resource_reindex(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { if (!descriptor_type_is_ssbo(nir_intrinsic_desc_type(intrin))) return false; b->cursor = nir_instr_remove(&intrin->instr); nir_def *addr = intrin->src[0].ssa; nir_def *index = intrin->src[1].ssa; nir_def *addr_high32; switch (ctx->ssbo_addr_format) { case nir_address_format_64bit_global_32bit_offset: case nir_address_format_64bit_bounded_global: addr_high32 = nir_channel(b, addr, 1); break; default: unreachable("Unknown address mode"); } nir_def *stride = nir_ushr_imm(b, addr_high32, 24); nir_def *offset = nir_imul(b, index, stride); addr = nir_build_addr_iadd(b, addr, ctx->ssbo_addr_format, nir_var_mem_ssbo, offset); nir_def_rewrite_uses(&intrin->def, addr); return true; } static bool lower_load_ssbo_descriptor(nir_builder *b, nir_intrinsic_instr *intrin, const struct lower_descriptors_ctx *ctx) { if (!descriptor_type_is_ssbo(nir_intrinsic_desc_type(intrin))) return false; b->cursor = nir_instr_remove(&intrin->instr); nir_def *addr = intrin->src[0].ssa; nir_def *base, *offset, *size = NULL; switch (ctx->ssbo_addr_format) { case nir_address_format_64bit_global_32bit_offset: { base = nir_pack_64_2x32(b, nir_trim_vector(b, addr, 2)); offset = nir_channel(b, addr, 3); break; } case nir_address_format_64bit_bounded_global: { base = nir_pack_64_2x32(b, nir_trim_vector(b, addr, 2)); size = nir_channel(b, addr, 2); offset = nir_channel(b, addr, 3); break; } default: unreachable("Unknown address mode"); } /* Mask off the binding stride */ base = nir_iand_imm(b, base, BITFIELD64_MASK(56)); nir_def *base_lo = nir_unpack_64_2x32_split_x(b, base); nir_def *base_hi = nir_unpack_64_2x32_split_y(b, base); nir_def *desc_root, *desc_global; nir_push_if(b, nir_ieq_imm(b, base_hi, ROOT_DESC_BASE_ADDR_HI)); { desc_root = nir_load_ubo(b, 4, 32, nir_imm_int(b, 0), nir_iadd(b, base_lo, offset), .align_mul = 16, .align_offset = 0, .range = ~0); if (size != NULL) { /* assert(binding_layout->array_size >= 1); */ nir_def *is_oob = nir_ult(b, nir_iadd_imm(b, size, -16), offset); desc_root = nir_bcsel(b, is_oob, nir_imm_zero(b, 4, 32), desc_root); } } nir_push_else(b, NULL); { if (size != NULL) { desc_global = nir_load_global_constant_bounded(b, 4, 32, base, offset, size, .align_mul = 16, .align_offset = 0); } else { desc_global = nir_load_global_constant_offset(b, 4, 32, base, offset, .align_mul = 16, .align_offset = 0); } } nir_pop_if(b, NULL); nir_def *desc = nir_if_phi(b, desc_root, desc_global); nir_def_rewrite_uses(&intrin->def, desc); return true; } static bool lower_ssbo_descriptor_instr(nir_builder *b, nir_instr *instr, void *_data) { const struct lower_descriptors_ctx *ctx = _data; if (instr->type != nir_instr_type_intrinsic) return false; nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr); switch (intrin->intrinsic) { case nir_intrinsic_vulkan_resource_index: return lower_ssbo_resource_index(b, intrin, ctx); case nir_intrinsic_vulkan_resource_reindex: return lower_ssbo_resource_reindex(b, intrin, ctx); case nir_intrinsic_load_vulkan_descriptor: return lower_load_ssbo_descriptor(b, intrin, ctx); default: return false; } } bool nvk_nir_lower_descriptors(nir_shader *nir, const struct nvk_physical_device *pdev, const struct vk_pipeline_robustness_state *rs, uint32_t set_layout_count, struct vk_descriptor_set_layout * const *set_layouts, struct nvk_cbuf_map *cbuf_map_out) { struct lower_descriptors_ctx ctx = { .dev_info = &pdev->info, .use_bindless_cbuf = nvk_use_bindless_cbuf(&pdev->info), .use_edb_buffer_views = nvk_use_edb_buffer_views(pdev), .clamp_desc_array_bounds = rs->storage_buffers != VK_PIPELINE_ROBUSTNESS_BUFFER_BEHAVIOR_DISABLED_EXT || rs->uniform_buffers != VK_PIPELINE_ROBUSTNESS_BUFFER_BEHAVIOR_DISABLED_EXT || rs->images != VK_PIPELINE_ROBUSTNESS_IMAGE_BEHAVIOR_DISABLED_EXT, .ssbo_addr_format = nvk_ssbo_addr_format(pdev, rs), .ubo_addr_format = nvk_ubo_addr_format(pdev, rs), }; assert(set_layout_count <= NVK_MAX_SETS); for (uint32_t s = 0; s < set_layout_count; s++) { if (set_layouts[s] != NULL) ctx.set_layouts[s] = vk_to_nvk_descriptor_set_layout(set_layouts[s]); } /* We run in four passes: * * 1. Find ranges of UBOs that we can promote to bound UBOs. Nothing is * actually lowered in this pass. It's just analysis. * * 2. Try to lower UBO loads to cbufs based on the map we just created. * We need to do this before the main lowering pass because it relies * on the original descriptor load intrinsics. * * 3. Attempt to lower everything with direct descriptors. This may fail * to lower some SSBO intrinsics when variable pointers are used. * * 4. Clean up any SSBO intrinsics which are left and lower them to * slightly less efficient but variable- pointers-correct versions. */ bool pass_lower_ubo = false; if (cbuf_map_out != NULL) { ctx.cbuf_map = cbuf_map_out; build_cbuf_map(nir, &ctx); pass_lower_ubo = nir_shader_intrinsics_pass(nir, lower_load_ubo_intrin, nir_metadata_control_flow, (void *)&ctx); } bool pass_lower_descriptors = nir_shader_instructions_pass(nir, try_lower_descriptors_instr, nir_metadata_control_flow, (void *)&ctx); bool pass_lower_ssbo = nir_shader_instructions_pass(nir, lower_ssbo_descriptor_instr, nir_metadata_control_flow, (void *)&ctx); return pass_lower_ubo || pass_lower_descriptors || pass_lower_ssbo; }