diff --git a/21.2-fixes.patch b/21.2-fixes.patch deleted file mode 100644 index 54b6448..0000000 --- a/21.2-fixes.patch +++ /dev/null @@ -1,4277 +0,0 @@ -diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp -index fa0bf6b65e2..2d1bb85e230 100644 ---- a/src/amd/compiler/aco_instruction_selection.cpp -+++ b/src/amd/compiler/aco_instruction_selection.cpp -@@ -40,6 +40,7 @@ - #include - #include - #include -+#include - #include - - namespace aco { -@@ -3509,6 +3510,14 @@ visit_alu_instr(isel_context* ctx, nir_alu_instr* instr) - case nir_op_fddy_fine: - case nir_op_fddx_coarse: - case nir_op_fddy_coarse: { -+ if (!nir_src_is_divergent(instr->src[0].src)) { -+ /* Source is the same in all lanes, so the derivative is zero. -+ * This also avoids emitting invalid IR. -+ */ -+ bld.copy(Definition(dst), Operand::zero()); -+ break; -+ } -+ - Temp src = as_vgpr(ctx, get_alu_src(ctx, instr->src[0])); - uint16_t dpp_ctrl1, dpp_ctrl2; - if (instr->op == nir_op_fddx_fine) { -diff --git a/src/amd/compiler/aco_lower_to_cssa.cpp b/src/amd/compiler/aco_lower_to_cssa.cpp -index db809867a70..cbd9873c32b 100644 ---- a/src/amd/compiler/aco_lower_to_cssa.cpp -+++ b/src/amd/compiler/aco_lower_to_cssa.cpp -@@ -384,7 +384,7 @@ struct ltg_node { - /* emit the copies in an order that does not - * create interferences within a merge-set */ - void --emit_copies_block(Builder bld, std::map& ltg, RegType type) -+emit_copies_block(Builder& bld, std::map& ltg, RegType type) - { - auto&& it = ltg.begin(); - while (it != ltg.end()) { -@@ -445,6 +445,9 @@ emit_parallelcopies(cssa_ctx& ctx) - continue; - - std::map ltg; -+ bool has_vgpr_copy = false; -+ bool has_sgpr_copy = false; -+ - /* first, try to coalesce all parallelcopies */ - for (const copy& cp : ctx.parallelcopies[i]) { - if (try_coalesce_copy(ctx, cp, i)) { -@@ -459,6 +462,10 @@ emit_parallelcopies(cssa_ctx& ctx) - uint32_t write_idx = ctx.merge_node_table[cp.def.tempId()].index; - assert(write_idx != -1u); - ltg[write_idx] = {cp, read_idx}; -+ -+ bool is_vgpr = cp.def.regClass().type() == RegType::vgpr; -+ has_vgpr_copy |= is_vgpr; -+ has_sgpr_copy |= !is_vgpr; - } - } - -@@ -475,19 +482,23 @@ emit_parallelcopies(cssa_ctx& ctx) - Builder bld(ctx.program); - Block& block = ctx.program->blocks[i]; - -- /* emit VGPR copies */ -- auto IsLogicalEnd = [](const aco_ptr& inst) -> bool -- { return inst->opcode == aco_opcode::p_logical_end; }; -- auto it = std::find_if(block.instructions.rbegin(), block.instructions.rend(), IsLogicalEnd); -- bld.reset(&block.instructions, std::prev(it.base())); -- emit_copies_block(bld, ltg, RegType::vgpr); -- -- /* emit SGPR copies */ -- aco_ptr branch = std::move(block.instructions.back()); -- block.instructions.pop_back(); -- bld.reset(&block.instructions); -- emit_copies_block(bld, ltg, RegType::sgpr); -- bld.insert(std::move(branch)); -+ if (has_vgpr_copy) { -+ /* emit VGPR copies */ -+ auto IsLogicalEnd = [](const aco_ptr& inst) -> bool -+ { return inst->opcode == aco_opcode::p_logical_end; }; -+ auto it = std::find_if(block.instructions.rbegin(), block.instructions.rend(), IsLogicalEnd); -+ bld.reset(&block.instructions, std::prev(it.base())); -+ emit_copies_block(bld, ltg, RegType::vgpr); -+ } -+ -+ if (has_sgpr_copy) { -+ /* emit SGPR copies */ -+ aco_ptr branch = std::move(block.instructions.back()); -+ block.instructions.pop_back(); -+ bld.reset(&block.instructions); -+ emit_copies_block(bld, ltg, RegType::sgpr); -+ bld.insert(std::move(branch)); -+ } - } - - /* finally, rename coalesced phi operands */ -diff --git a/src/amd/compiler/aco_lower_to_hw_instr.cpp b/src/amd/compiler/aco_lower_to_hw_instr.cpp -index 8a9db76abad..1a83dc51512 100644 ---- a/src/amd/compiler/aco_lower_to_hw_instr.cpp -+++ b/src/amd/compiler/aco_lower_to_hw_instr.cpp -@@ -2153,16 +2153,11 @@ lower_to_hw_instr(Program* program) - } - } else { - assert(dst.regClass() == v2b); -- aco_ptr sdwa{create_instruction( -- aco_opcode::v_mov_b32, -- (Format)((uint16_t)Format::VOP1 | (uint16_t)Format::SDWA), 1, 1)}; -- sdwa->operands[0] = op; -- sdwa->definitions[0] = -- Definition(dst.physReg().advance(-dst.physReg().byte()), v1); -- sdwa->sel[0] = sdwa_uword; -- sdwa->dst_sel = sdwa_ubyte0 + dst.physReg().byte() + index; -- sdwa->dst_preserve = 1; -- bld.insert(std::move(sdwa)); -+ Operand sdwa_op = Operand(op.physReg().advance(-op.physReg().byte()), -+ RegClass::get(op.regClass().type(), 4)); -+ bld.vop2_sdwa(aco_opcode::v_lshlrev_b32, dst, Operand::c32(offset), sdwa_op) -+ .instr->sdwa() -+ .sel[1] = sdwa_ubyte0 + op.physReg().byte(); - } - break; - } -diff --git a/src/amd/compiler/aco_optimizer.cpp b/src/amd/compiler/aco_optimizer.cpp -index 3a09e939b9a..05f99c1cb0c 100644 ---- a/src/amd/compiler/aco_optimizer.cpp -+++ b/src/amd/compiler/aco_optimizer.cpp -@@ -929,6 +929,9 @@ apply_extract(opt_ctx& ctx, aco_ptr& instr, unsigned idx, ssa_info& - instr->vop3().opsel |= 1 << idx; - } - -+ instr->operands[idx].set16bit(false); -+ instr->operands[idx].set24bit(false); -+ - ctx.info[tmp.id()].label &= ~label_insert; - /* label_vopc seems to be the only one worth keeping at the moment */ - for (Definition& def : instr->definitions) -@@ -1116,7 +1119,7 @@ label_instruction(opt_ctx& ctx, aco_ptr& instr) - } - unsigned bits = get_operand_size(instr, i); - if (info.is_constant(bits) && alu_can_accept_constant(instr->opcode, i) && -- (!instr->isSDWA() || ctx.program->chip_class >= GFX9)) { -+ (!instr->isSDWA() || ctx.program->chip_class >= GFX9) && !instr->isDPP()) { - Operand op = get_constant_op(ctx, info, bits); - perfwarn(ctx.program, instr->opcode == aco_opcode::v_cndmask_b32 && i == 2, - "v_cndmask_b32 with a constant selector", instr.get()); -@@ -3577,6 +3580,14 @@ combine_instruction(opt_ctx& ctx, aco_ptr& instr) - bool - to_uniform_bool_instr(opt_ctx& ctx, aco_ptr& instr) - { -+ /* Check every operand to make sure they are suitable. */ -+ for (Operand& op : instr->operands) { -+ if (!op.isTemp()) -+ return false; -+ if (!ctx.info[op.tempId()].is_uniform_bool() && !ctx.info[op.tempId()].is_uniform_bitwise()) -+ return false; -+ } -+ - switch (instr->opcode) { - case aco_opcode::s_and_b32: - case aco_opcode::s_and_b64: instr->opcode = aco_opcode::s_and_b32; break; -diff --git a/src/amd/compiler/aco_optimizer_postRA.cpp b/src/amd/compiler/aco_optimizer_postRA.cpp -index d086eff7cef..e19d8e1f6f7 100644 ---- a/src/amd/compiler/aco_optimizer_postRA.cpp -+++ b/src/amd/compiler/aco_optimizer_postRA.cpp -@@ -70,8 +70,9 @@ save_reg_writes(pr_opt_ctx& ctx, aco_ptr& instr) - if (def.regClass().is_subdword()) - idx = clobbered; - -+ assert((r + dw_size) <= max_reg_cnt); - assert(def.size() == dw_size || def.regClass().is_subdword()); -- std::fill(&ctx.instr_idx_by_regs[r], &ctx.instr_idx_by_regs[r + dw_size], idx); -+ std::fill(&ctx.instr_idx_by_regs[r], &ctx.instr_idx_by_regs[r] + dw_size, idx); - } - } - -diff --git a/src/amd/compiler/aco_spill.cpp b/src/amd/compiler/aco_spill.cpp -index 40a4d8c0fc0..205a687b7e4 100644 ---- a/src/amd/compiler/aco_spill.cpp -+++ b/src/amd/compiler/aco_spill.cpp -@@ -193,14 +193,15 @@ next_uses_per_block(spill_ctx& ctx, unsigned block_idx, std::set& work - aco_ptr& instr = block->instructions[idx]; - assert(instr->opcode == aco_opcode::p_linear_phi || instr->opcode == aco_opcode::p_phi); - -- if (!instr->definitions[0].isTemp()) { -- idx--; -- continue; -+ std::pair distance{block_idx, 0}; -+ -+ auto it = instr->definitions[0].isTemp() ? next_uses.find(instr->definitions[0].getTemp()) -+ : next_uses.end(); -+ if (it != next_uses.end()) { -+ distance = it->second; -+ next_uses.erase(it); - } - -- auto it = next_uses.find(instr->definitions[0].getTemp()); -- std::pair distance = -- it == next_uses.end() ? std::make_pair(block_idx, 0u) : it->second; - for (unsigned i = 0; i < instr->operands.size(); i++) { - unsigned pred_idx = - instr->opcode == aco_opcode::p_phi ? block->logical_preds[i] : block->linear_preds[i]; -@@ -212,7 +213,6 @@ next_uses_per_block(spill_ctx& ctx, unsigned block_idx, std::set& work - ctx.next_use_distances_end[pred_idx][instr->operands[i].getTemp()] = distance; - } - } -- next_uses.erase(instr->definitions[0].getTemp()); - idx--; - } - -diff --git a/src/amd/compiler/tests/test_to_hw_instr.cpp b/src/amd/compiler/tests/test_to_hw_instr.cpp -index 4e641111a16..4c9b55a13a2 100644 ---- a/src/amd/compiler/tests/test_to_hw_instr.cpp -+++ b/src/amd/compiler/tests/test_to_hw_instr.cpp -@@ -640,15 +640,15 @@ BEGIN_TEST(to_hw_instr.insert) - //>> p_unit_test 2 - bld.pseudo(aco_opcode::p_unit_test, Operand::c32(2u)); - //~gfx7! v2b: %_:v[0][0:16] = v_bfe_u32 %_:v[1][0:16], 0, 8 -- //~gfx[^7]! v1: %_:v[0] = v_mov_b32 %_:v[1][0:16] dst_sel:ubyte0 dst_preserve -+ //~gfx[^7]! v2b: %0:v[0][0:16] = v_lshlrev_b32 0, %0:v[1][0:7] dst_preserve - INS(0, 0) -- //~gfx[^7]! v1: %_:v[0] = v_mov_b32 %_:v[1][0:16] dst_sel:ubyte2 dst_preserve -+ //~gfx[^7]! v2b: %0:v[0][16:32] = v_lshlrev_b32 0, %0:v[1][0:7] dst_preserve - if (i != GFX7) - INS(0, 2) - //~gfx7! v2b: %_:v[0][0:16] = v_lshlrev_b32 8, %_:v[1][0:16] -- //~gfx[^7]! v1: %_:v[0] = v_mov_b32 %_:v[1][0:16] dst_sel:ubyte1 dst_preserve -+ //~gfx[^7]! v2b: %0:v[0][0:16] = v_lshlrev_b32 8, %0:v[1][0:7] dst_preserve - INS(1, 0) -- //~gfx[^7]! v1: %_:v[0] = v_mov_b32 %_:v[1][0:16] dst_sel:ubyte3 dst_preserve -+ //~gfx[^7]! v2b: %0:v[0][16:32] = v_lshlrev_b32 8, %0:v[1][0:7] dst_preserve - if (i != GFX7) - INS(1, 2) - -diff --git a/src/amd/vulkan/radv_image.c b/src/amd/vulkan/radv_image.c -index f53c46059ea..56dea6c36ff 100644 ---- a/src/amd/vulkan/radv_image.c -+++ b/src/amd/vulkan/radv_image.c -@@ -226,7 +226,7 @@ radv_use_dcc_for_image(struct radv_device *device, const struct radv_image *imag - * decompressing a lot anyway we might as well not have DCC. - */ - if ((pCreateInfo->usage & VK_IMAGE_USAGE_STORAGE_BIT) && -- (!radv_image_use_dcc_image_stores(device, image) || -+ (device->physical_device->rad_info.chip_class < GFX10 || - radv_formats_is_atomic_allowed(pCreateInfo->pNext, format, pCreateInfo->flags))) - return false; - -@@ -276,7 +276,20 @@ radv_use_dcc_for_image(struct radv_device *device, const struct radv_image *imag - bool - radv_image_use_dcc_image_stores(const struct radv_device *device, const struct radv_image *image) - { -- return device->physical_device->rad_info.chip_class >= GFX10; -+ /* DCC image stores is only available for GFX10+. */ -+ if (device->physical_device->rad_info.chip_class < GFX10) -+ return false; -+ -+ if ((device->physical_device->rad_info.family == CHIP_NAVI12 || -+ device->physical_device->rad_info.family == CHIP_NAVI14) && -+ !image->planes[0].surface.u.gfx9.color.dcc.independent_128B_blocks) { -+ /* Do not enable DCC image stores because INDEPENDENT_128B_BLOCKS is required, and 64B is used -+ * for displayable DCC on NAVI12-14. -+ */ -+ return false; -+ } -+ -+ return true; - } - - /* -diff --git a/src/amd/vulkan/radv_meta_bufimage.c b/src/amd/vulkan/radv_meta_bufimage.c -index 36d7637a82a..39136991284 100644 ---- a/src/amd/vulkan/radv_meta_bufimage.c -+++ b/src/amd/vulkan/radv_meta_bufimage.c -@@ -1287,18 +1287,22 @@ fail_itob: - - static void - create_iview(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_blit2d_surf *surf, -- struct radv_image_view *iview) -+ struct radv_image_view *iview, VkFormat format, VkImageAspectFlagBits aspects) - { - VkImageViewType view_type = cmd_buffer->device->physical_device->rad_info.chip_class < GFX9 - ? VK_IMAGE_VIEW_TYPE_2D - : radv_meta_get_view_type(surf->image); -+ -+ if (format == VK_FORMAT_UNDEFINED) -+ format = surf->format; -+ - radv_image_view_init(iview, cmd_buffer->device, - &(VkImageViewCreateInfo){ - .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO, - .image = radv_image_to_handle(surf->image), - .viewType = view_type, -- .format = surf->format, -- .subresourceRange = {.aspectMask = surf->aspect_mask, -+ .format = format, -+ .subresourceRange = {.aspectMask = aspects, - .baseMipLevel = surf->level, - .levelCount = 1, - .baseArrayLayer = surf->layer, -@@ -1439,7 +1443,7 @@ radv_meta_image_to_buffer(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_b - struct radv_image_view src_view; - struct radv_buffer_view dst_view; - -- create_iview(cmd_buffer, src, &src_view); -+ create_iview(cmd_buffer, src, &src_view, VK_FORMAT_UNDEFINED, src->aspect_mask); - create_bview(cmd_buffer, dst->buffer, dst->offset, dst->format, &dst_view); - itob_bind_descriptors(cmd_buffer, &src_view, &dst_view); - -@@ -1585,7 +1589,7 @@ radv_meta_buffer_to_image_cs(struct radv_cmd_buffer *cmd_buffer, - } - - create_bview(cmd_buffer, src->buffer, src->offset, src->format, &src_view); -- create_iview(cmd_buffer, dst, &dst_view); -+ create_iview(cmd_buffer, dst, &dst_view, VK_FORMAT_UNDEFINED, dst->aspect_mask); - btoi_bind_descriptors(cmd_buffer, &src_view, &dst_view); - - if (device->physical_device->rad_info.chip_class >= GFX9 && dst->image->type == VK_IMAGE_TYPE_3D) -@@ -1740,27 +1744,36 @@ radv_meta_image_to_image_cs(struct radv_cmd_buffer *cmd_buffer, struct radv_meta - return; - } - -- create_iview(cmd_buffer, src, &src_view); -- create_iview(cmd_buffer, dst, &dst_view); -- -- itoi_bind_descriptors(cmd_buffer, &src_view, &dst_view); -- -- VkPipeline pipeline = cmd_buffer->device->meta_state.itoi.pipeline[samples_log2]; -- if (device->physical_device->rad_info.chip_class >= GFX9 && -- (src->image->type == VK_IMAGE_TYPE_3D || dst->image->type == VK_IMAGE_TYPE_3D)) -- pipeline = cmd_buffer->device->meta_state.itoi.pipeline_3d; -- radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE, -- pipeline); -- -- for (unsigned r = 0; r < num_rects; ++r) { -- unsigned push_constants[6] = { -- rects[r].src_x, rects[r].src_y, src->layer, rects[r].dst_x, rects[r].dst_y, dst->layer, -- }; -- radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer), -- device->meta_state.itoi.img_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, -- 24, push_constants); -- -- radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1); -+ u_foreach_bit(i, dst->aspect_mask) { -+ unsigned aspect_mask = 1u << i; -+ VkFormat depth_format = 0; -+ if (aspect_mask == VK_IMAGE_ASPECT_STENCIL_BIT) -+ depth_format = vk_format_stencil_only(dst->image->vk_format); -+ else if (aspect_mask == VK_IMAGE_ASPECT_DEPTH_BIT) -+ depth_format = vk_format_depth_only(dst->image->vk_format); -+ -+ create_iview(cmd_buffer, src, &src_view, depth_format, aspect_mask); -+ create_iview(cmd_buffer, dst, &dst_view, depth_format, aspect_mask); -+ -+ itoi_bind_descriptors(cmd_buffer, &src_view, &dst_view); -+ -+ VkPipeline pipeline = cmd_buffer->device->meta_state.itoi.pipeline[samples_log2]; -+ if (device->physical_device->rad_info.chip_class >= GFX9 && -+ (src->image->type == VK_IMAGE_TYPE_3D || dst->image->type == VK_IMAGE_TYPE_3D)) -+ pipeline = cmd_buffer->device->meta_state.itoi.pipeline_3d; -+ radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE, -+ pipeline); -+ -+ for (unsigned r = 0; r < num_rects; ++r) { -+ unsigned push_constants[6] = { -+ rects[r].src_x, rects[r].src_y, src->layer, rects[r].dst_x, rects[r].dst_y, dst->layer, -+ }; -+ radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer), -+ device->meta_state.itoi.img_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, -+ 24, push_constants); -+ -+ radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1); -+ } - } - } - -@@ -1865,7 +1878,7 @@ radv_meta_clear_image_cs(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_bl - return; - } - -- create_iview(cmd_buffer, dst, &dst_iview); -+ create_iview(cmd_buffer, dst, &dst_iview, VK_FORMAT_UNDEFINED, dst->aspect_mask); - cleari_bind_descriptors(cmd_buffer, &dst_iview); - - VkPipeline pipeline = cmd_buffer->device->meta_state.cleari.pipeline[samples_log2]; -diff --git a/src/amd/vulkan/radv_meta_clear.c b/src/amd/vulkan/radv_meta_clear.c -index 65543c0fdf9..ba2f76e3f03 100644 ---- a/src/amd/vulkan/radv_meta_clear.c -+++ b/src/amd/vulkan/radv_meta_clear.c -@@ -986,6 +986,14 @@ radv_can_fast_clear_depth(struct radv_cmd_buffer *cmd_buffer, const struct radv_ - !radv_is_fast_clear_stencil_allowed(clear_value)))) - return false; - -+ if (iview->image->info.levels > 1) { -+ uint32_t last_level = iview->base_mip + iview->level_count - 1; -+ if (last_level >= iview->image->planes[0].surface.num_meta_levels) { -+ /* Do not fast clears if one level can't be fast cleared. */ -+ return false; -+ } -+ } -+ - return true; - } - -diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c -index 3dfaa44c4bc..9fd261a670f 100644 ---- a/src/amd/vulkan/radv_pipeline.c -+++ b/src/amd/vulkan/radv_pipeline.c -@@ -254,7 +254,7 @@ radv_pipeline_init_scratch(const struct radv_device *device, struct radv_pipelin - - max_stage_waves = - MIN2(max_stage_waves, 4 * device->physical_device->rad_info.num_good_compute_units * -- (256 / pipeline->shaders[i]->config.num_vgprs)); -+ radv_get_max_waves(device, pipeline->shaders[i], i)); - max_waves = MAX2(max_waves, max_stage_waves); - } - } -diff --git a/src/amd/vulkan/radv_query.c b/src/amd/vulkan/radv_query.c -index b0b8453cf4c..f5353e8dacb 100644 ---- a/src/amd/vulkan/radv_query.c -+++ b/src/amd/vulkan/radv_query.c -@@ -854,7 +854,8 @@ radv_query_shader(struct radv_cmd_buffer *cmd_buffer, VkPipeline *pipeline, - old_predicating = cmd_buffer->state.predicating; - cmd_buffer->state.predicating = false; - -- struct radv_buffer dst_buffer = {.bo = dst_bo, .offset = dst_offset, .size = dst_stride * count}; -+ uint64_t dst_buffer_size = count == 1 ? src_stride : dst_stride * count; -+ struct radv_buffer dst_buffer = {.bo = dst_bo, .offset = dst_offset, .size = dst_buffer_size}; - - struct radv_buffer src_buffer = { - .bo = src_bo, -diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c -index 03262e6078a..8124bfb6fa3 100644 ---- a/src/amd/vulkan/radv_shader.c -+++ b/src/amd/vulkan/radv_shader.c -@@ -1819,7 +1819,7 @@ radv_get_max_workgroup_size(enum chip_class chip_class, gl_shader_stage stage, - } - - unsigned --radv_get_max_waves(struct radv_device *device, struct radv_shader_variant *variant, -+radv_get_max_waves(const struct radv_device *device, struct radv_shader_variant *variant, - gl_shader_stage stage) - { - struct radeon_info *info = &device->physical_device->rad_info; -diff --git a/src/amd/vulkan/radv_shader.h b/src/amd/vulkan/radv_shader.h -index cabf6845a87..df647a01775 100644 ---- a/src/amd/vulkan/radv_shader.h -+++ b/src/amd/vulkan/radv_shader.h -@@ -469,7 +469,7 @@ struct radv_shader_variant *radv_create_trap_handler_shader(struct radv_device * - - void radv_shader_variant_destroy(struct radv_device *device, struct radv_shader_variant *variant); - --unsigned radv_get_max_waves(struct radv_device *device, struct radv_shader_variant *variant, -+unsigned radv_get_max_waves(const struct radv_device *device, struct radv_shader_variant *variant, - gl_shader_stage stage); - - unsigned radv_get_max_workgroup_size(enum chip_class chip_class, gl_shader_stage stage, -diff --git a/src/broadcom/ci/piglit-v3d-rpi4-fails.txt b/src/broadcom/ci/piglit-v3d-rpi4-fails.txt -index f335cf8f016..e613955739b 100644 ---- a/src/broadcom/ci/piglit-v3d-rpi4-fails.txt -+++ b/src/broadcom/ci/piglit-v3d-rpi4-fails.txt -@@ -208,7 +208,33 @@ spec@ext_framebuffer_object@fbo-blending-formats,Fail - spec@ext_framebuffer_object@fbo-blending-formats@GL_RGB10,Fail - spec@ext_framebuffer_object@getteximage-formats init-by-clear-and-render,Fail - spec@ext_framebuffer_object@getteximage-formats init-by-rendering,Fail --spec@ext_image_dma_buf_import@ext_image_dma_buf_import-export-tex,Fail -+spec@ext_gpu_shader4@execution@texelfetch@fs-texelfetch-isampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetch@fs-texelfetch-sampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetch@fs-texelfetch-usampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetchoffset@fs-texelfetch-isampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetchoffset@fs-texelfetch-sampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetchoffset@fs-texelfetch-usampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetchoffset@vs-texelfetch-isampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetchoffset@vs-texelfetch-sampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetchoffset@vs-texelfetch-usampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetch@vs-texelfetch-isampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetch@vs-texelfetch-sampler1darray,Fail -+spec@ext_gpu_shader4@execution@texelfetch@vs-texelfetch-usampler1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texture() 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texture() 1darrayshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texture(bias) 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texture(bias) 1darrayshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texture() cubeshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturegrad 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturegrad 1darrayshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturegradoffset 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturegradoffset 1darrayshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturelod 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturelod 1darrayshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturelodoffset 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4texturelodoffset 1darrayshadow,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4textureoffset 1darray,Fail -+spec@ext_gpu_shader4@tex-miplevel-selection gpu4textureoffset 1darrayshadow,Fail - spec@ext_packed_depth_stencil@texwrap formats bordercolor,Fail - spec@ext_packed_depth_stencil@texwrap formats bordercolor@GL_DEPTH24_STENCIL8- border color only,Fail - spec@ext_packed_depth_stencil@texwrap formats bordercolor-swizzled,Fail -diff --git a/src/compiler/glsl/ast.h b/src/compiler/glsl/ast.h -index c6b578cb894..0a5b94bb1ce 100644 ---- a/src/compiler/glsl/ast.h -+++ b/src/compiler/glsl/ast.h -@@ -1195,6 +1195,8 @@ public: - ast_node *condition; - ast_expression *rest_expression; - -+ exec_list rest_instructions; -+ - ast_node *body; - - /** -diff --git a/src/compiler/glsl/ast_to_hir.cpp b/src/compiler/glsl/ast_to_hir.cpp -index 370f6934bd4..cd9a16a4754 100644 ---- a/src/compiler/glsl/ast_to_hir.cpp -+++ b/src/compiler/glsl/ast_to_hir.cpp -@@ -1703,6 +1703,7 @@ ast_expression::do_hir(exec_list *instructions, - if ((op[0]->type == glsl_type::error_type || - op[1]->type == glsl_type::error_type)) { - error_emitted = true; -+ result = ir_rvalue::error_value(ctx); - break; - } - -@@ -1740,6 +1741,14 @@ ast_expression::do_hir(exec_list *instructions, - op[0] = this->subexpressions[0]->hir(instructions, state); - op[1] = this->subexpressions[1]->hir(instructions, state); - -+ /* Break out if operand types were not parsed successfully. */ -+ if ((op[0]->type == glsl_type::error_type || -+ op[1]->type == glsl_type::error_type)) { -+ error_emitted = true; -+ result = ir_rvalue::error_value(ctx); -+ break; -+ } -+ - orig_type = op[0]->type; - type = modulus_result_type(op[0], op[1], state, &loc); - -@@ -1770,6 +1779,15 @@ ast_expression::do_hir(exec_list *instructions, - this->subexpressions[0]->set_is_lhs(true); - op[0] = this->subexpressions[0]->hir(instructions, state); - op[1] = this->subexpressions[1]->hir(instructions, state); -+ -+ /* Break out if operand types were not parsed successfully. */ -+ if ((op[0]->type == glsl_type::error_type || -+ op[1]->type == glsl_type::error_type)) { -+ error_emitted = true; -+ result = ir_rvalue::error_value(ctx); -+ break; -+ } -+ - type = shift_result_type(op[0]->type, op[1]->type, this->oper, state, - &loc); - ir_rvalue *temp_rhs = new(ctx) ir_expression(operations[this->oper], -@@ -1790,6 +1808,14 @@ ast_expression::do_hir(exec_list *instructions, - op[0] = this->subexpressions[0]->hir(instructions, state); - op[1] = this->subexpressions[1]->hir(instructions, state); - -+ /* Break out if operand types were not parsed successfully. */ -+ if ((op[0]->type == glsl_type::error_type || -+ op[1]->type == glsl_type::error_type)) { -+ error_emitted = true; -+ result = ir_rvalue::error_value(ctx); -+ break; -+ } -+ - orig_type = op[0]->type; - type = bit_logic_result_type(op[0], op[1], this->oper, state, &loc); - -@@ -6531,8 +6557,8 @@ ast_jump_statement::hir(exec_list *instructions, - if (state->loop_nesting_ast != NULL && - mode == ast_continue && !state->switch_state.is_switch_innermost) { - if (state->loop_nesting_ast->rest_expression) { -- state->loop_nesting_ast->rest_expression->hir(instructions, -- state); -+ clone_ir_list(ctx, instructions, -+ &state->loop_nesting_ast->rest_instructions); - } - if (state->loop_nesting_ast->mode == - ast_iteration_statement::ast_do_while) { -@@ -6780,8 +6806,8 @@ ast_switch_statement::hir(exec_list *instructions, - - if (state->loop_nesting_ast != NULL) { - if (state->loop_nesting_ast->rest_expression) { -- state->loop_nesting_ast->rest_expression->hir(&irif->then_instructions, -- state); -+ clone_ir_list(ctx, &irif->then_instructions, -+ &state->loop_nesting_ast->rest_instructions); - } - if (state->loop_nesting_ast->mode == - ast_iteration_statement::ast_do_while) { -@@ -6830,8 +6856,11 @@ ir_rvalue * - ast_switch_body::hir(exec_list *instructions, - struct _mesa_glsl_parse_state *state) - { -- if (stmts != NULL) -+ if (stmts != NULL) { -+ state->symbols->push_scope(); - stmts->hir(instructions, state); -+ state->symbols->pop_scope(); -+ } - - /* Switch bodies do not have r-values. */ - return NULL; -@@ -7135,11 +7164,21 @@ ast_iteration_statement::hir(exec_list *instructions, - if (mode != ast_do_while) - condition_to_hir(&stmt->body_instructions, state); - -- if (body != NULL) -+ if (rest_expression != NULL) -+ rest_expression->hir(&rest_instructions, state); -+ -+ if (body != NULL) { -+ if (mode == ast_do_while) -+ state->symbols->push_scope(); -+ - body->hir(& stmt->body_instructions, state); - -+ if (mode == ast_do_while) -+ state->symbols->pop_scope(); -+ } -+ - if (rest_expression != NULL) -- rest_expression->hir(& stmt->body_instructions, state); -+ stmt->body_instructions.append_list(&rest_instructions); - - if (mode == ast_do_while) - condition_to_hir(&stmt->body_instructions, state); -diff --git a/src/compiler/glsl/glsl_parser.yy b/src/compiler/glsl/glsl_parser.yy -index ec66e680a2b..4111c45c97d 100644 ---- a/src/compiler/glsl/glsl_parser.yy -+++ b/src/compiler/glsl/glsl_parser.yy -@@ -2743,7 +2743,7 @@ iteration_statement: - NULL, $3, NULL, $5); - $$->set_location_range(@1, @4); - } -- | DO statement WHILE '(' expression ')' ';' -+ | DO statement_no_new_scope WHILE '(' expression ')' ';' - { - void *ctx = state->linalloc; - $$ = new(ctx) ast_iteration_statement(ast_iteration_statement::ast_do_while, -diff --git a/src/compiler/glsl/link_varyings.cpp b/src/compiler/glsl/link_varyings.cpp -index 9954a731479..997513a636d 100644 ---- a/src/compiler/glsl/link_varyings.cpp -+++ b/src/compiler/glsl/link_varyings.cpp -@@ -660,9 +660,11 @@ validate_explicit_variable_location(struct gl_context *ctx, - glsl_struct_field *field = &type_without_array->fields.structure[i]; - unsigned field_location = field->location - - (field->patch ? VARYING_SLOT_PATCH0 : VARYING_SLOT_VAR0); -+ unsigned field_slots = field->type->count_attribute_slots(false); - if (!check_location_aliasing(explicit_locations, var, - field_location, -- 0, field_location + 1, -+ 0, -+ field_location + field_slots, - field->type, - field->interpolation, - field->centroid, -diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h -index 3de53b51388..dd6cb11094e 100644 ---- a/src/compiler/nir/nir.h -+++ b/src/compiler/nir/nir.h -@@ -2749,6 +2749,32 @@ nir_block_ends_in_jump(nir_block *block) - nir_block_last_instr(block)->type == nir_instr_type_jump; - } - -+static inline bool -+nir_block_ends_in_return_or_halt(nir_block *block) -+{ -+ if (exec_list_is_empty(&block->instr_list)) -+ return false; -+ -+ nir_instr *instr = nir_block_last_instr(block); -+ if (instr->type != nir_instr_type_jump) -+ return false; -+ -+ nir_jump_instr *jump_instr = nir_instr_as_jump(instr); -+ return jump_instr->type == nir_jump_return || -+ jump_instr->type == nir_jump_halt; -+} -+ -+static inline bool -+nir_block_ends_in_break(nir_block *block) -+{ -+ if (exec_list_is_empty(&block->instr_list)) -+ return false; -+ -+ nir_instr *instr = nir_block_last_instr(block); -+ return instr->type == nir_instr_type_jump && -+ nir_instr_as_jump(instr)->type == nir_jump_break; -+} -+ - #define nir_foreach_instr(instr, block) \ - foreach_list_typed(nir_instr, instr, node, &(block)->instr_list) - #define nir_foreach_instr_reverse(instr, block) \ -diff --git a/src/compiler/nir/nir_constant_expressions.py b/src/compiler/nir/nir_constant_expressions.py -index 606e974353f..c1097de7c67 100644 ---- a/src/compiler/nir/nir_constant_expressions.py -+++ b/src/compiler/nir/nir_constant_expressions.py -@@ -68,8 +68,6 @@ template = """\ - #include "util/bigmath.h" - #include "nir_constant_expressions.h" - --#define MAX_UINT_FOR_SIZE(bits) (UINT64_MAX >> (64 - (bits))) -- - /** - * \brief Checks if the provided value is a denorm and flushes it to zero. - */ -diff --git a/src/compiler/nir/nir_loop_analyze.h b/src/compiler/nir/nir_loop_analyze.h -index 7b4ed66ee58..18c23051717 100644 ---- a/src/compiler/nir/nir_loop_analyze.h -+++ b/src/compiler/nir/nir_loop_analyze.h -@@ -92,15 +92,4 @@ nir_is_trivial_loop_if(nir_if *nif, nir_block *break_block) - return true; - } - --static inline bool --nir_block_ends_in_break(nir_block *block) --{ -- if (exec_list_is_empty(&block->instr_list)) -- return false; -- -- nir_instr *instr = nir_block_last_instr(block); -- return instr->type == nir_instr_type_jump && -- nir_instr_as_jump(instr)->type == nir_jump_break; --} -- - #endif /* NIR_LOOP_ANALYZE_H */ -diff --git a/src/compiler/nir/nir_lower_bit_size.c b/src/compiler/nir/nir_lower_bit_size.c -index 5473ea7c0c5..2c082f71f38 100644 ---- a/src/compiler/nir/nir_lower_bit_size.c -+++ b/src/compiler/nir/nir_lower_bit_size.c -@@ -74,13 +74,30 @@ lower_alu_instr(nir_builder *bld, nir_alu_instr *alu, unsigned bit_size) - nir_ssa_def *lowered_dst = NULL; - if (op == nir_op_imul_high || op == nir_op_umul_high) { - assert(dst_bit_size * 2 <= bit_size); -- nir_ssa_def *lowered_dst = nir_imul(bld, srcs[0], srcs[1]); -+ lowered_dst = nir_imul(bld, srcs[0], srcs[1]); - if (nir_op_infos[op].output_type & nir_type_uint) - lowered_dst = nir_ushr_imm(bld, lowered_dst, dst_bit_size); - else - lowered_dst = nir_ishr_imm(bld, lowered_dst, dst_bit_size); - } else { - lowered_dst = nir_build_alu_src_arr(bld, op, srcs); -+ -+ /* The add_sat and sub_sat instructions need to clamp the result to the -+ * range of the original type. -+ */ -+ if (op == nir_op_iadd_sat || op == nir_op_isub_sat) { -+ const int64_t int_max = u_intN_max(dst_bit_size); -+ const int64_t int_min = u_intN_min(dst_bit_size); -+ -+ lowered_dst = nir_iclamp(bld, lowered_dst, -+ nir_imm_intN_t(bld, int_min, bit_size), -+ nir_imm_intN_t(bld, int_max, bit_size)); -+ } else if (op == nir_op_uadd_sat || op == nir_op_usub_sat) { -+ const uint64_t uint_max = u_uintN_max(dst_bit_size); -+ -+ lowered_dst = nir_umin(bld, lowered_dst, -+ nir_imm_intN_t(bld, uint_max, bit_size)); -+ } - } - - -diff --git a/src/compiler/nir/nir_lower_io_to_vector.c b/src/compiler/nir/nir_lower_io_to_vector.c -index 433e5ccff10..13d692e72be 100644 ---- a/src/compiler/nir/nir_lower_io_to_vector.c -+++ b/src/compiler/nir/nir_lower_io_to_vector.c -@@ -632,6 +632,7 @@ nir_vectorize_tess_levels_impl(nir_function_impl *impl) - } else { - b.cursor = nir_after_instr(instr); - nir_ssa_def *val = &intrin->dest.ssa; -+ val->num_components = intrin->num_components; - nir_ssa_def *comp = nir_channel(&b, val, index); - nir_ssa_def_rewrite_uses_after(val, comp, comp->parent_instr); - } -diff --git a/src/compiler/nir/nir_opcodes.py b/src/compiler/nir/nir_opcodes.py -index f9330c0af2c..872116c7ed3 100644 ---- a/src/compiler/nir/nir_opcodes.py -+++ b/src/compiler/nir/nir_opcodes.py -@@ -487,7 +487,7 @@ for (int bit = 31; bit >= 0; bit--) { - } - """) - --unop_convert("ifind_msb_rev", tint32, tuint, """ -+unop_convert("ifind_msb_rev", tint32, tint, """ - dst = -1; - if (src0 != 0 && src0 != -1) { - for (int bit = 0; bit < 31; bit++) { -@@ -634,7 +634,7 @@ binop("iadd_sat", tint, _2src_commutative, """ - (src0 < src0 + src1 ? (1ull << (bit_size - 1)) : src0 + src1) - """) - binop("uadd_sat", tuint, _2src_commutative, -- "(src0 + src1) < src0 ? MAX_UINT_FOR_SIZE(sizeof(src0) * 8) : (src0 + src1)") -+ "(src0 + src1) < src0 ? u_uintN_max(sizeof(src0) * 8) : (src0 + src1)") - binop("isub_sat", tint, "", """ - src1 < 0 ? - (src0 - src1 < src0 ? (1ull << (bit_size - 1)) - 1 : src0 - src1) : -diff --git a/src/compiler/nir/nir_opt_peephole_select.c b/src/compiler/nir/nir_opt_peephole_select.c -index 5eeb5f66b94..72b6249cf43 100644 ---- a/src/compiler/nir/nir_opt_peephole_select.c -+++ b/src/compiler/nir/nir_opt_peephole_select.c -@@ -381,6 +381,17 @@ nir_opt_peephole_select_block(nir_block *block, nir_shader *shader, - if (prev_node->type != nir_cf_node_if) - return false; - -+ nir_block *prev_block = nir_cf_node_as_block(nir_cf_node_prev(prev_node)); -+ -+ /* If the last instruction before this if/else block is a jump, we can't -+ * append stuff after it because it would break a bunch of assumption about -+ * control flow (nir_validate expects the successor of a return/halt jump -+ * to be the end of the function, which might not match the successor of -+ * the if/else blocks). -+ */ -+ if (nir_block_ends_in_return_or_halt(prev_block)) -+ return false; -+ - nir_if *if_stmt = nir_cf_node_as_if(prev_node); - - /* first, try to collapse the if */ -@@ -422,8 +433,6 @@ nir_opt_peephole_select_block(nir_block *block, nir_shader *shader, - * selects. - */ - -- nir_block *prev_block = nir_cf_node_as_block(nir_cf_node_prev(prev_node)); -- - /* First, we move the remaining instructions from the blocks to the - * block before. We have already guaranteed that this is safe by - * calling block_check_for_allowed_instrs() -diff --git a/src/compiler/nir/nir_range_analysis.c b/src/compiler/nir/nir_range_analysis.c -index 4e37881526f..6b4d86c1bbf 100644 ---- a/src/compiler/nir/nir_range_analysis.c -+++ b/src/compiler/nir/nir_range_analysis.c -@@ -1292,7 +1292,15 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht, - nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(scalar.def->parent_instr); - switch (intrin->intrinsic) { - case nir_intrinsic_load_local_invocation_index: -- if (shader->info.workgroup_size_variable) { -+ /* The local invocation index is used under the hood by RADV for -+ * some non-compute-like shaders (eg. LS and NGG). These technically -+ * run in workgroups on the HW, even though this fact is not exposed -+ * by the API. -+ * They can safely use the same code path here as variable sized -+ * compute-like shader stages. -+ */ -+ if (!gl_shader_stage_uses_workgroup(shader->info.stage) || -+ shader->info.workgroup_size_variable) { - res = config->max_workgroup_invocations - 1; - } else { - res = (shader->info.workgroup_size[0] * -diff --git a/src/compiler/nir/nir_search.c b/src/compiler/nir/nir_search.c -index 272a3f6444a..437a24b9b02 100644 ---- a/src/compiler/nir/nir_search.c -+++ b/src/compiler/nir/nir_search.c -@@ -368,7 +368,7 @@ match_value(const nir_search_value *value, nir_alu_instr *instr, unsigned src, - case nir_type_uint: - case nir_type_bool: { - unsigned bit_size = nir_src_bit_size(instr->src[src].src); -- uint64_t mask = bit_size == 64 ? UINT64_MAX : (1ull << bit_size) - 1; -+ uint64_t mask = u_uintN_max(bit_size); - for (unsigned i = 0; i < num_components; ++i) { - uint64_t val = nir_src_comp_as_uint(instr->src[src].src, - new_swizzle[i]); -diff --git a/src/compiler/nir/nir_serialize.c b/src/compiler/nir/nir_serialize.c -index cc376cbcb6b..cce2d2e83ad 100644 ---- a/src/compiler/nir/nir_serialize.c -+++ b/src/compiler/nir/nir_serialize.c -@@ -1814,6 +1814,7 @@ static void - write_if(write_ctx *ctx, nir_if *nif) - { - write_src(ctx, &nif->condition); -+ blob_write_uint8(ctx->blob, nif->control); - - write_cf_list(ctx, &nif->then_list); - write_cf_list(ctx, &nif->else_list); -@@ -1825,6 +1826,7 @@ read_if(read_ctx *ctx, struct exec_list *cf_list) - nir_if *nif = nir_if_create(ctx->nir); - - read_src(ctx, &nif->condition, nif); -+ nif->control = blob_read_uint8(ctx->blob); - - nir_cf_node_insert_end(cf_list, &nif->cf_node); - -@@ -1835,6 +1837,7 @@ read_if(read_ctx *ctx, struct exec_list *cf_list) - static void - write_loop(write_ctx *ctx, nir_loop *loop) - { -+ blob_write_uint8(ctx->blob, loop->control); - write_cf_list(ctx, &loop->body); - } - -@@ -1845,6 +1848,7 @@ read_loop(read_ctx *ctx, struct exec_list *cf_list) - - nir_cf_node_insert_end(cf_list, &loop->cf_node); - -+ loop->control = blob_read_uint8(ctx->blob); - read_cf_list(ctx, &loop->body); - } - -diff --git a/src/freedreno/ir3/ir3_lower_parallelcopy.c b/src/freedreno/ir3/ir3_lower_parallelcopy.c -index 81087d694ef..8807dd2d157 100644 ---- a/src/freedreno/ir3/ir3_lower_parallelcopy.c -+++ b/src/freedreno/ir3/ir3_lower_parallelcopy.c -@@ -282,7 +282,7 @@ static void - split_32bit_copy(struct copy_ctx *ctx, struct copy_entry *entry) - { - assert(!entry->done); -- assert(!(entry->flags & (IR3_REG_IMMED | IR3_REG_CONST))); -+ assert(!(entry->src.flags & (IR3_REG_IMMED | IR3_REG_CONST))); - assert(copy_entry_size(entry) == 2); - struct copy_entry *new_entry = &ctx->entries[ctx->entry_count++]; - -@@ -362,7 +362,7 @@ _handle_copies(struct ir3_compiler *compiler, struct ir3_instruction *instr, - - if (((ctx->physreg_use_count[entry->dst] == 0 || - ctx->physreg_use_count[entry->dst + 1] == 0)) && -- !(entry->flags & (IR3_REG_IMMED | IR3_REG_CONST))) { -+ !(entry->src.flags & (IR3_REG_IMMED | IR3_REG_CONST))) { - split_32bit_copy(ctx, entry); - progress = true; - } -@@ -451,6 +451,8 @@ _handle_copies(struct ir3_compiler *compiler, struct ir3_instruction *instr, - entry->src.reg + (blocking->src.reg - entry->dst); - } - } -+ -+ entry->done = true; - } - } - -diff --git a/src/freedreno/ir3/ir3_ra.c b/src/freedreno/ir3/ir3_ra.c -index 6463b62ed2a..b92155cdc3b 100644 ---- a/src/freedreno/ir3/ir3_ra.c -+++ b/src/freedreno/ir3/ir3_ra.c -@@ -548,6 +548,18 @@ ra_file_mark_killed(struct ra_file *file, struct ra_interval *interval) - interval->is_killed = true; - } - -+static void -+ra_file_unmark_killed(struct ra_file *file, struct ra_interval *interval) -+{ -+ assert(!interval->interval.parent); -+ -+ for (physreg_t i = interval->physreg_start; i < interval->physreg_end; i++) { -+ BITSET_CLEAR(file->available, i); -+ } -+ -+ interval->is_killed = false; -+} -+ - static physreg_t - ra_interval_get_physreg(const struct ra_interval *interval) - { -@@ -950,6 +962,12 @@ static physreg_t - find_best_gap(struct ra_file *file, unsigned file_size, unsigned size, - unsigned align, bool is_source) - { -+ /* This can happen if we create a very large merge set. Just bail out in that -+ * case. -+ */ -+ if (size > file_size) -+ return (physreg_t) ~0; -+ - BITSET_WORD *available = - is_source ? file->available_to_evict : file->available; - -@@ -1311,15 +1329,11 @@ handle_collect(struct ra_ctx *ctx, struct ir3_instruction *instr) - */ - physreg_t dst_fixed = (physreg_t)~0u; - -- for (unsigned i = 0; i < instr->srcs_count; i++) { -- if (!ra_reg_is_src(instr->srcs[i])) -- continue; -- -- if (instr->srcs[i]->flags & IR3_REG_FIRST_KILL) { -- mark_src_killed(ctx, instr->srcs[i]); -+ ra_foreach_src (src, instr) { -+ if (src->flags & IR3_REG_FIRST_KILL) { -+ mark_src_killed(ctx, src); - } - -- struct ir3_register *src = instr->srcs[i]; - struct ra_interval *interval = &ctx->intervals[src->def->name]; - - if (src->def->merge_set != dst_set || interval->is_killed) -@@ -1347,11 +1361,7 @@ handle_collect(struct ra_ctx *ctx, struct ir3_instruction *instr) - allocate_dst(ctx, instr->dsts[0]); - - /* Remove the temporary is_killed we added */ -- for (unsigned i = 0; i < instr->srcs_count; i++) { -- if (!ra_reg_is_src(instr->srcs[i])) -- continue; -- -- struct ir3_register *src = instr->srcs[i]; -+ ra_foreach_src (src, instr) { - struct ra_interval *interval = &ctx->intervals[src->def->name]; - while (interval->interval.parent != NULL) { - interval = ir3_reg_interval_to_ra_interval(interval->interval.parent); -@@ -1359,8 +1369,9 @@ handle_collect(struct ra_ctx *ctx, struct ir3_instruction *instr) - - /* Filter out cases where it actually should be killed */ - if (interval != &ctx->intervals[src->def->name] || -- !(src->flags & IR3_REG_KILL)) -- interval->is_killed = false; -+ !(src->flags & IR3_REG_KILL)) { -+ ra_file_unmark_killed(ra_get_file(ctx, src), interval); -+ } - } - - ra_foreach_src_rev (src, instr) { -diff --git a/src/freedreno/vulkan/tu_cmd_buffer.c b/src/freedreno/vulkan/tu_cmd_buffer.c -index ee69aa6d94c..a71a7d9cc47 100644 ---- a/src/freedreno/vulkan/tu_cmd_buffer.c -+++ b/src/freedreno/vulkan/tu_cmd_buffer.c -@@ -1497,10 +1497,6 @@ tu_BeginCommandBuffer(VkCommandBuffer commandBuffer, - memset(&cmd_buffer->state, 0, sizeof(cmd_buffer->state)); - cmd_buffer->state.index_size = 0xff; /* dirty restart index */ - -- cmd_buffer->state.last_vs_params.first_instance = -1; -- cmd_buffer->state.last_vs_params.params_offset = -1; -- cmd_buffer->state.last_vs_params.vertex_offset = -1; -- - tu_cache_init(&cmd_buffer->state.cache); - tu_cache_init(&cmd_buffer->state.renderpass_cache); - cmd_buffer->usage_flags = pBeginInfo->flags; -@@ -2092,7 +2088,8 @@ tu_CmdBindPipeline(VkCommandBuffer commandBuffer, - assert(pipelineBindPoint == VK_PIPELINE_BIND_POINT_GRAPHICS); - - cmd->state.pipeline = pipeline; -- cmd->state.dirty |= TU_CMD_DIRTY_DESC_SETS_LOAD | TU_CMD_DIRTY_SHADER_CONSTS | TU_CMD_DIRTY_LRZ; -+ cmd->state.dirty |= TU_CMD_DIRTY_DESC_SETS_LOAD | TU_CMD_DIRTY_SHADER_CONSTS | -+ TU_CMD_DIRTY_LRZ | TU_CMD_DIRTY_VS_PARAMS; - - /* note: this also avoids emitting draw states before renderpass clears, - * which may use the 3D clear path (for MSAA cases) -@@ -3854,14 +3851,17 @@ tu6_emit_vs_params(struct tu_cmd_buffer *cmd, - uint32_t vertex_offset, - uint32_t first_instance) - { -- uint32_t offset = vs_params_offset(cmd); -- -- if (offset == cmd->state.last_vs_params.params_offset && -+ /* Beside re-emitting params when they are changed, we should re-emit -+ * them after constants are invalidated via HLSQ_INVALIDATE_CMD. -+ */ -+ if (!(cmd->state.dirty & (TU_CMD_DIRTY_DRAW_STATE | TU_CMD_DIRTY_VS_PARAMS)) && - vertex_offset == cmd->state.last_vs_params.vertex_offset && - first_instance == cmd->state.last_vs_params.first_instance) { - return; - } - -+ uint32_t offset = vs_params_offset(cmd); -+ - struct tu_cs cs; - VkResult result = tu_cs_begin_sub_stream(&cmd->sub_cs, 3 + (offset ? 8 : 0), &cs); - if (result != VK_SUCCESS) { -@@ -3889,7 +3889,6 @@ tu6_emit_vs_params(struct tu_cmd_buffer *cmd, - tu_cs_emit(&cs, 0); - } - -- cmd->state.last_vs_params.params_offset = offset; - cmd->state.last_vs_params.vertex_offset = vertex_offset; - cmd->state.last_vs_params.first_instance = first_instance; - -diff --git a/src/freedreno/vulkan/tu_private.h b/src/freedreno/vulkan/tu_private.h -index 683eeb89725..51bd33a8e66 100644 ---- a/src/freedreno/vulkan/tu_private.h -+++ b/src/freedreno/vulkan/tu_private.h -@@ -896,7 +896,6 @@ struct tu_lrz_state - }; - - struct tu_vs_params { -- uint32_t params_offset; - uint32_t vertex_offset; - uint32_t first_instance; - }; -diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir.c b/src/gallium/auxiliary/gallivm/lp_bld_nir.c -index 6795090053e..249e8f63a8e 100644 ---- a/src/gallium/auxiliary/gallivm/lp_bld_nir.c -+++ b/src/gallium/auxiliary/gallivm/lp_bld_nir.c -@@ -2110,6 +2110,26 @@ static void visit_tex(struct lp_build_nir_context *bld_base, nir_tex_instr *inst - params.lod = explicit_lod; - params.ms_index = ms_index; - bld_base->tex(bld_base, ¶ms); -+ -+ if (nir_dest_bit_size(instr->dest) != 32) { -+ assert(nir_dest_bit_size(instr->dest) == 16); -+ LLVMTypeRef vec_type; -+ switch (nir_alu_type_get_base_type(instr->dest_type)) { -+ case nir_type_int: -+ vec_type = bld_base->int16_bld.vec_type; -+ break; -+ case nir_type_uint: -+ vec_type = bld_base->uint16_bld.vec_type; -+ break; -+ default: -+ unreachable("unexpected alu type"); -+ } -+ for (int i = 0; i < nir_dest_num_components(instr->dest); ++i) { -+ texel[i] = LLVMBuildBitCast(builder, texel[i], bld_base->int_bld.vec_type, ""); -+ texel[i] = LLVMBuildTrunc(builder, texel[i], vec_type, ""); -+ } -+ } -+ - assign_dest(bld_base, &instr->dest, texel); - } - -diff --git a/src/gallium/auxiliary/indices/u_primconvert.c b/src/gallium/auxiliary/indices/u_primconvert.c -index d8704237e49..62956910aa8 100644 ---- a/src/gallium/auxiliary/indices/u_primconvert.c -+++ b/src/gallium/auxiliary/indices/u_primconvert.c -@@ -179,10 +179,12 @@ util_primconvert_draw_vbo(struct primconvert_context *pc, - src = (const uint8_t *)src; - - /* if the resulting primitive type is not supported by the driver for primitive restart, -+ * or if the original primitive type was not supported by the driver, - * the draw needs to be rewritten to not use primitive restart - */ - if (info->primitive_restart && -- !(pc->cfg.restart_primtypes_mask & BITFIELD_BIT(mode))) { -+ (!(pc->cfg.restart_primtypes_mask & BITFIELD_BIT(mode)) || -+ !(pc->cfg.primtypes_mask & BITFIELD_BIT(info->mode)))) { - /* step 1: rewrite draw to not use primitive primitive restart; - * this pre-filters degenerate primitives - */ -diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c -index 24a18ec904f..3bb5f1f8bae 100644 ---- a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c -+++ b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c -@@ -579,13 +579,13 @@ void nir_tgsi_scan_shader(const struct nir_shader *nir, - info->indirect_files |= 1 << TGSI_FILE_INPUT; - info->file_max[TGSI_FILE_INPUT] = info->num_inputs - 1; - } else { -- int max = -1; -+ int max = info->file_max[TGSI_FILE_INPUT] = -1; - nir_foreach_shader_in_variable(var, nir) { -- int slots = glsl_count_attribute_slots(var->type, false); -- int tmax = var->data.driver_location + slots - 1; -- if (tmax > max) -- max = tmax; -- info->file_max[TGSI_FILE_INPUT] = max; -+ int slots = glsl_count_attribute_slots(var->type, false); -+ int tmax = var->data.driver_location + slots - 1; -+ if (tmax > max) -+ max = tmax; -+ info->file_max[TGSI_FILE_INPUT] = max; - } - } - -diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c -index e60b92e867d..3e9c63bc444 100644 ---- a/src/gallium/auxiliary/nir/tgsi_to_nir.c -+++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c -@@ -430,6 +430,8 @@ ttn_emit_declaration(struct ttn_compile *c) - if (var->data.location == VARYING_SLOT_FOGC || - var->data.location == VARYING_SLOT_PSIZ) { - var->type = glsl_float_type(); -+ } else if (var->data.location == VARYING_SLOT_LAYER) { -+ var->type = glsl_int_type(); - } - } - -@@ -2220,8 +2222,9 @@ ttn_add_output_stores(struct ttn_compile *c) - else if (var->data.location == FRAG_RESULT_SAMPLE_MASK) - store_value = nir_channel(b, store_value, 0); - } else { -- /* FOGC and PSIZ are scalar values */ -+ /* FOGC, LAYER, and PSIZ are scalar values */ - if (var->data.location == VARYING_SLOT_FOGC || -+ var->data.location == VARYING_SLOT_LAYER || - var->data.location == VARYING_SLOT_PSIZ) { - store_value = nir_channel(b, store_value, 0); - } -diff --git a/src/gallium/drivers/crocus/crocus_context.h b/src/gallium/drivers/crocus/crocus_context.h -index 15b41079ce1..a907af36665 100644 ---- a/src/gallium/drivers/crocus/crocus_context.h -+++ b/src/gallium/drivers/crocus/crocus_context.h -@@ -587,6 +587,7 @@ struct crocus_context { - - bool primitive_restart; - unsigned cut_index; -+ enum pipe_prim_type reduced_prim_mode:8; - enum pipe_prim_type prim_mode:8; - bool prim_is_points_or_lines; - uint8_t vertices_per_patch; -diff --git a/src/gallium/drivers/crocus/crocus_draw.c b/src/gallium/drivers/crocus/crocus_draw.c -index feef4e78ecf..cdfe6a63b26 100644 ---- a/src/gallium/drivers/crocus/crocus_draw.c -+++ b/src/gallium/drivers/crocus/crocus_draw.c -@@ -139,11 +139,18 @@ crocus_update_draw_info(struct crocus_context *ice, - if (ice->state.prim_mode != mode) { - ice->state.prim_mode = mode; - -+ enum pipe_prim_type reduced = u_reduced_prim(mode); -+ if (ice->state.reduced_prim_mode != reduced) { -+ if (screen->devinfo.ver < 6) -+ ice->state.dirty |= CROCUS_DIRTY_GEN4_CLIP_PROG | CROCUS_DIRTY_GEN4_SF_PROG; -+ /* if the reduced prim changes the WM needs updating. */ -+ ice->state.stage_dirty |= CROCUS_STAGE_DIRTY_UNCOMPILED_FS; -+ ice->state.reduced_prim_mode = reduced; -+ } -+ - if (screen->devinfo.ver == 8) - ice->state.dirty |= CROCUS_DIRTY_GEN8_VF_TOPOLOGY; - -- if (screen->devinfo.ver < 6) -- ice->state.dirty |= CROCUS_DIRTY_GEN4_CLIP_PROG | CROCUS_DIRTY_GEN4_SF_PROG; - if (screen->devinfo.ver <= 6) - ice->state.dirty |= CROCUS_DIRTY_GEN4_FF_GS_PROG; - -diff --git a/src/gallium/drivers/crocus/crocus_program.c b/src/gallium/drivers/crocus/crocus_program.c -index d96edb7152d..4da1970e8cb 100644 ---- a/src/gallium/drivers/crocus/crocus_program.c -+++ b/src/gallium/drivers/crocus/crocus_program.c -@@ -2082,7 +2082,7 @@ crocus_update_compiled_clip(struct crocus_context *ice) - memcpy(key.interp_mode, wm_prog_data->interp_mode, sizeof(key.interp_mode)); - } - -- key.primitive = u_reduced_prim(ice->state.prim_mode); -+ key.primitive = ice->state.reduced_prim_mode; - key.attrs = ice->shaders.last_vue_map->slots_valid; - - struct pipe_rasterizer_state *rs_state = crocus_get_rast_state(ice); -@@ -2230,7 +2230,7 @@ crocus_update_compiled_sf(struct crocus_context *ice) - - key.attrs = ice->shaders.last_vue_map->slots_valid; - -- switch (u_reduced_prim(ice->state.prim_mode)) { -+ switch (ice->state.reduced_prim_mode) { - case GL_TRIANGLES: - default: - if (key.attrs & BITFIELD64_BIT(VARYING_SLOT_EDGE)) -diff --git a/src/gallium/drivers/crocus/crocus_state.c b/src/gallium/drivers/crocus/crocus_state.c -index 07f72c3609e..1df647d81c4 100644 ---- a/src/gallium/drivers/crocus/crocus_state.c -+++ b/src/gallium/drivers/crocus/crocus_state.c -@@ -2026,6 +2026,9 @@ crocus_create_rasterizer_state(struct pipe_context *ctx, - sf.LineEndCapAntialiasingRegionWidth = - state->line_smooth ? _10pixels : _05pixels; - sf.LastPixelEnable = state->line_last_pixel; -+#if GFX_VER <= 7 -+ sf.AntialiasingEnable = state->line_smooth; -+#endif - #if GFX_VER == 8 - struct crocus_screen *screen = (struct crocus_screen *)ctx->screen; - if (screen->devinfo.is_cherryview) -@@ -4774,7 +4777,7 @@ crocus_populate_fs_key(const struct crocus_context *ice, - - uint32_t line_aa = BRW_WM_AA_NEVER; - if (rast->cso.line_smooth) { -- int reduced_prim = u_reduced_prim(ice->state.prim_mode); -+ int reduced_prim = ice->state.reduced_prim_mode; - if (reduced_prim == PIPE_PRIM_LINES) - line_aa = BRW_WM_AA_ALWAYS; - else if (reduced_prim == PIPE_PRIM_TRIANGLES) { -@@ -4939,7 +4942,7 @@ emit_surface_state(struct crocus_batch *batch, - struct crocus_resource *res, - const struct isl_surf *in_surf, - bool adjust_surf, -- struct isl_view *view, -+ struct isl_view *in_view, - bool writeable, - enum isl_aux_usage aux_usage, - bool blend_enable, -@@ -4956,23 +4959,24 @@ emit_surface_state(struct crocus_batch *batch, - reloc |= RELOC_WRITE; - - struct isl_surf surf = *in_surf; -+ struct isl_view view = *in_view; - if (adjust_surf) { -- if (res->base.b.target == PIPE_TEXTURE_3D && view->array_len == 1) { -+ if (res->base.b.target == PIPE_TEXTURE_3D && view.array_len == 1) { - isl_surf_get_image_surf(isl_dev, in_surf, -- view->base_level, 0, -- view->base_array_layer, -+ view.base_level, 0, -+ view.base_array_layer, - &surf, &offset, - &tile_x_sa, &tile_y_sa); -- view->base_array_layer = 0; -- view->base_level = 0; -+ view.base_array_layer = 0; -+ view.base_level = 0; - } else if (res->base.b.target == PIPE_TEXTURE_CUBE && devinfo->ver == 4) { - isl_surf_get_image_surf(isl_dev, in_surf, -- view->base_level, view->base_array_layer, -+ view.base_level, view.base_array_layer, - 0, - &surf, &offset, - &tile_x_sa, &tile_y_sa); -- view->base_array_layer = 0; -- view->base_level = 0; -+ view.base_array_layer = 0; -+ view.base_level = 0; - } else if (res->base.b.target == PIPE_TEXTURE_1D_ARRAY) - surf.dim = ISL_SURF_DIM_2D; - } -@@ -4991,7 +4995,7 @@ emit_surface_state(struct crocus_batch *batch, - - isl_surf_fill_state(isl_dev, surf_state, - .surf = &surf, -- .view = view, -+ .view = &view, - .address = crocus_state_reloc(batch, - addr_offset + isl_dev->ss.addr_offset, - res->bo, offset, reloc), -@@ -7016,11 +7020,15 @@ crocus_upload_dirty_render_state(struct crocus_context *ice, - sf.DestinationOriginHorizontalBias = 0.5; - sf.DestinationOriginVerticalBias = 0.5; - -+ sf.LineEndCapAntialiasingRegionWidth = -+ cso_state->line_smooth ? _10pixels : _05pixels; - sf.LastPixelEnable = cso_state->line_last_pixel; -+ sf.AntialiasingEnable = cso_state->line_smooth; -+ - sf.LineWidth = get_line_width(cso_state); - sf.PointWidth = cso_state->point_size; - sf.PointWidthSource = cso_state->point_size_per_vertex ? Vertex : State; --#if GFX_VERx10 == 45 || GFX_VER >= 5 -+#if GFX_VERx10 >= 45 - sf.AALineDistanceMode = AALINEDISTANCE_TRUE; - #endif - sf.ViewportTransformEnable = true; -@@ -9230,6 +9238,7 @@ genX(crocus_init_state)(struct crocus_context *ice) - ice->state.sample_mask = 0xff; - ice->state.num_viewports = 1; - ice->state.prim_mode = PIPE_PRIM_MAX; -+ ice->state.reduced_prim_mode = PIPE_PRIM_MAX; - ice->state.genx = calloc(1, sizeof(struct crocus_genx_state)); - ice->draw.derived_params.drawid = -1; - -diff --git a/src/gallium/drivers/etnaviv/etnaviv_resource.c b/src/gallium/drivers/etnaviv/etnaviv_resource.c -index aa47be8ed07..6f77b829151 100644 ---- a/src/gallium/drivers/etnaviv/etnaviv_resource.c -+++ b/src/gallium/drivers/etnaviv/etnaviv_resource.c -@@ -639,8 +639,7 @@ etna_resource_get_param(struct pipe_screen *pscreen, - enum pipe_resource_param param, - unsigned usage, uint64_t *value) - { -- switch (param) { -- case PIPE_RESOURCE_PARAM_NPLANES: { -+ if (param == PIPE_RESOURCE_PARAM_NPLANES) { - unsigned count = 0; - - for (struct pipe_resource *cur = prsc; cur; cur = cur->next) -@@ -648,6 +647,25 @@ etna_resource_get_param(struct pipe_screen *pscreen, - *value = count; - return true; - } -+ -+ struct pipe_resource *cur = prsc; -+ for (int i = 0; i < plane; i++) { -+ cur = cur->next; -+ if (!cur) -+ return false; -+ } -+ struct etna_resource *rsc = etna_resource(cur); -+ -+ switch (param) { -+ case PIPE_RESOURCE_PARAM_STRIDE: -+ *value = rsc->levels[level].stride; -+ return true; -+ case PIPE_RESOURCE_PARAM_OFFSET: -+ *value = rsc->levels[level].offset; -+ return true; -+ case PIPE_RESOURCE_PARAM_MODIFIER: -+ *value = layout_to_modifier(rsc->layout); -+ return true; - default: - return false; - } -diff --git a/src/gallium/drivers/freedreno/freedreno_util.h b/src/gallium/drivers/freedreno/freedreno_util.h -index f8cf9b6cb19..c989262cc74 100644 ---- a/src/gallium/drivers/freedreno/freedreno_util.h -+++ b/src/gallium/drivers/freedreno/freedreno_util.h -@@ -398,7 +398,7 @@ emit_marker(struct fd_ringbuffer *ring, int scratch_idx) - if (reg == HW_QUERY_BASE_REG) - return; - if (__EMIT_MARKER) { -- OUT_WFI5(ring); -+ OUT_WFI(ring); - OUT_PKT0(ring, reg, 1); - OUT_RING(ring, p_atomic_inc_return(&marker_cnt)); - } -diff --git a/src/gallium/drivers/freedreno/meson.build b/src/gallium/drivers/freedreno/meson.build -index 9f1747646e1..eac04aa4d97 100644 ---- a/src/gallium/drivers/freedreno/meson.build -+++ b/src/gallium/drivers/freedreno/meson.build -@@ -273,6 +273,7 @@ libfreedreno = static_library( - cpp_args : [freedreno_cpp_args], - gnu_symbol_visibility : 'hidden', - dependencies : libfreedreno_dependencies, -+ override_options : ['cpp_std=c++17'], - ) - - driver_freedreno = declare_dependency( -diff --git a/src/gallium/drivers/iris/iris_clear.c b/src/gallium/drivers/iris/iris_clear.c -index cc619b46a0c..a59ba735cbc 100644 ---- a/src/gallium/drivers/iris/iris_clear.c -+++ b/src/gallium/drivers/iris/iris_clear.c -@@ -185,8 +185,8 @@ convert_clear_color(enum pipe_format format, - unsigned bits = util_format_get_component_bits( - format, UTIL_FORMAT_COLORSPACE_RGB, i); - if (bits > 0 && bits < 32) { -- int32_t max = (1 << (bits - 1)) - 1; -- int32_t min = -(1 << (bits - 1)); -+ int32_t max = u_intN_max(bits); -+ int32_t min = u_intN_min(bits); - override_color.i32[i] = CLAMP(override_color.i32[i], min, max); - } - } -diff --git a/src/gallium/drivers/iris/iris_context.h b/src/gallium/drivers/iris/iris_context.h -index 0f78e7d82fa..3b54a6b680d 100644 ---- a/src/gallium/drivers/iris/iris_context.h -+++ b/src/gallium/drivers/iris/iris_context.h -@@ -202,10 +202,15 @@ struct iris_base_prog_key { - unsigned program_string_id; - }; - -+/** -+ * Note, we need to take care to have padding explicitly declared -+ * for key since we will directly memcmp the whole struct. -+ */ - struct iris_vue_prog_key { - struct iris_base_prog_key base; - - unsigned nr_userclip_plane_consts:4; -+ unsigned padding:28; - }; - - struct iris_vs_prog_key { -diff --git a/src/gallium/drivers/iris/iris_state.c b/src/gallium/drivers/iris/iris_state.c -index 2f0831a1078..de04c11f093 100644 ---- a/src/gallium/drivers/iris/iris_state.c -+++ b/src/gallium/drivers/iris/iris_state.c -@@ -2415,10 +2415,15 @@ iris_create_sampler_view(struct pipe_context *ctx, - if (tmpl->target != PIPE_BUFFER) { - isv->view.base_level = tmpl->u.tex.first_level; - isv->view.levels = tmpl->u.tex.last_level - tmpl->u.tex.first_level + 1; -- // XXX: do I need to port f9fd0cf4790cb2a530e75d1a2206dbb9d8af7cb2? -- isv->view.base_array_layer = tmpl->u.tex.first_layer; -- isv->view.array_len = -- tmpl->u.tex.last_layer - tmpl->u.tex.first_layer + 1; -+ -+ if (tmpl->target == PIPE_TEXTURE_3D) { -+ isv->view.base_array_layer = 0; -+ isv->view.array_len = 1; -+ } else { -+ isv->view.base_array_layer = tmpl->u.tex.first_layer; -+ isv->view.array_len = -+ tmpl->u.tex.last_layer - tmpl->u.tex.first_layer + 1; -+ } - - if (iris_resource_unfinished_aux_import(isv->res)) - iris_resource_finish_aux_import(&screen->base, isv->res); -diff --git a/src/gallium/drivers/llvmpipe/lp_cs_tpool.c b/src/gallium/drivers/llvmpipe/lp_cs_tpool.c -index ea284468512..dd28dbafcbd 100644 ---- a/src/gallium/drivers/llvmpipe/lp_cs_tpool.c -+++ b/src/gallium/drivers/llvmpipe/lp_cs_tpool.c -@@ -121,6 +121,7 @@ lp_cs_tpool_queue_task(struct lp_cs_tpool *pool, - for (unsigned t = 0; t < num_iters; t++) { - work(data, t, &lmem); - } -+ FREE(lmem.local_mem_ptr); - return NULL; - } - task = CALLOC_STRUCT(lp_cs_tpool_task); -diff --git a/src/gallium/drivers/llvmpipe/lp_state_cs.c b/src/gallium/drivers/llvmpipe/lp_state_cs.c -index 93e5d0cca6f..9a3291ded46 100644 ---- a/src/gallium/drivers/llvmpipe/lp_state_cs.c -+++ b/src/gallium/drivers/llvmpipe/lp_state_cs.c -@@ -1439,6 +1439,9 @@ lp_csctx_destroy(struct lp_cs_context *csctx) - for (i = 0; i < ARRAY_SIZE(csctx->ssbos); i++) { - pipe_resource_reference(&csctx->ssbos[i].current.buffer, NULL); - } -+ for (i = 0; i < ARRAY_SIZE(csctx->images); i++) { -+ pipe_resource_reference(&csctx->images[i].current.resource, NULL); -+ } - FREE(csctx); - } - -diff --git a/src/gallium/drivers/panfrost/ci/deqp-panfrost-g72-fails.txt b/src/gallium/drivers/panfrost/ci/deqp-panfrost-g72-fails.txt -index 03a9ff5e49a..aee45fb997a 100644 ---- a/src/gallium/drivers/panfrost/ci/deqp-panfrost-g72-fails.txt -+++ b/src/gallium/drivers/panfrost/ci/deqp-panfrost-g72-fails.txt -@@ -1,91 +1,3 @@ --dEQP-GLES31.functional.blend_equation_advanced.barrier.colorburn,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.colordodge,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.darken,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.difference,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.exclusion,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.hardlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.hsl_color,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.hsl_hue,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.hsl_luminosity,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.hsl_saturation,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.lighten,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.multiply,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.overlay,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.screen,Fail --dEQP-GLES31.functional.blend_equation_advanced.barrier.softlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.colorburn,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.colordodge,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.darken,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.difference,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.exclusion,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.hardlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.hsl_color,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.hsl_hue,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.hsl_luminosity,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.hsl_saturation,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.lighten,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.multiply,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.overlay,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.screen,Fail --dEQP-GLES31.functional.blend_equation_advanced.basic.softlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.colorburn,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.colordodge,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.darken,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.difference,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.exclusion,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.hardlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.hsl_color,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.hsl_hue,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.hsl_luminosity,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.hsl_saturation,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.lighten,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.colorburn,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.colordodge,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.darken,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.difference,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.exclusion,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.hardlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.hsl_color,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.hsl_hue,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.hsl_luminosity,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.hsl_saturation,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.lighten,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.multiply,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.overlay,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.screen,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent_msaa.softlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.multiply,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.overlay,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.screen,Fail --dEQP-GLES31.functional.blend_equation_advanced.coherent.softlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.colorburn,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.colordodge,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.darken,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.difference,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.exclusion,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.hsl_hue,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.hsl_luminosity,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.hsl_saturation,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.lighten,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.multiply,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.overlay,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.screen,Fail --dEQP-GLES31.functional.blend_equation_advanced.msaa.softlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.colorburn,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.colordodge,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.darken,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.difference,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.exclusion,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.hardlight,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.hsl_color,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.hsl_hue,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.hsl_luminosity,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.hsl_saturation,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.lighten,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.multiply,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.overlay,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.screen,Fail --dEQP-GLES31.functional.blend_equation_advanced.srgb.softlight,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_advanced_blend_eq_buffer_advanced_blend_eq,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_blend_eq_buffer_advanced_blend_eq,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_separate_blend_eq_buffer_advanced_blend_eq,Fail -@@ -113,47 +25,6 @@ dEQP-GLES31.functional.separate_shader.random.79,Fail - dEQP-GLES31.functional.separate_shader.random.80,Fail - dEQP-GLES31.functional.separate_shader.random.82,Fail - dEQP-GLES31.functional.separate_shader.random.89,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.fragment_discard,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.framebuffer_texture_layer,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.framebuffer_texture_level,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.last_frag_data,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.multiple_assignment,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.texel_fetch,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r11f_g11f_b10f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r16f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r16i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r16ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r32f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r32i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r32ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r8,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.r8ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg16f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg16i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg16ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg32f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg32i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg32ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg8,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg8i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rg8ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb10_a2,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb10_a2ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb16f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb565,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb5_a1,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb8,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba16f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba16i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba16ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba32f,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba32i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba32ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba4,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba8,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba8i,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgba8ui,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.srgb8_alpha8,Fail - dEQP-GLES31.functional.ssbo.layout.random.all_shared_buffer.21,Crash - dEQP-GLES31.functional.ssbo.layout.random.all_shared_buffer.36,Crash - dEQP-GLES31.functional.texture.border_clamp.depth_compare_mode.depth24_stencil8.gather_size_npot,Fail -diff --git a/src/gallium/drivers/panfrost/ci/deqp-panfrost-t860-fails.txt b/src/gallium/drivers/panfrost/ci/deqp-panfrost-t860-fails.txt -index 9219c97eb5b..4fe0791c622 100644 ---- a/src/gallium/drivers/panfrost/ci/deqp-panfrost-t860-fails.txt -+++ b/src/gallium/drivers/panfrost/ci/deqp-panfrost-t860-fails.txt -@@ -12,31 +12,12 @@ dEQP-GLES3.functional.fbo.blit.rect.nearest_consistency_min_reverse_src_dst_x,Fa - dEQP-GLES3.functional.fbo.blit.rect.nearest_consistency_min_reverse_src_dst_y,Fail - dEQP-GLES3.functional.fbo.blit.rect.nearest_consistency_min_reverse_src_x,Fail - dEQP-GLES3.functional.fbo.blit.rect.nearest_consistency_min_reverse_src_y,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_blend_eq_buffer_blend_eq,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_blend_eq_buffer_separate_blend_eq,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_blend_func_buffer_blend_func,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_blend_func_buffer_separate_blend_func,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_color_mask_buffer_color_mask,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_disable_buffer_disable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_disable_buffer_enable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_enable_buffer_disable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_enable_buffer_enable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_separate_blend_eq_buffer_blend_eq,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_separate_blend_eq_buffer_separate_blend_eq,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_separate_blend_func_buffer_blend_func,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_common.common_separate_blend_func_buffer_separate_blend_func,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_blend_eq_buffer_blend_eq,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_blend_eq_buffer_separate_blend_eq,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_blend_func_buffer_blend_func,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_blend_func_buffer_separate_blend_func,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_disable_buffer_disable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_disable_buffer_enable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_enable_buffer_disable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_enable_buffer_enable,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_separate_blend_eq_buffer_blend_eq,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_separate_blend_eq_buffer_separate_blend_eq,Fail - dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_separate_blend_func_buffer_blend_func,Fail --dEQP-GLES31.functional.draw_buffers_indexed.overwrite_indexed.common_separate_blend_func_buffer_separate_blend_func,Fail - dEQP-GLES31.functional.draw_buffers_indexed.random.max_implementation_draw_buffers.0,Fail - dEQP-GLES31.functional.draw_buffers_indexed.random.max_implementation_draw_buffers.1,Fail - dEQP-GLES31.functional.draw_buffers_indexed.random.max_implementation_draw_buffers.10,Fail -@@ -81,10 +62,8 @@ dEQP-GLES31.functional.shaders.builtin_functions.integer.findmsb.uvec2_lowp_comp - dEQP-GLES31.functional.shaders.builtin_functions.integer.findmsb.uvec3_lowp_compute,Fail - dEQP-GLES31.functional.shaders.builtin_functions.integer.imulextended.ivec3_highp_fragment,Fail - dEQP-GLES31.functional.shaders.builtin_functions.integer.umulextended.uvec3_highp_fragment,Fail --dEQP-GLES31.functional.shaders.framebuffer_fetch.framebuffer_format.rgb10_a2,Fail - dEQP-GLES31.functional.shaders.opaque_type_indexing.ssbo.const_expression_vertex,Fail - dEQP-GLES31.functional.shaders.opaque_type_indexing.ssbo.const_literal_vertex,Fail --dEQP-GLES31.functional.shaders.opaque_type_indexing.ubo.const_expression_fragment,Fail - dEQP-GLES31.functional.shaders.opaque_type_indexing.ubo.const_expression_vertex,Fail - dEQP-GLES31.functional.shaders.opaque_type_indexing.ubo.const_literal_fragment,Fail - dEQP-GLES31.functional.shaders.opaque_type_indexing.ubo.const_literal_vertex,Fail -diff --git a/src/gallium/drivers/panfrost/ci/gitlab-ci.yml b/src/gallium/drivers/panfrost/ci/gitlab-ci.yml -index 01f91016200..9af4cdd3b0b 100644 ---- a/src/gallium/drivers/panfrost/ci/gitlab-ci.yml -+++ b/src/gallium/drivers/panfrost/ci/gitlab-ci.yml -@@ -154,7 +154,7 @@ panfrost-g52-gles31:arm64: - - .lava-meson-g12b-a311d-khadas-vim3 - variables: - DEQP_VER: gles31 -- PAN_MESA_DEBUG: "deqp,sync" -+ PAN_MESA_DEBUG: "deqp,sync,indirect" - DEQP_PARALLEL: 6 - DEQP_EXPECTED_RENDERER: G52 - -diff --git a/src/gallium/drivers/panfrost/pan_cmdstream.c b/src/gallium/drivers/panfrost/pan_cmdstream.c -index d31a7466a9d..4d3a442eac9 100644 ---- a/src/gallium/drivers/panfrost/pan_cmdstream.c -+++ b/src/gallium/drivers/panfrost/pan_cmdstream.c -@@ -695,7 +695,7 @@ panfrost_emit_frag_shader_meta(struct panfrost_batch *batch) - PAN_DESC_ARRAY(rt_count, BLEND)); - } - -- mali_ptr blend_shaders[PIPE_MAX_COLOR_BUFS]; -+ mali_ptr blend_shaders[PIPE_MAX_COLOR_BUFS] = { 0 }; - unsigned shader_offset = 0; - struct panfrost_bo *shader_bo = NULL; - -@@ -3078,8 +3078,8 @@ panfrost_draw_vbo(struct pipe_context *pipe, - if (!panfrost_render_condition_check(ctx)) - return; - -- /* Emulate indirect draws when debugging */ -- if (dev->debug & PAN_DBG_NOINDIRECT && indirect && indirect->buffer) { -+ /* Emulate indirect draws unless we're using the experimental path */ -+ if (!(dev->debug & PAN_DBG_INDIRECT) && indirect && indirect->buffer) { - assert(num_draws == 1); - util_draw_indirect(pipe, info, indirect); - return; -diff --git a/src/gallium/drivers/panfrost/pan_context.c b/src/gallium/drivers/panfrost/pan_context.c -index f8a4326acc0..b617c41241b 100644 ---- a/src/gallium/drivers/panfrost/pan_context.c -+++ b/src/gallium/drivers/panfrost/pan_context.c -@@ -297,6 +297,8 @@ panfrost_create_shader_state( - struct panfrost_device *dev = pan_device(pctx->screen); - so->base = *cso; - -+ simple_mtx_init(&so->lock, mtx_plain); -+ - /* Token deep copy to prevent memory corruption */ - - if (cso->type == PIPE_SHADER_IR_TGSI) -@@ -337,6 +339,8 @@ panfrost_delete_shader_state( - panfrost_bo_unreference(shader_state->linkage.bo); - } - -+ simple_mtx_destroy(&cso->lock); -+ - free(cso->variants); - free(so); - } -@@ -364,8 +368,6 @@ panfrost_variant_matches( - struct panfrost_shader_state *variant, - enum pipe_shader_type type) - { -- struct panfrost_device *dev = pan_device(ctx->base.screen); -- - if (variant->info.stage == MESA_SHADER_FRAGMENT && - variant->info.fs.outputs_read) { - struct pipe_framebuffer_state *fb = &ctx->pipe_framebuffer; -@@ -377,10 +379,7 @@ panfrost_variant_matches( - if ((fb->nr_cbufs > i) && fb->cbufs[i]) - fmt = fb->cbufs[i]->format; - -- const struct util_format_description *desc = -- util_format_description(fmt); -- -- if (pan_format_class_load(desc, dev->quirks) == PAN_FORMAT_NATIVE) -+ if (panfrost_blendable_formats_v6[fmt].internal) - fmt = PIPE_FORMAT_NONE; - - if (variant->rt_formats[i] != fmt) -@@ -442,7 +441,6 @@ panfrost_bind_shader_state( - enum pipe_shader_type type) - { - struct panfrost_context *ctx = pan_context(pctx); -- struct panfrost_device *dev = pan_device(ctx->base.screen); - ctx->shader[type] = hwcso; - - ctx->dirty |= PAN_DIRTY_TLS_SIZE; -@@ -455,6 +453,8 @@ panfrost_bind_shader_state( - signed variant = -1; - struct panfrost_shader_variants *variants = (struct panfrost_shader_variants *) hwcso; - -+ simple_mtx_lock(&variants->lock); -+ - for (unsigned i = 0; i < variants->variant_count; ++i) { - if (panfrost_variant_matches(ctx, &variants->variants[i], type)) { - variant = i; -@@ -498,10 +498,7 @@ panfrost_bind_shader_state( - if ((fb->nr_cbufs > i) && fb->cbufs[i]) - fmt = fb->cbufs[i]->format; - -- const struct util_format_description *desc = -- util_format_description(fmt); -- -- if (pan_format_class_load(desc, dev->quirks) == PAN_FORMAT_NATIVE) -+ if (panfrost_blendable_formats_v6[fmt].internal) - fmt = PIPE_FORMAT_NONE; - - v->rt_formats[i] = fmt; -@@ -535,6 +532,11 @@ panfrost_bind_shader_state( - update_so_info(&shader_state->stream_output, - shader_state->info.outputs_written); - } -+ -+ /* TODO: it would be more efficient to release the lock before -+ * compiling instead of after, but that can race if thread A compiles a -+ * variant while thread B searches for that same variant */ -+ simple_mtx_unlock(&variants->lock); - } - - static void * -@@ -792,6 +794,8 @@ panfrost_destroy(struct pipe_context *pipe) - { - struct panfrost_context *panfrost = pan_context(pipe); - -+ _mesa_hash_table_destroy(panfrost->writers, NULL); -+ - if (panfrost->blitter) - util_blitter_destroy(panfrost->blitter); - -@@ -1124,6 +1128,9 @@ panfrost_create_context(struct pipe_screen *screen, void *priv, unsigned flags) - - ctx->blitter = util_blitter_create(gallium); - -+ ctx->writers = _mesa_hash_table_create(gallium, _mesa_hash_pointer, -+ _mesa_key_pointer_equal); -+ - assert(ctx->blitter); - - /* Prepare for render! */ -diff --git a/src/gallium/drivers/panfrost/pan_context.h b/src/gallium/drivers/panfrost/pan_context.h -index b2ad9af36ba..6febeb8d4cf 100644 ---- a/src/gallium/drivers/panfrost/pan_context.h -+++ b/src/gallium/drivers/panfrost/pan_context.h -@@ -44,6 +44,7 @@ - #include "pipe/p_state.h" - #include "util/u_blitter.h" - #include "util/hash_table.h" -+#include "util/simple_mtx.h" - - #include "midgard/midgard_compile.h" - #include "compiler/shader_enums.h" -@@ -140,8 +141,14 @@ struct panfrost_context { - struct { - uint64_t seqnum; - struct panfrost_batch slots[PAN_MAX_BATCHES]; -+ -+ /** Set of active batches for faster traversal */ -+ BITSET_DECLARE(active, PAN_MAX_BATCHES); - } batches; - -+ /* Map from resources to panfrost_batches */ -+ struct hash_table *writers; -+ - /* Bound job batch */ - struct panfrost_batch *batch; - -@@ -290,6 +297,9 @@ struct panfrost_shader_variants { - struct pipe_compute_state cbase; - }; - -+ /** Lock for the variants array */ -+ simple_mtx_t lock; -+ - struct panfrost_shader_state *variants; - unsigned variant_space; - -diff --git a/src/gallium/drivers/panfrost/pan_job.c b/src/gallium/drivers/panfrost/pan_job.c -index 3dadf45e72b..d84d604361e 100644 ---- a/src/gallium/drivers/panfrost/pan_job.c -+++ b/src/gallium/drivers/panfrost/pan_job.c -@@ -40,6 +40,9 @@ - #include "decode.h" - #include "panfrost-quirks.h" - -+#define foreach_batch(ctx, idx) \ -+ BITSET_FOREACH_SET(idx, ctx->batches.active, PAN_MAX_BATCHES) -+ - static unsigned - panfrost_batch_idx(struct panfrost_batch *batch) - { -@@ -65,7 +68,8 @@ panfrost_batch_init(struct panfrost_context *ctx, - batch->maxx = batch->maxy = 0; - - util_copy_framebuffer_state(&batch->key, key); -- util_dynarray_init(&batch->resources, NULL); -+ batch->resources =_mesa_set_create(NULL, _mesa_hash_pointer, -+ _mesa_key_pointer_equal); - - /* Preallocate the main pool, since every batch has at least one job - * structure so it will be used */ -@@ -125,16 +129,20 @@ panfrost_batch_cleanup(struct panfrost_batch *batch) - panfrost_bo_unreference(bo); - } - -- util_dynarray_foreach(&batch->resources, struct panfrost_resource *, rsrc) { -- BITSET_CLEAR((*rsrc)->track.users, batch_idx); -+ set_foreach_remove(batch->resources, entry) { -+ struct panfrost_resource *rsrc = (void *) entry->key; -+ -+ if (_mesa_hash_table_search(ctx->writers, rsrc)) { -+ _mesa_hash_table_remove_key(ctx->writers, rsrc); -+ rsrc->track.nr_writers--; -+ } - -- if ((*rsrc)->track.writer == batch) -- (*rsrc)->track.writer = NULL; -+ rsrc->track.nr_users--; - -- pipe_resource_reference((struct pipe_resource **) rsrc, NULL); -+ pipe_resource_reference((struct pipe_resource **) &rsrc, NULL); - } - -- util_dynarray_fini(&batch->resources); -+ _mesa_set_destroy(batch->resources, NULL); - panfrost_pool_cleanup(&batch->pool); - panfrost_pool_cleanup(&batch->invisible_pool); - -@@ -143,6 +151,7 @@ panfrost_batch_cleanup(struct panfrost_batch *batch) - util_sparse_array_finish(&batch->bos); - - memset(batch, 0, sizeof(*batch)); -+ BITSET_CLEAR(ctx->batches.active, batch_idx); - } - - static void -@@ -177,6 +186,9 @@ panfrost_get_batch(struct panfrost_context *ctx, - - panfrost_batch_init(ctx, key, batch); - -+ unsigned batch_idx = panfrost_batch_idx(batch); -+ BITSET_SET(ctx->batches.active, batch_idx); -+ - return batch; - } - -@@ -262,33 +274,40 @@ panfrost_batch_update_access(struct panfrost_batch *batch, - { - struct panfrost_context *ctx = batch->ctx; - uint32_t batch_idx = panfrost_batch_idx(batch); -- struct panfrost_batch *writer = rsrc->track.writer; -+ struct hash_entry *entry = _mesa_hash_table_search(ctx->writers, rsrc); -+ struct panfrost_batch *writer = entry ? entry->data : NULL; -+ bool found = false; - -- if (unlikely(!BITSET_TEST(rsrc->track.users, batch_idx))) { -- BITSET_SET(rsrc->track.users, batch_idx); -+ _mesa_set_search_or_add(batch->resources, rsrc, &found); - -- /* Reference the resource on the batch */ -- struct pipe_resource **dst = util_dynarray_grow(&batch->resources, -- struct pipe_resource *, 1); -+ if (!found) { -+ /* Cache number of batches accessing a resource */ -+ rsrc->track.nr_users++; - -- *dst = NULL; -- pipe_resource_reference(dst, &rsrc->base); -+ /* Reference the resource on the batch */ -+ pipe_reference(NULL, &rsrc->base.reference); - } - - /* Flush users if required */ - if (writes || ((writer != NULL) && (writer != batch))) { - unsigned i; -- BITSET_FOREACH_SET(i, rsrc->track.users, PAN_MAX_BATCHES) { -+ foreach_batch(ctx, i) { -+ struct panfrost_batch *batch = &ctx->batches.slots[i]; -+ - /* Skip the entry if this our batch. */ - if (i == batch_idx) - continue; - -- panfrost_batch_submit(&ctx->batches.slots[i], 0, 0); -+ /* Submit if it's a user */ -+ if (_mesa_set_search(batch->resources, rsrc)) -+ panfrost_batch_submit(batch, 0, 0); - } - } - -- if (writes) -- rsrc->track.writer = batch; -+ if (writes) { -+ _mesa_hash_table_insert(ctx->writers, rsrc, batch); -+ rsrc->track.nr_writers++; -+ } - } - - static void -@@ -919,9 +938,10 @@ void - panfrost_flush_writer(struct panfrost_context *ctx, - struct panfrost_resource *rsrc) - { -- if (rsrc->track.writer) { -- panfrost_batch_submit(rsrc->track.writer, ctx->syncobj, ctx->syncobj); -- rsrc->track.writer = NULL; -+ struct hash_entry *entry = _mesa_hash_table_search(ctx->writers, rsrc); -+ -+ if (entry) { -+ panfrost_batch_submit(entry->data, ctx->syncobj, ctx->syncobj); - } - } - -@@ -930,13 +950,14 @@ panfrost_flush_batches_accessing_rsrc(struct panfrost_context *ctx, - struct panfrost_resource *rsrc) - { - unsigned i; -- BITSET_FOREACH_SET(i, rsrc->track.users, PAN_MAX_BATCHES) { -- panfrost_batch_submit(&ctx->batches.slots[i], -- ctx->syncobj, ctx->syncobj); -- } -+ foreach_batch(ctx, i) { -+ struct panfrost_batch *batch = &ctx->batches.slots[i]; - -- assert(!BITSET_COUNT(rsrc->track.users)); -- rsrc->track.writer = NULL; -+ if (!_mesa_set_search(batch->resources, rsrc)) -+ continue; -+ -+ panfrost_batch_submit(batch, ctx->syncobj, ctx->syncobj); -+ } - } - - void -@@ -969,7 +990,7 @@ panfrost_batch_clear(struct panfrost_batch *batch, - continue; - - enum pipe_format format = ctx->pipe_framebuffer.cbufs[i]->format; -- pan_pack_color(batch->clear_color[i], color, format); -+ pan_pack_color(batch->clear_color[i], color, format, false); - } - } - -diff --git a/src/gallium/drivers/panfrost/pan_job.h b/src/gallium/drivers/panfrost/pan_job.h -index 2e5af79d73c..6f05af31e6b 100644 ---- a/src/gallium/drivers/panfrost/pan_job.h -+++ b/src/gallium/drivers/panfrost/pan_job.h -@@ -130,8 +130,8 @@ struct panfrost_batch { - mali_ptr uniform_buffers[PIPE_SHADER_TYPES]; - mali_ptr push_uniforms[PIPE_SHADER_TYPES]; - -- /* Referenced resources for cleanup */ -- struct util_dynarray resources; -+ /* Referenced resources */ -+ struct set *resources; - }; - - /* Functions for managing the above */ -diff --git a/src/gallium/drivers/panfrost/pan_resource.c b/src/gallium/drivers/panfrost/pan_resource.c -index b56e5c428b1..131699ac8d9 100644 ---- a/src/gallium/drivers/panfrost/pan_resource.c -+++ b/src/gallium/drivers/panfrost/pan_resource.c -@@ -67,7 +67,7 @@ panfrost_resource_from_handle(struct pipe_screen *pscreen, - - assert(whandle->type == WINSYS_HANDLE_TYPE_FD); - -- rsc = rzalloc(pscreen, struct panfrost_resource); -+ rsc = CALLOC_STRUCT(panfrost_resource); - if (!rsc) - return NULL; - -@@ -98,7 +98,7 @@ panfrost_resource_from_handle(struct pipe_screen *pscreen, - crc_mode, &explicit_layout); - - if (!valid) { -- ralloc_free(rsc); -+ FREE(rsc); - return NULL; - } - -@@ -107,7 +107,7 @@ panfrost_resource_from_handle(struct pipe_screen *pscreen, - * memory space to mmap it etc. - */ - if (!rsc->image.data.bo) { -- ralloc_free(rsc); -+ FREE(rsc); - return NULL; - } - if (rsc->image.layout.crc_mode == PAN_IMAGE_CRC_OOB) -@@ -183,6 +183,30 @@ panfrost_resource_get_handle(struct pipe_screen *pscreen, - return false; - } - -+static bool -+panfrost_resource_get_param(struct pipe_screen *pscreen, -+ struct pipe_context *pctx, struct pipe_resource *prsc, -+ unsigned plane, unsigned layer, unsigned level, -+ enum pipe_resource_param param, -+ unsigned usage, uint64_t *value) -+{ -+ struct panfrost_resource *rsrc = (struct panfrost_resource *) prsc; -+ -+ switch (param) { -+ case PIPE_RESOURCE_PARAM_STRIDE: -+ *value = rsrc->image.layout.slices[level].line_stride; -+ return true; -+ case PIPE_RESOURCE_PARAM_OFFSET: -+ *value = rsrc->image.layout.slices[level].offset; -+ return true; -+ case PIPE_RESOURCE_PARAM_MODIFIER: -+ *value = rsrc->image.layout.modifier; -+ return true; -+ default: -+ return false; -+ } -+} -+ - static void - panfrost_flush_resource(struct pipe_context *pctx, struct pipe_resource *prsc) - { -@@ -526,7 +550,7 @@ panfrost_resource_set_damage_region(struct pipe_screen *screen, - pres->damage.tile_map.stride * - DIV_ROUND_UP(res->height0, 32); - pres->damage.tile_map.data = -- ralloc_size(pres, pres->damage.tile_map.size); -+ malloc(pres->damage.tile_map.size); - } - - memset(pres->damage.tile_map.data, 0, pres->damage.tile_map.size); -@@ -610,7 +634,7 @@ panfrost_resource_create_with_modifier(struct pipe_screen *screen, - (PIPE_BIND_DISPLAY_TARGET | PIPE_BIND_SCANOUT | PIPE_BIND_SHARED))) - return panfrost_create_scanout_res(screen, template, modifier); - -- struct panfrost_resource *so = rzalloc(screen, struct panfrost_resource); -+ struct panfrost_resource *so = CALLOC_STRUCT(panfrost_resource); - so->base = *template; - so->base.screen = screen; - -@@ -648,7 +672,7 @@ panfrost_resource_create_with_modifier(struct pipe_screen *screen, - panfrost_resource_set_damage_region(screen, &so->base, 0, NULL); - - if (template->bind & PIPE_BIND_INDEX_BUFFER) -- so->index_cache = rzalloc(so, struct panfrost_minmax_cache); -+ so->index_cache = CALLOC_STRUCT(panfrost_minmax_cache); - - return (struct pipe_resource *)so; - } -@@ -699,8 +723,11 @@ panfrost_resource_destroy(struct pipe_screen *screen, - if (rsrc->image.crc.bo) - panfrost_bo_unreference(rsrc->image.crc.bo); - -+ free(rsrc->index_cache); -+ free(rsrc->damage.tile_map.data); -+ - util_range_destroy(&rsrc->valid_buffer_range); -- ralloc_free(rsrc); -+ free(rsrc); - } - - /* Most of the time we can do CPU-side transfers, but sometimes we need to use -@@ -843,7 +870,7 @@ panfrost_ptr_map(struct pipe_context *pctx, - - bool valid = BITSET_TEST(rsrc->valid.data, level); - -- if ((usage & PIPE_MAP_READ) && (valid || rsrc->track.writer)) { -+ if ((usage & PIPE_MAP_READ) && (valid || rsrc->track.nr_writers > 0)) { - pan_blit_to_staging(pctx, transfer); - panfrost_flush_writer(ctx, staging); - panfrost_bo_wait(staging->image.data.bo, INT64_MAX, false); -@@ -867,7 +894,7 @@ panfrost_ptr_map(struct pipe_context *pctx, - (usage & PIPE_MAP_WRITE) && - !(resource->target == PIPE_BUFFER - && !util_ranges_intersect(&rsrc->valid_buffer_range, box->x, box->x + box->width)) && -- BITSET_COUNT(rsrc->track.users) != 0) { -+ rsrc->track.nr_users > 0) { - - /* When a resource to be modified is already being used by a - * pending batch, it is often faster to copy the whole BO than -@@ -886,7 +913,7 @@ panfrost_ptr_map(struct pipe_context *pctx, - * not ready yet (still accessed by one of the already flushed - * batches), we try to allocate a new one to avoid waiting. - */ -- if (BITSET_COUNT(rsrc->track.users) || -+ if (rsrc->track.nr_users > 0 || - !panfrost_bo_wait(bo, 0, true)) { - /* We want the BO to be MMAPed. */ - uint32_t flags = bo->flags & ~PAN_BO_DELAY_MMAP; -@@ -1314,6 +1341,7 @@ panfrost_resource_screen_init(struct pipe_screen *pscreen) - pscreen->resource_destroy = u_transfer_helper_resource_destroy; - pscreen->resource_from_handle = panfrost_resource_from_handle; - pscreen->resource_get_handle = panfrost_resource_get_handle; -+ pscreen->resource_get_param = panfrost_resource_get_param; - pscreen->transfer_helper = u_transfer_helper_create(&transfer_vtbl, - true, false, - fake_rgtc, true); -diff --git a/src/gallium/drivers/panfrost/pan_resource.h b/src/gallium/drivers/panfrost/pan_resource.h -index a3cd7bf14e1..3102229f9ca 100644 ---- a/src/gallium/drivers/panfrost/pan_resource.h -+++ b/src/gallium/drivers/panfrost/pan_resource.h -@@ -48,8 +48,14 @@ struct panfrost_resource { - } damage; - - struct { -- struct panfrost_batch *writer; -- BITSET_DECLARE(users, PAN_MAX_BATCHES); -+ /** Number of batches accessing this resource. Used to check if -+ * a resource is in use. */ -+ _Atomic unsigned nr_users; -+ -+ /** Number of batches writing this resource. Note that only one -+ * batch per context may write a resource, so this is the -+ * number of contexts that have an active writer. */ -+ _Atomic unsigned nr_writers; - } track; - - struct renderonly_scanout *scanout; -diff --git a/src/gallium/drivers/panfrost/pan_screen.c b/src/gallium/drivers/panfrost/pan_screen.c -index 4b1fabc1809..593f158087d 100644 ---- a/src/gallium/drivers/panfrost/pan_screen.c -+++ b/src/gallium/drivers/panfrost/pan_screen.c -@@ -67,7 +67,9 @@ static const struct debug_named_value panfrost_debug_options[] = { - {"noafbc", PAN_DBG_NO_AFBC, "Disable AFBC support"}, - {"nocrc", PAN_DBG_NO_CRC, "Disable transaction elimination"}, - {"msaa16", PAN_DBG_MSAA16, "Enable MSAA 8x and 16x support"}, -- {"noindirect", PAN_DBG_NOINDIRECT, "Emulate indirect draws on the CPU"}, -+ {"indirect", PAN_DBG_INDIRECT, "Use experimental compute kernel for indirect draws"}, -+ {"linear", PAN_DBG_LINEAR, "Force linear textures"}, -+ {"nocache", PAN_DBG_NO_CACHE, "Disable BO cache"}, - DEBUG_NAMED_VALUE_END - }; - -@@ -216,11 +218,11 @@ panfrost_get_param(struct pipe_screen *screen, enum pipe_cap param) - return 1; - - case PIPE_CAP_MAX_TEXTURE_2D_SIZE: -- return 4096; -+ return 1 << (MAX_MIP_LEVELS - 1); -+ - case PIPE_CAP_MAX_TEXTURE_3D_LEVELS: -- return 13; - case PIPE_CAP_MAX_TEXTURE_CUBE_LEVELS: -- return 13; -+ return MAX_MIP_LEVELS; - - case PIPE_CAP_TGSI_FS_COORD_ORIGIN_LOWER_LEFT: - /* Hardware is natively upper left */ -@@ -699,7 +701,8 @@ panfrost_destroy_screen(struct pipe_screen *pscreen) - panfrost_pool_cleanup(&screen->blitter.desc_pool); - pan_blend_shaders_cleanup(dev); - -- screen->vtbl.screen_destroy(pscreen); -+ if (screen->vtbl.screen_destroy) -+ screen->vtbl.screen_destroy(pscreen); - - if (dev->ro) - dev->ro->destroy(dev->ro); -@@ -836,10 +839,6 @@ panfrost_create_screen(int fd, struct renderonly *ro) - if (dev->arch == 7) - dev->quirks |= MIDGARD_NO_AFBC; - -- /* XXX: Indirect draws on Midgard need debugging, emulate for now */ -- if (dev->arch < 6) -- dev->debug |= PAN_DBG_NOINDIRECT; -- - dev->ro = ro; - - /* Check if we're loading against a supported GPU model. */ -diff --git a/src/gallium/drivers/radeonsi/si_descriptors.c b/src/gallium/drivers/radeonsi/si_descriptors.c -index 2391a9355b5..275a3132f93 100644 ---- a/src/gallium/drivers/radeonsi/si_descriptors.c -+++ b/src/gallium/drivers/radeonsi/si_descriptors.c -@@ -360,7 +360,10 @@ void si_set_mutable_tex_desc_fields(struct si_screen *sscreen, struct si_texture - - state[6] |= S_00A018_META_PIPE_ALIGNED(meta.pipe_aligned) | - S_00A018_META_DATA_ADDRESS_LO(meta_va >> 8) | -- S_00A018_WRITE_COMPRESS_ENABLE((access & SI_IMAGE_ACCESS_DCC_WRITE) != 0); -+ /* DCC image stores require INDEPENDENT_128B_BLOCKS, which is not set -+ * with displayable DCC on Navi12-14 due to DCN limitations. */ -+ S_00A018_WRITE_COMPRESS_ENABLE(tex->surface.u.gfx9.color.dcc.independent_128B_blocks && -+ access & SI_IMAGE_ACCESS_DCC_WRITE); - } - - state[7] = meta_va >> 16; -diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c -index 0cdd3a9f539..b0fcbfe031a 100644 ---- a/src/gallium/drivers/radeonsi/si_shader.c -+++ b/src/gallium/drivers/radeonsi/si_shader.c -@@ -832,7 +832,9 @@ static unsigned si_get_shader_binary_size(struct si_screen *screen, struct si_sh - { - struct ac_rtld_binary rtld; - si_shader_binary_open(screen, shader, &rtld); -- return rtld.exec_size; -+ uint64_t size = rtld.exec_size; -+ ac_rtld_close(&rtld); -+ return size; - } - - static bool si_get_external_symbol(void *data, const char *name, uint64_t *value) -diff --git a/src/gallium/drivers/v3d/v3d_resource.c b/src/gallium/drivers/v3d/v3d_resource.c -index 62da3147620..d48479460db 100644 ---- a/src/gallium/drivers/v3d/v3d_resource.c -+++ b/src/gallium/drivers/v3d/v3d_resource.c -@@ -396,6 +396,21 @@ v3d_resource_destroy(struct pipe_screen *pscreen, - free(rsc); - } - -+static uint64_t -+v3d_resource_modifier(struct v3d_resource *rsc) -+{ -+ if (rsc->tiled) { -+ /* A shared tiled buffer should always be allocated as UIF, -+ * not UBLINEAR or LT. -+ */ -+ assert(rsc->slices[0].tiling == V3D_TILING_UIF_XOR || -+ rsc->slices[0].tiling == V3D_TILING_UIF_NO_XOR); -+ return DRM_FORMAT_MOD_BROADCOM_UIF; -+ } else { -+ return DRM_FORMAT_MOD_LINEAR; -+ } -+} -+ - static bool - v3d_resource_get_handle(struct pipe_screen *pscreen, - struct pipe_context *pctx, -@@ -409,6 +424,7 @@ v3d_resource_get_handle(struct pipe_screen *pscreen, - - whandle->stride = rsc->slices[0].stride; - whandle->offset = 0; -+ whandle->modifier = v3d_resource_modifier(rsc); - - /* If we're passing some reference to our BO out to some other part of - * the system, then we can't do any optimizations about only us being -@@ -416,17 +432,6 @@ v3d_resource_get_handle(struct pipe_screen *pscreen, - */ - bo->private = false; - -- if (rsc->tiled) { -- /* A shared tiled buffer should always be allocated as UIF, -- * not UBLINEAR or LT. -- */ -- assert(rsc->slices[0].tiling == V3D_TILING_UIF_XOR || -- rsc->slices[0].tiling == V3D_TILING_UIF_NO_XOR); -- whandle->modifier = DRM_FORMAT_MOD_BROADCOM_UIF; -- } else { -- whandle->modifier = DRM_FORMAT_MOD_LINEAR; -- } -- - switch (whandle->type) { - case WINSYS_HANDLE_TYPE_SHARED: - return v3d_bo_flink(bo, &whandle->handle); -@@ -448,6 +453,30 @@ v3d_resource_get_handle(struct pipe_screen *pscreen, - return false; - } - -+static bool -+v3d_resource_get_param(struct pipe_screen *pscreen, -+ struct pipe_context *pctx, struct pipe_resource *prsc, -+ unsigned plane, unsigned layer, unsigned level, -+ enum pipe_resource_param param, -+ unsigned usage, uint64_t *value) -+{ -+ struct v3d_resource *rsc = v3d_resource(prsc); -+ -+ switch (param) { -+ case PIPE_RESOURCE_PARAM_STRIDE: -+ *value = rsc->slices[level].stride; -+ return true; -+ case PIPE_RESOURCE_PARAM_OFFSET: -+ *value = 0; -+ return true; -+ case PIPE_RESOURCE_PARAM_MODIFIER: -+ *value = v3d_resource_modifier(rsc); -+ return true; -+ default: -+ return false; -+ } -+} -+ - #define PAGE_UB_ROWS (V3D_UIFCFG_PAGE_SIZE / V3D_UIFBLOCK_ROW_SIZE) - #define PAGE_UB_ROWS_TIMES_1_5 ((PAGE_UB_ROWS * 3) >> 1) - #define PAGE_CACHE_UB_ROWS (V3D_PAGE_CACHE_SIZE / V3D_UIFBLOCK_ROW_SIZE) -@@ -1148,6 +1177,7 @@ v3d_resource_screen_init(struct pipe_screen *pscreen) - pscreen->resource_create = u_transfer_helper_resource_create; - pscreen->resource_from_handle = v3d_resource_from_handle; - pscreen->resource_get_handle = v3d_resource_get_handle; -+ pscreen->resource_get_param = v3d_resource_get_param; - pscreen->resource_destroy = u_transfer_helper_resource_destroy; - pscreen->transfer_helper = u_transfer_helper_create(&transfer_vtbl, - true, false, -diff --git a/src/gallium/drivers/vc4/vc4_resource.c b/src/gallium/drivers/vc4/vc4_resource.c -index af61c4860b8..052588e49f6 100644 ---- a/src/gallium/drivers/vc4/vc4_resource.c -+++ b/src/gallium/drivers/vc4/vc4_resource.c -@@ -283,6 +283,15 @@ vc4_resource_destroy(struct pipe_screen *pscreen, - free(rsc); - } - -+static uint64_t -+vc4_resource_modifier(struct vc4_resource *rsc) -+{ -+ if (rsc->tiled) -+ return DRM_FORMAT_MOD_BROADCOM_VC4_T_TILED; -+ else -+ return DRM_FORMAT_MOD_LINEAR; -+} -+ - static bool - vc4_resource_get_handle(struct pipe_screen *pscreen, - struct pipe_context *pctx, -@@ -295,6 +304,7 @@ vc4_resource_get_handle(struct pipe_screen *pscreen, - - whandle->stride = rsc->slices[0].stride; - whandle->offset = 0; -+ whandle->modifier = vc4_resource_modifier(rsc); - - /* If we're passing some reference to our BO out to some other part of - * the system, then we can't do any optimizations about only us being -@@ -302,11 +312,6 @@ vc4_resource_get_handle(struct pipe_screen *pscreen, - */ - rsc->bo->private = false; - -- if (rsc->tiled) -- whandle->modifier = DRM_FORMAT_MOD_BROADCOM_VC4_T_TILED; -- else -- whandle->modifier = DRM_FORMAT_MOD_LINEAR; -- - switch (whandle->type) { - case WINSYS_HANDLE_TYPE_SHARED: - if (screen->ro) { -@@ -334,6 +339,30 @@ vc4_resource_get_handle(struct pipe_screen *pscreen, - return false; - } - -+static bool -+vc4_resource_get_param(struct pipe_screen *pscreen, -+ struct pipe_context *pctx, struct pipe_resource *prsc, -+ unsigned plane, unsigned layer, unsigned level, -+ enum pipe_resource_param param, -+ unsigned usage, uint64_t *value) -+{ -+ struct vc4_resource *rsc = vc4_resource(prsc); -+ -+ switch (param) { -+ case PIPE_RESOURCE_PARAM_STRIDE: -+ *value = rsc->slices[level].stride; -+ return true; -+ case PIPE_RESOURCE_PARAM_OFFSET: -+ *value = 0; -+ return true; -+ case PIPE_RESOURCE_PARAM_MODIFIER: -+ *value = vc4_resource_modifier(rsc); -+ return true; -+ default: -+ return false; -+ } -+} -+ - static void - vc4_setup_slices(struct vc4_resource *rsc, const char *caller) - { -@@ -1119,6 +1148,7 @@ vc4_resource_screen_init(struct pipe_screen *pscreen) - vc4_resource_create_with_modifiers; - pscreen->resource_from_handle = vc4_resource_from_handle; - pscreen->resource_get_handle = vc4_resource_get_handle; -+ pscreen->resource_get_param = vc4_resource_get_param; - pscreen->resource_destroy = vc4_resource_destroy; - pscreen->transfer_helper = u_transfer_helper_create(&transfer_vtbl, - false, false, -diff --git a/src/gallium/drivers/zink/zink_program.c b/src/gallium/drivers/zink/zink_program.c -index e5e736c72df..81af56b6176 100644 ---- a/src/gallium/drivers/zink/zink_program.c -+++ b/src/gallium/drivers/zink/zink_program.c -@@ -679,8 +679,9 @@ zink_destroy_gfx_program(struct zink_screen *screen, - if (prog->shaders[i]) { - _mesa_set_remove_key(prog->shaders[i]->programs, prog); - prog->shaders[i] = NULL; -- destroy_shader_cache(screen, &prog->base.shader_cache[i]); - } -+ destroy_shader_cache(screen, &prog->base.shader_cache[i]); -+ ralloc_free(prog->nir[i]); - } - - for (int i = 0; i < ARRAY_SIZE(prog->pipelines); ++i) { -@@ -823,7 +824,7 @@ zink_get_gfx_pipeline(struct zink_context *ctx, - memcpy(&pc_entry->state, state, sizeof(*state)); - pc_entry->pipeline = pipeline; - -- entry = _mesa_hash_table_insert_pre_hashed(prog->pipelines[vkmode], state->final_hash, state, pc_entry); -+ entry = _mesa_hash_table_insert_pre_hashed(prog->pipelines[vkmode], state->final_hash, pc_entry, pc_entry); - assert(entry); - } - -@@ -862,7 +863,7 @@ zink_get_compute_pipeline(struct zink_screen *screen, - memcpy(&pc_entry->state, state, sizeof(*state)); - pc_entry->pipeline = pipeline; - -- entry = _mesa_hash_table_insert_pre_hashed(comp->pipelines, state->hash, state, pc_entry); -+ entry = _mesa_hash_table_insert_pre_hashed(comp->pipelines, state->hash, pc_entry, pc_entry); - assert(entry); - } - -diff --git a/src/gallium/frontends/lavapipe/lvp_device.c b/src/gallium/frontends/lavapipe/lvp_device.c -index 58ea8238a5d..a575ab167cb 100644 ---- a/src/gallium/frontends/lavapipe/lvp_device.c -+++ b/src/gallium/frontends/lavapipe/lvp_device.c -@@ -957,7 +957,9 @@ VKAPI_ATTR void VKAPI_CALL lvp_GetPhysicalDeviceProperties2( - case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_LINE_RASTERIZATION_PROPERTIES_EXT: { - VkPhysicalDeviceLineRasterizationPropertiesEXT *properties = - (VkPhysicalDeviceLineRasterizationPropertiesEXT *)ext; -- properties->lineSubPixelPrecisionBits = 4; -+ properties->lineSubPixelPrecisionBits = -+ pdevice->pscreen->get_param(pdevice->pscreen, -+ PIPE_CAP_RASTERIZER_SUBPIXEL_BITS); - break; - } - case VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SUBGROUP_PROPERTIES: { -diff --git a/src/gallium/frontends/osmesa/osmesa.c b/src/gallium/frontends/osmesa/osmesa.c -index 963888971af..91a250d6421 100644 ---- a/src/gallium/frontends/osmesa/osmesa.c -+++ b/src/gallium/frontends/osmesa/osmesa.c -@@ -781,8 +781,11 @@ OSMesaMakeCurrent(OSMesaContext osmesa, void *buffer, GLenum type, - if (osmesa->current_buffer && - (osmesa->current_buffer->visual.color_format != color_format || - osmesa->current_buffer->visual.depth_stencil_format != osmesa->depth_stencil_format || -- osmesa->current_buffer->visual.accum_format != osmesa->accum_format)) { -+ osmesa->current_buffer->visual.accum_format != osmesa->accum_format || -+ osmesa->current_buffer->width != width || -+ osmesa->current_buffer->height != height)) { - osmesa_destroy_buffer(osmesa->current_buffer); -+ osmesa->current_buffer = NULL; - } - - if (!osmesa->current_buffer) { -diff --git a/src/gallium/winsys/radeon/drm/radeon_drm_winsys.c b/src/gallium/winsys/radeon/drm/radeon_drm_winsys.c -index ef9ec590d5d..51ff7bc74be 100644 ---- a/src/gallium/winsys/radeon/drm/radeon_drm_winsys.c -+++ b/src/gallium/winsys/radeon/drm/radeon_drm_winsys.c -@@ -315,6 +315,7 @@ static bool do_winsys_init(struct radeon_drm_winsys *ws) - - /* Check for UVD and VCE */ - ws->info.has_video_hw.uvd_decode = false; -+ ws->info.has_video_hw.vce_encode = false; - ws->info.vce_fw_version = 0x00000000; - if (ws->info.drm_minor >= 32) { - uint32_t value = RADEON_CS_RING_UVD; -@@ -332,6 +333,7 @@ static bool do_winsys_init(struct radeon_drm_winsys *ws) - "VCE FW version", &value)) { - ws->info.vce_fw_version = value; - ws->info.num_rings[RING_VCE] = 1; -+ ws->info.has_video_hw.vce_encode = true; - } - } - } -diff --git a/src/gallium/winsys/svga/drm/vmw_buffer.c b/src/gallium/winsys/svga/drm/vmw_buffer.c -index d537c8be96e..6c235a9a486 100644 ---- a/src/gallium/winsys/svga/drm/vmw_buffer.c -+++ b/src/gallium/winsys/svga/drm/vmw_buffer.c -@@ -357,32 +357,30 @@ vmw_svga_winsys_buffer_map(struct svga_winsys_screen *sws, - enum pipe_map_flags flags) - { - void *map; -+ enum pb_usage_flags pb_flags = 0; - - (void)sws; - if (flags & PIPE_MAP_UNSYNCHRONIZED) - flags &= ~PIPE_MAP_DONTBLOCK; - -- /* NOTE: we're passing PIPE_MAP_x flags instead of -- * PB_USAGE_x flags here. We should probably fix that. -- */ -- STATIC_ASSERT((unsigned) PB_USAGE_CPU_READ == -- (unsigned) PIPE_MAP_READ); -- STATIC_ASSERT((unsigned) PB_USAGE_CPU_WRITE == -- (unsigned) PIPE_MAP_WRITE); -- STATIC_ASSERT((unsigned) PB_USAGE_GPU_READ == -- (unsigned) PIPE_MAP_DIRECTLY); -- STATIC_ASSERT((unsigned) PB_USAGE_DONTBLOCK == -- (unsigned) PIPE_MAP_DONTBLOCK); -- STATIC_ASSERT((unsigned) PB_USAGE_UNSYNCHRONIZED == -- (unsigned) PIPE_MAP_UNSYNCHRONIZED); -- STATIC_ASSERT((unsigned) PB_USAGE_PERSISTENT == -- (unsigned) PIPE_MAP_PERSISTENT); -- -- map = pb_map(vmw_pb_buffer(buf), flags & PB_USAGE_ALL, NULL); -+ if (flags & PIPE_MAP_READ) -+ pb_flags |= PB_USAGE_CPU_READ; -+ if (flags & PIPE_MAP_WRITE) -+ pb_flags |= PB_USAGE_CPU_WRITE; -+ if (flags & PIPE_MAP_DIRECTLY) -+ pb_flags |= PB_USAGE_GPU_READ; -+ if (flags & PIPE_MAP_DONTBLOCK) -+ pb_flags |= PB_USAGE_DONTBLOCK; -+ if (flags & PIPE_MAP_UNSYNCHRONIZED) -+ pb_flags |= PB_USAGE_UNSYNCHRONIZED; -+ if (flags & PIPE_MAP_PERSISTENT) -+ pb_flags |= PB_USAGE_PERSISTENT; -+ -+ map = pb_map(vmw_pb_buffer(buf), pb_flags, NULL); - - #ifdef DEBUG - if (map != NULL) -- debug_flush_map(buf->fbuf, flags); -+ debug_flush_map(buf->fbuf, pb_flags); - #endif - - return map; -diff --git a/src/glx/glxext.c b/src/glx/glxext.c -index 91d021e710e..07bb42e22fe 100644 ---- a/src/glx/glxext.c -+++ b/src/glx/glxext.c -@@ -861,6 +861,7 @@ AllocAndFetchScreenConfigs(Display * dpy, struct glx_display * priv) - _X_HIDDEN struct glx_display * - __glXInitialize(Display * dpy) - { -+ XExtCodes *codes; - struct glx_display *dpyPriv, *d; - #if defined(GLX_DIRECT_RENDERING) && !defined(GLX_USE_APPLEGL) - Bool glx_direct, glx_accel; -@@ -883,8 +884,13 @@ __glXInitialize(Display * dpy) - if (!dpyPriv) - return NULL; - -- dpyPriv->codes = *XInitExtension(dpy, __glXExtensionName); -+ codes = XInitExtension(dpy, __glXExtensionName); -+ if (!codes) { -+ free(dpyPriv); -+ return NULL; -+ } - -+ dpyPriv->codes = *codes; - dpyPriv->dpy = dpy; - - /* This GLX implementation requires X_GLXQueryExtensionsString -diff --git a/src/intel/blorp/blorp_blit.c b/src/intel/blorp/blorp_blit.c -index 17538e4f963..abd99ae271a 100644 ---- a/src/intel/blorp/blorp_blit.c -+++ b/src/intel/blorp/blorp_blit.c -@@ -1649,8 +1649,8 @@ blorp_surf_retile_w_to_y(const struct isl_device *isl_dev, - blorp_surf_fake_interleaved_msaa(isl_dev, info); - } - -- if (isl_dev->info->ver == 6) { -- /* Gfx6 stencil buffers have a very large alignment coming in from the -+ if (isl_dev->info->ver == 6 || isl_dev->info->ver == 7) { -+ /* Gfx6-7 stencil buffers have a very large alignment coming in from the - * miptree. It's out-of-bounds for what the surface state can handle. - * Since we have a single layer and level, it doesn't really matter as - * long as we don't pass a bogus value into isl_surf_fill_state(). -diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp -index 2735a430ec3..8fbb35cf7c0 100644 ---- a/src/intel/compiler/brw_fs.cpp -+++ b/src/intel/compiler/brw_fs.cpp -@@ -4871,6 +4871,20 @@ lower_fb_read_logical_send(const fs_builder &bld, fs_inst *inst) - } - } - -+ /* BSpec 12470 (Gfx8-11), BSpec 47842 (Gfx12+) : -+ * -+ * "Must be zero for Render Target Read message." -+ * -+ * For bits : -+ * - 14 : Stencil Present to Render Target -+ * - 13 : Source Depth Present to Render Target -+ * - 12 : oMask to Render Target -+ * - 11 : Source0 Alpha Present to Render Target -+ */ -+ ubld.group(1, 0).AND(component(header, 0), -+ component(header, 0), -+ brw_imm_ud(~INTEL_MASK(14, 11))); -+ - inst->resize_sources(1); - inst->src[0] = header; - inst->opcode = FS_OPCODE_FB_READ; -diff --git a/src/intel/isl/isl_format.c b/src/intel/isl/isl_format.c -index 60f4409e2b0..bd9c63d140c 100644 ---- a/src/intel/isl/isl_format.c -+++ b/src/intel/isl/isl_format.c -@@ -1188,11 +1188,11 @@ pack_channel(const union isl_color_value *value, unsigned i, - } - break; - case ISL_UINT: -- packed = MIN(value->u32[i], MAX_UINT(layout->bits)); -+ packed = MIN(value->u32[i], u_uintN_max(layout->bits)); - break; - case ISL_SINT: -- packed = MIN(MAX(value->u32[i], MIN_INT(layout->bits)), -- MAX_INT(layout->bits)); -+ packed = MIN(MAX(value->u32[i], u_intN_min(layout->bits)), -+ u_intN_max(layout->bits)); - break; - - default: -@@ -1202,7 +1202,7 @@ pack_channel(const union isl_color_value *value, unsigned i, - unsigned dword = layout->start_bit / 32; - unsigned bit = layout->start_bit % 32; - assert(bit + layout->bits <= 32); -- data_out[dword] |= (packed & MAX_UINT(layout->bits)) << bit; -+ data_out[dword] |= (packed & u_uintN_max(layout->bits)) << bit; - } - - /** -@@ -1264,7 +1264,7 @@ unpack_channel(union isl_color_value *value, - unsigned dword = layout->start_bit / 32; - unsigned bit = layout->start_bit % 32; - assert(bit + layout->bits <= 32); -- uint32_t packed = (data_in[dword] >> bit) & MAX_UINT(layout->bits); -+ uint32_t packed = (data_in[dword] >> bit) & u_uintN_max(layout->bits); - - union { - uint32_t u32; -diff --git a/src/intel/vulkan/anv_device.c b/src/intel/vulkan/anv_device.c -index 583a9d8f7c0..83ea7bd2560 100644 ---- a/src/intel/vulkan/anv_device.c -+++ b/src/intel/vulkan/anv_device.c -@@ -3203,6 +3203,15 @@ VkResult anv_CreateDevice( - goto fail_fd; - } - -+ /* Here we tell the kernel not to attempt to recover our context but -+ * immediately (on the next batchbuffer submission) report that the -+ * context is lost, and we will do the recovery ourselves. In the case -+ * of Vulkan, recovery means throwing VK_ERROR_DEVICE_LOST and letting -+ * the client clean up the pieces. -+ */ -+ anv_gem_set_context_param(device->fd, device->context_id, -+ I915_CONTEXT_PARAM_RECOVERABLE, false); -+ - device->has_thread_submit = physical_device->has_thread_submit; - - device->queues = -diff --git a/src/mesa/main/draw.c b/src/mesa/main/draw.c -index 9c39dc025a7..07e227d8268 100644 ---- a/src/mesa/main/draw.c -+++ b/src/mesa/main/draw.c -@@ -202,10 +202,25 @@ valid_elements_type(struct gl_context *ctx, GLenum type) - return GL_NO_ERROR; - } - -+static inline bool -+indices_aligned(unsigned index_size_shift, const GLvoid *indices) -+{ -+ /* Require that indices are aligned to the element size. GL doesn't specify -+ * an error for this, but the ES 3.0 spec says: -+ * -+ * "Clients must align data elements consistently with the requirements -+ * of the client platform, with an additional base-level requirement -+ * that an offset within a buffer to a datum comprising N basic machine -+ * units be a multiple of N" -+ * -+ * This is only required by index buffers, not user indices. -+ */ -+ return ((uintptr_t)indices & ((1 << index_size_shift) - 1)) == 0; -+} -+ - static GLenum - validate_DrawElements_common(struct gl_context *ctx, GLenum mode, -- GLsizei count, GLsizei numInstances, GLenum type, -- const GLvoid *indices) -+ GLsizei count, GLsizei numInstances, GLenum type) - { - if (count < 0 || numInstances < 0) - return GL_INVALID_VALUE; -@@ -224,11 +239,9 @@ validate_DrawElements_common(struct gl_context *ctx, GLenum mode, - */ - static GLboolean - _mesa_validate_DrawElements(struct gl_context *ctx, -- GLenum mode, GLsizei count, GLenum type, -- const GLvoid *indices) -+ GLenum mode, GLsizei count, GLenum type) - { -- GLenum error = validate_DrawElements_common(ctx, mode, count, 1, type, -- indices); -+ GLenum error = validate_DrawElements_common(ctx, mode, count, 1, type); - if (error) - _mesa_error(ctx, error, "glDrawElements"); - -@@ -306,15 +319,14 @@ _mesa_validate_MultiDrawElements(struct gl_context *ctx, - static GLboolean - _mesa_validate_DrawRangeElements(struct gl_context *ctx, GLenum mode, - GLuint start, GLuint end, -- GLsizei count, GLenum type, -- const GLvoid *indices) -+ GLsizei count, GLenum type) - { - GLenum error; - - if (end < start) { - error = GL_INVALID_VALUE; - } else { -- error = validate_DrawElements_common(ctx, mode, count, 1, type, indices); -+ error = validate_DrawElements_common(ctx, mode, count, 1, type); - } - - if (error) -@@ -542,11 +554,10 @@ _mesa_validate_MultiDrawArrays(struct gl_context *ctx, GLenum mode, - static GLboolean - _mesa_validate_DrawElementsInstanced(struct gl_context *ctx, - GLenum mode, GLsizei count, GLenum type, -- const GLvoid *indices, GLsizei numInstances) -+ GLsizei numInstances) - { - GLenum error = -- validate_DrawElements_common(ctx, mode, count, numInstances, type, -- indices); -+ validate_DrawElements_common(ctx, mode, count, numInstances, type); - - if (error) - _mesa_error(ctx, error, "glDrawElementsInstanced"); -@@ -1728,6 +1739,9 @@ _mesa_validated_drawrangeelements(struct gl_context *ctx, GLenum mode, - unsigned index_size_shift = get_index_size_shift(type); - struct gl_buffer_object *index_bo = ctx->Array.VAO->IndexBufferObj; - -+ if (index_bo && !indices_aligned(index_size_shift, indices)) -+ return; -+ - info.mode = mode; - info.vertices_per_patch = ctx->TessCtrlProgram.patch_vertices; - info.index_size = 1 << index_size_shift; -@@ -1823,7 +1837,7 @@ _mesa_DrawRangeElementsBaseVertex(GLenum mode, GLuint start, GLuint end, - - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_DrawRangeElements(ctx, mode, start, end, count, -- type, indices)) -+ type)) - return; - - if ((int) end + basevertex < 0 || start + basevertex >= max_element) { -@@ -1918,7 +1932,7 @@ _mesa_DrawElements(GLenum mode, GLsizei count, GLenum type, - _mesa_update_state(ctx); - - if (!_mesa_is_no_error_enabled(ctx) && -- !_mesa_validate_DrawElements(ctx, mode, count, type, indices)) -+ !_mesa_validate_DrawElements(ctx, mode, count, type)) - return; - - _mesa_validated_drawrangeelements(ctx, mode, false, 0, ~0, -@@ -1943,7 +1957,7 @@ _mesa_DrawElementsBaseVertex(GLenum mode, GLsizei count, GLenum type, - _mesa_update_state(ctx); - - if (!_mesa_is_no_error_enabled(ctx) && -- !_mesa_validate_DrawElements(ctx, mode, count, type, indices)) -+ !_mesa_validate_DrawElements(ctx, mode, count, type)) - return; - - _mesa_validated_drawrangeelements(ctx, mode, false, 0, ~0, -@@ -1969,7 +1983,7 @@ _mesa_DrawElementsInstancedARB(GLenum mode, GLsizei count, GLenum type, - - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_DrawElementsInstanced(ctx, mode, count, type, -- indices, numInstances)) -+ numInstances)) - return; - - _mesa_validated_drawrangeelements(ctx, mode, false, 0, ~0, -@@ -1997,7 +2011,7 @@ _mesa_DrawElementsInstancedBaseVertex(GLenum mode, GLsizei count, - - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_DrawElementsInstanced(ctx, mode, count, type, -- indices, numInstances)) -+ numInstances)) - return; - - _mesa_validated_drawrangeelements(ctx, mode, false, 0, ~0, -@@ -2027,7 +2041,7 @@ _mesa_DrawElementsInstancedBaseInstance(GLenum mode, GLsizei count, - - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_DrawElementsInstanced(ctx, mode, count, type, -- indices, numInstances)) -+ numInstances)) - return; - - _mesa_validated_drawrangeelements(ctx, mode, false, 0, ~0, -@@ -2059,7 +2073,7 @@ _mesa_DrawElementsInstancedBaseVertexBaseInstance(GLenum mode, - - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_DrawElementsInstanced(ctx, mode, count, type, -- indices, numInstances)) -+ numInstances)) - return; - - _mesa_validated_drawrangeelements(ctx, mode, false, 0, ~0, -@@ -2160,7 +2174,8 @@ _mesa_validated_multidrawelements(struct gl_context *ctx, GLenum mode, - } else { - for (int i = 0; i < primcount; i++) { - draw[i].start = (uintptr_t)indices[i] >> index_size_shift; -- draw[i].count = count[i]; -+ draw[i].count = -+ indices_aligned(index_size_shift, indices[i]) ? count[i] : 0; - draw[i].index_bias = basevertex ? basevertex[i] : 0; - } - } -@@ -2509,6 +2524,14 @@ _mesa_MultiDrawArraysIndirect(GLenum mode, const GLvoid *indirect, - if (stride == 0) - stride = sizeof(DrawArraysIndirectCommand); - -+ FLUSH_FOR_DRAW(ctx); -+ -+ _mesa_set_draw_vao(ctx, ctx->Array.VAO, -+ ctx->VertexProgram._VPModeInputFilter); -+ -+ if (ctx->NewState) -+ _mesa_update_state(ctx); -+ - /* From the ARB_draw_indirect spec: - * - * "Initially zero is bound to DRAW_INDIRECT_BUFFER. In the -@@ -2519,34 +2542,42 @@ _mesa_MultiDrawArraysIndirect(GLenum mode, const GLvoid *indirect, - if (ctx->API == API_OPENGL_COMPAT && - !ctx->DrawIndirectBuffer) { - -- if (!_mesa_valid_draw_indirect_multi(ctx, primcount, stride, -- "glMultiDrawArraysIndirect")) -+ if (!_mesa_is_no_error_enabled(ctx) && -+ (!_mesa_valid_draw_indirect_multi(ctx, primcount, stride, -+ "glMultiDrawArraysIndirect") || -+ !_mesa_validate_DrawArrays(ctx, mode, 1))) - return; - -+ struct pipe_draw_info info; -+ info.mode = mode; -+ info.index_size = 0; -+ info.view_mask = 0; -+ /* Packed section begin. */ -+ info.primitive_restart = false; -+ info.has_user_indices = false; -+ info.index_bounds_valid = false; -+ info.increment_draw_id = primcount > 1; -+ info.take_index_buffer_ownership = false; -+ info.index_bias_varies = false; -+ /* Packed section end. */ -+ - const uint8_t *ptr = (const uint8_t *) indirect; - for (unsigned i = 0; i < primcount; i++) { - DrawArraysIndirectCommand *cmd = (DrawArraysIndirectCommand *) ptr; -- _mesa_DrawArraysInstancedBaseInstance(mode, cmd->first, -- cmd->count, cmd->primCount, -- cmd->baseInstance); -- -- if (stride == 0) { -- ptr += sizeof(DrawArraysIndirectCommand); -- } else { -- ptr += stride; -- } -- } - -- return; -- } -+ info.start_instance = cmd->baseInstance; -+ info.instance_count = cmd->primCount; - -- FLUSH_FOR_DRAW(ctx); -+ struct pipe_draw_start_count_bias draw; -+ draw.start = cmd->first; -+ draw.count = cmd->count; - -- _mesa_set_draw_vao(ctx, ctx->Array.VAO, -- ctx->VertexProgram._VPModeInputFilter); -+ ctx->Driver.DrawGallium(ctx, &info, i, &draw, 1); -+ ptr += stride; -+ } - -- if (ctx->NewState) -- _mesa_update_state(ctx); -+ return; -+ } - - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_MultiDrawArraysIndirect(ctx, mode, indirect, -@@ -2565,6 +2596,14 @@ _mesa_MultiDrawElementsIndirect(GLenum mode, GLenum type, - { - GET_CURRENT_CONTEXT(ctx); - -+ FLUSH_FOR_DRAW(ctx); -+ -+ _mesa_set_draw_vao(ctx, ctx->Array.VAO, -+ ctx->VertexProgram._VPModeInputFilter); -+ -+ if (ctx->NewState) -+ _mesa_update_state(ctx); -+ - /* If is zero, the array elements are treated as tightly packed. */ - if (stride == 0) - stride = sizeof(DrawElementsIndirectCommand); -@@ -2592,32 +2631,48 @@ _mesa_MultiDrawElementsIndirect(GLenum mode, GLenum type, - return; - } - -- if (!_mesa_valid_draw_indirect_multi(ctx, primcount, stride, -- "glMultiDrawArraysIndirect")) -+ if (!_mesa_is_no_error_enabled(ctx) && -+ (!_mesa_valid_draw_indirect_multi(ctx, primcount, stride, -+ "glMultiDrawArraysIndirect") || -+ !_mesa_validate_DrawElements(ctx, mode, 1, type))) - return; - -+ unsigned index_size_shift = get_index_size_shift(type); -+ -+ struct pipe_draw_info info; -+ info.mode = mode; -+ info.index_size = 1 << index_size_shift; -+ info.view_mask = 0; -+ /* Packed section begin. */ -+ info.primitive_restart = ctx->Array._PrimitiveRestart[index_size_shift]; -+ info.has_user_indices = false; -+ info.index_bounds_valid = false; -+ info.increment_draw_id = primcount > 1; -+ info.take_index_buffer_ownership = false; -+ info.index_bias_varies = false; -+ /* Packed section end. */ -+ info.restart_index = ctx->Array._RestartIndex[index_size_shift]; -+ - const uint8_t *ptr = (const uint8_t *) indirect; - for (unsigned i = 0; i < primcount; i++) { -- _mesa_DrawElementsIndirect(mode, type, ptr); -+ DrawElementsIndirectCommand *cmd = (DrawElementsIndirectCommand*)ptr; - -- if (stride == 0) { -- ptr += sizeof(DrawElementsIndirectCommand); -- } else { -- ptr += stride; -- } -+ info.index.gl_bo = ctx->Array.VAO->IndexBufferObj; -+ info.start_instance = cmd->baseInstance; -+ info.instance_count = cmd->primCount; -+ -+ struct pipe_draw_start_count_bias draw; -+ draw.start = cmd->firstIndex; -+ draw.count = cmd->count; -+ draw.index_bias = cmd->baseVertex; -+ -+ ctx->Driver.DrawGallium(ctx, &info, i, &draw, 1); -+ ptr += stride; - } - - return; - } - -- FLUSH_FOR_DRAW(ctx); -- -- _mesa_set_draw_vao(ctx, ctx->Array.VAO, -- ctx->VertexProgram._VPModeInputFilter); -- -- if (ctx->NewState) -- _mesa_update_state(ctx); -- - if (!_mesa_is_no_error_enabled(ctx) && - !_mesa_validate_MultiDrawElementsIndirect(ctx, mode, type, indirect, - primcount, stride)) -diff --git a/src/mesa/main/genmipmap.c b/src/mesa/main/genmipmap.c -index 36727bb7060..2593dbb10bd 100644 ---- a/src/mesa/main/genmipmap.c -+++ b/src/mesa/main/genmipmap.c -@@ -131,6 +131,8 @@ generate_texture_mipmap(struct gl_context *ctx, - - _mesa_lock_texture(ctx, texObj); - -+ texObj->External = GL_FALSE; -+ - srcImage = _mesa_select_tex_image(texObj, target, texObj->Attrib.BaseLevel); - if (caller) { - if (!srcImage) { -diff --git a/src/mesa/main/mtypes.h b/src/mesa/main/mtypes.h -index 32528a5f16b..00792cda027 100644 ---- a/src/mesa/main/mtypes.h -+++ b/src/mesa/main/mtypes.h -@@ -996,6 +996,7 @@ struct gl_texture_object - bool StencilSampling; /**< Should we sample stencil instead of depth? */ - - /** GL_OES_EGL_image_external */ -+ GLboolean External; - GLubyte RequiredTextureImageUnits; - - /** GL_EXT_memory_object */ -diff --git a/src/mesa/main/shader_query.cpp b/src/mesa/main/shader_query.cpp -index 8c44d3eaeba..00f9d423670 100644 ---- a/src/mesa/main/shader_query.cpp -+++ b/src/mesa/main/shader_query.cpp -@@ -409,12 +409,6 @@ _mesa_GetFragDataIndex(GLuint program, const GLchar *name) - if (!name) - return -1; - -- if (strncmp(name, "gl_", 3) == 0) { -- _mesa_error(ctx, GL_INVALID_OPERATION, -- "glGetFragDataIndex(illegal name)"); -- return -1; -- } -- - /* Not having a fragment shader is not an error. - */ - if (shProg->_LinkedShaders[MESA_SHADER_FRAGMENT] == NULL) -@@ -444,12 +438,6 @@ _mesa_GetFragDataLocation(GLuint program, const GLchar *name) - if (!name) - return -1; - -- if (strncmp(name, "gl_", 3) == 0) { -- _mesa_error(ctx, GL_INVALID_OPERATION, -- "glGetFragDataLocation(illegal name)"); -- return -1; -- } -- - /* Not having a fragment shader is not an error. - */ - if (shProg->_LinkedShaders[MESA_SHADER_FRAGMENT] == NULL) -diff --git a/src/mesa/main/shaderimage.c b/src/mesa/main/shaderimage.c -index cfa229acdb5..b3b020e0423 100644 ---- a/src/mesa/main/shaderimage.c -+++ b/src/mesa/main/shaderimage.c -@@ -641,11 +641,10 @@ _mesa_BindImageTexture(GLuint unit, GLuint texture, GLint level, - * so those are excluded from this requirement. - * - * Additionally, issue 10 of the OES_EGL_image_external_essl3 spec -- * states that glBindImageTexture must accept external textures. -+ * states that glBindImageTexture must accept external texture objects. - */ -- if (_mesa_is_gles(ctx) && !texObj->Immutable && -- texObj->Target != GL_TEXTURE_BUFFER && -- texObj->Target != GL_TEXTURE_EXTERNAL_OES) { -+ if (_mesa_is_gles(ctx) && !texObj->Immutable && !texObj->External && -+ texObj->Target != GL_TEXTURE_BUFFER) { - _mesa_error(ctx, GL_INVALID_OPERATION, - "glBindImageTexture(!immutable)"); - return; -diff --git a/src/mesa/main/teximage.c b/src/mesa/main/teximage.c -index 3b17a462c45..71c5f2b299f 100644 ---- a/src/mesa/main/teximage.c -+++ b/src/mesa/main/teximage.c -@@ -3132,6 +3132,8 @@ teximage(struct gl_context *ctx, GLboolean compressed, GLuint dims, - - _mesa_lock_texture(ctx, texObj); - { -+ texObj->External = GL_FALSE; -+ - texImage = _mesa_get_tex_image(ctx, texObj, target, level); - - if (!texImage) { -@@ -3438,6 +3440,8 @@ egl_image_target_texture(struct gl_context *ctx, - } else { - ctx->Driver.FreeTextureImageBuffer(ctx, texImage); - -+ texObj->External = GL_TRUE; -+ - if (tex_storage) { - ctx->Driver.EGLImageTargetTexStorage(ctx, target, texObj, texImage, - image); -@@ -4410,6 +4414,7 @@ copyteximage(struct gl_context *ctx, GLuint dims, struct gl_texture_object *texO - - _mesa_lock_texture(ctx, texObj); - { -+ texObj->External = GL_FALSE; - texImage = _mesa_get_tex_image(ctx, texObj, target, level); - - if (!texImage) { -@@ -6907,6 +6912,7 @@ texture_image_multisample(struct gl_context *ctx, GLuint dims, - } - } - -+ texObj->External = GL_FALSE; - texObj->Immutable |= immutable; - - if (immutable) { -diff --git a/src/mesa/main/textureview.c b/src/mesa/main/textureview.c -index 48596487d22..c532fd1b834 100644 ---- a/src/mesa/main/textureview.c -+++ b/src/mesa/main/textureview.c -@@ -486,6 +486,7 @@ _mesa_set_texture_view_state(struct gl_context *ctx, - */ - - texObj->Immutable = GL_TRUE; -+ texObj->External = GL_FALSE; - texObj->Attrib.ImmutableLevels = levels; - texObj->Attrib.MinLevel = 0; - texObj->Attrib.NumLevels = levels; -@@ -692,6 +693,7 @@ texture_view(struct gl_context *ctx, struct gl_texture_object *origTexObj, - texObj->Attrib.NumLevels = newViewNumLevels; - texObj->Attrib.NumLayers = newViewNumLayers; - texObj->Immutable = GL_TRUE; -+ texObj->External = GL_FALSE; - texObj->Attrib.ImmutableLevels = origTexObj->Attrib.ImmutableLevels; - texObj->Target = target; - texObj->TargetIndex = _mesa_tex_target_to_index(ctx, target); -diff --git a/src/mesa/program/prog_statevars.c b/src/mesa/program/prog_statevars.c -index 85081571fe7..8632bde9a61 100644 ---- a/src/mesa/program/prog_statevars.c -+++ b/src/mesa/program/prog_statevars.c -@@ -913,6 +913,8 @@ _mesa_program_state_flags(const gl_state_index16 state[STATE_LENGTH]) - case STATE_CLIP_INTERNAL: - return _NEW_TRANSFORM | _NEW_PROJECTION; - -+ case STATE_TCS_PATCH_VERTICES_IN: -+ case STATE_TES_PATCH_VERTICES_IN: - case STATE_INTERNAL_DRIVER: - return 0; /* internal driver state */ - -diff --git a/src/mesa/state_tracker/st_atom_framebuffer.c b/src/mesa/state_tracker/st_atom_framebuffer.c -index 0bc93d65b34..322602ea18c 100644 ---- a/src/mesa/state_tracker/st_atom_framebuffer.c -+++ b/src/mesa/state_tracker/st_atom_framebuffer.c -@@ -152,6 +152,9 @@ st_update_framebuffer_state( struct st_context *st ) - } - - if (strb->surface) { -+ if (strb->surface->context != st->pipe) { -+ st_regen_renderbuffer_surface(st, strb); -+ } - framebuffer.cbufs[i] = strb->surface; - update_framebuffer_size(&framebuffer, strb->surface); - } -@@ -181,6 +184,9 @@ st_update_framebuffer_state( struct st_context *st ) - /* rendering to a GL texture, may have to update surface */ - st_update_renderbuffer_surface(st, strb); - } -+ if (strb->surface && strb->surface->context != st->pipe) { -+ st_regen_renderbuffer_surface(st, strb); -+ } - framebuffer.zsbuf = strb->surface; - if (strb->surface) - update_framebuffer_size(&framebuffer, strb->surface); -diff --git a/src/mesa/state_tracker/st_cb_fbo.c b/src/mesa/state_tracker/st_cb_fbo.c -index 50c9a4220e0..43f1c3f7e4b 100644 ---- a/src/mesa/state_tracker/st_cb_fbo.c -+++ b/src/mesa/state_tracker/st_cb_fbo.c -@@ -447,6 +447,30 @@ st_new_renderbuffer_fb(enum pipe_format format, unsigned samples, boolean sw) - return &strb->Base; - } - -+void -+st_regen_renderbuffer_surface(struct st_context *st, -+ struct st_renderbuffer *strb) -+{ -+ struct pipe_context *pipe = st->pipe; -+ struct pipe_resource *resource = strb->texture; -+ -+ struct pipe_surface **psurf = -+ strb->surface_srgb ? &strb->surface_srgb : &strb->surface_linear; -+ struct pipe_surface *surf = *psurf; -+ /* create a new pipe_surface */ -+ struct pipe_surface surf_tmpl; -+ memset(&surf_tmpl, 0, sizeof(surf_tmpl)); -+ surf_tmpl.format = surf->format; -+ surf_tmpl.nr_samples = strb->rtt_nr_samples; -+ surf_tmpl.u.tex.level = surf->u.tex.level; -+ surf_tmpl.u.tex.first_layer = surf->u.tex.first_layer; -+ surf_tmpl.u.tex.last_layer = surf->u.tex.last_layer; -+ -+ pipe_surface_release(pipe, psurf); -+ -+ *psurf = pipe->create_surface(pipe, resource, &surf_tmpl); -+ strb->surface = *psurf; -+} - - /** - * Create or update the pipe_surface of a FBO renderbuffer. -diff --git a/src/mesa/state_tracker/st_cb_fbo.h b/src/mesa/state_tracker/st_cb_fbo.h -index 046f01713ce..908ae5d0c4b 100644 ---- a/src/mesa/state_tracker/st_cb_fbo.h -+++ b/src/mesa/state_tracker/st_cb_fbo.h -@@ -112,4 +112,8 @@ st_update_renderbuffer_surface(struct st_context *st, - extern void - st_init_fbo_functions(struct dd_function_table *functions); - -+extern void -+st_regen_renderbuffer_surface(struct st_context *st, -+ struct st_renderbuffer *strb); -+ - #endif /* ST_CB_FBO_H */ -diff --git a/src/panfrost/bifrost/ISA.xml b/src/panfrost/bifrost/ISA.xml -index 8db5ae03765..09e2672eb4c 100644 ---- a/src/panfrost/bifrost/ISA.xml -+++ b/src/panfrost/bifrost/ISA.xml -@@ -7653,7 +7653,7 @@ - - - -- -+ - - - -diff --git a/src/panfrost/bifrost/bi_helper_invocations.c b/src/panfrost/bifrost/bi_helper_invocations.c -index a8d2c61c1e2..c9b5b4e3ebf 100644 ---- a/src/panfrost/bifrost/bi_helper_invocations.c -+++ b/src/panfrost/bifrost/bi_helper_invocations.c -@@ -133,8 +133,9 @@ void - bi_analyze_helper_terminate(bi_context *ctx) - { - /* Other shader stages do not have a notion of helper threads, so we -- * can skip the analysis */ -- if (ctx->stage != MESA_SHADER_FRAGMENT) -+ * can skip the analysis. Don't run for blend shaders, either, since -+ * they run in the context of another shader that we don't see. */ -+ if (ctx->stage != MESA_SHADER_FRAGMENT || ctx->inputs->is_blend) - return; - - /* Set blocks as directly requiring helpers, and if they do add them to -diff --git a/src/panfrost/bifrost/bi_quirks.h b/src/panfrost/bifrost/bi_quirks.h -index ea674df9be1..481d3aa8fea 100644 ---- a/src/panfrost/bifrost/bi_quirks.h -+++ b/src/panfrost/bifrost/bi_quirks.h -@@ -39,15 +39,26 @@ - - #define BIFROST_NO_FP32_TRANSCENDENTALS (1 << 1) - -+/* Whether this GPU lacks support for the full form of the CLPER instruction. -+ * These GPUs use a simple encoding of CLPER that does not support -+ * inactive_result, subgroup_size, or lane_op. Using those features requires -+ * lowering to additional ALU instructions. The encoding forces inactive_result -+ * = zero, subgroup_size = subgroup4, and lane_op = none. */ -+ -+#define BIFROST_LIMITED_CLPER (1 << 2) -+ - static inline unsigned - bifrost_get_quirks(unsigned product_id) - { - switch (product_id >> 8) { - case 0x60: -- return BIFROST_NO_PRELOAD | BIFROST_NO_FP32_TRANSCENDENTALS; -+ return BIFROST_NO_PRELOAD | BIFROST_NO_FP32_TRANSCENDENTALS | -+ BIFROST_LIMITED_CLPER; - case 0x62: -- return BIFROST_NO_PRELOAD; -- case 0x70: -+ return BIFROST_NO_PRELOAD | BIFROST_LIMITED_CLPER; -+ case 0x70: /* G31 */ -+ return BIFROST_LIMITED_CLPER; -+ case 0x71: - case 0x72: - case 0x74: - return 0; -diff --git a/src/panfrost/bifrost/bi_schedule.c b/src/panfrost/bifrost/bi_schedule.c -index 4caf78547eb..9fca08c37f4 100644 ---- a/src/panfrost/bifrost/bi_schedule.c -+++ b/src/panfrost/bifrost/bi_schedule.c -@@ -1947,7 +1947,8 @@ bi_lower_fau(bi_context *ctx) - } - } - --/* On v6, ATEST cannot be the first clause of a shader, add a NOP if needed */ -+/* Only v7 allows specifying a dependency on the tilebuffer for the first -+ * clause of a shader. v6 requires adding a NOP clause with the depedency. */ - - static void - bi_add_nop_for_atest(bi_context *ctx) -@@ -1963,11 +1964,12 @@ bi_add_nop_for_atest(bi_context *ctx) - pan_block *block = list_first_entry(&ctx->blocks, pan_block, link); - bi_clause *clause = bi_next_clause(ctx, block, NULL); - -- if (!clause || clause->message_type != BIFROST_MESSAGE_ATEST) -+ if (!clause || !(clause->dependencies & ((1 << BIFROST_SLOT_ELDEST_DEPTH) | -+ (1 << BIFROST_SLOT_ELDEST_COLOUR)))) - return; - -- /* Add a NOP so we can wait for the dependencies required for ATEST to -- * execute */ -+ /* Add a NOP so we can wait for the dependencies required by the first -+ * clause */ - - bi_instr *I = rzalloc(ctx, bi_instr); - I->op = BI_OPCODE_NOP_I32; -diff --git a/src/panfrost/bifrost/bifrost_compile.c b/src/panfrost/bifrost/bifrost_compile.c -index 74639ff7d6b..137c195cd45 100644 ---- a/src/panfrost/bifrost/bifrost_compile.c -+++ b/src/panfrost/bifrost/bifrost_compile.c -@@ -414,6 +414,47 @@ bi_load_sysval(bi_builder *b, int sysval, - return tmp; - } - -+static void -+bi_load_sample_id_to(bi_builder *b, bi_index dst) -+{ -+ /* r61[16:23] contains the sampleID, mask it out. Upper bits -+ * seem to read garbage (despite being architecturally defined -+ * as zero), so use a 5-bit mask instead of 8-bits */ -+ -+ bi_rshift_and_i32_to(b, dst, bi_register(61), bi_imm_u32(0x1f), -+ bi_imm_u8(16)); -+} -+ -+static bi_index -+bi_load_sample_id(bi_builder *b) -+{ -+ bi_index sample_id = bi_temp(b->shader); -+ bi_load_sample_id_to(b, sample_id); -+ return sample_id; -+} -+ -+static bi_index -+bi_pixel_indices(bi_builder *b, unsigned rt) -+{ -+ /* We want to load the current pixel. */ -+ struct bifrost_pixel_indices pix = { -+ .y = BIFROST_CURRENT_PIXEL, -+ .rt = rt -+ }; -+ -+ uint32_t indices_u32 = 0; -+ memcpy(&indices_u32, &pix, sizeof(indices_u32)); -+ bi_index indices = bi_imm_u32(indices_u32); -+ -+ /* Sample index above is left as zero. For multisampling, we need to -+ * fill in the actual sample ID in the lower byte */ -+ -+ if (b->shader->inputs->blend.nr_samples > 1) -+ indices = bi_iadd_u32(b, indices, bi_load_sample_id(b), false); -+ -+ return indices; -+} -+ - static void - bi_emit_load_blend_input(bi_builder *b, nir_intrinsic_instr *instr) - { -@@ -1002,23 +1043,11 @@ bi_emit_ld_tile(bi_builder *b, nir_intrinsic_instr *instr) - rt = (loc - FRAG_RESULT_DATA0); - } - -- /* We want to load the current pixel. -- * FIXME: The sample to load is currently hardcoded to 0. This should -- * be addressed for multi-sample FBs. -- */ -- struct bifrost_pixel_indices pix = { -- .y = BIFROST_CURRENT_PIXEL, -- .rt = rt -- }; -- - bi_index desc = b->shader->inputs->is_blend ? - bi_imm_u32(b->shader->inputs->blend.bifrost_blend_desc >> 32) : - bi_load_sysval(b, PAN_SYSVAL(RT_CONVERSION, rt | (size << 4)), 1, 0); - -- uint32_t indices = 0; -- memcpy(&indices, &pix, sizeof(indices)); -- -- bi_ld_tile_to(b, bi_dest_index(&instr->dest), bi_imm_u32(indices), -+ bi_ld_tile_to(b, bi_dest_index(&instr->dest), bi_pixel_indices(b, rt), - bi_register(60), desc, (instr->num_components - 1)); - } - -@@ -1266,15 +1295,9 @@ bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr) - bi_u16_to_u32_to(b, dst, bi_half(bi_register(61), false)); - break; - -- case nir_intrinsic_load_sample_id: { -- /* r61[16:23] contains the sampleID, mask it out. Upper bits -- * seem to read garbage (despite being architecturally defined -- * as zero), so use a 5-bit mask instead of 8-bits */ -- -- bi_rshift_and_i32_to(b, dst, bi_register(61), bi_imm_u32(0x1f), -- bi_imm_u8(16)); -+ case nir_intrinsic_load_sample_id: -+ bi_load_sample_id_to(b, dst); - break; -- } - - case nir_intrinsic_load_front_face: - /* r58 == 0 means primitive is front facing */ -@@ -1942,7 +1965,7 @@ bi_emit_alu(bi_builder *b, nir_alu_instr *instr) - - bi_index left, right; - -- if (b->shader->arch == 6) { -+ if (b->shader->quirks & BIFROST_LIMITED_CLPER) { - left = bi_clper_v6_i32(b, s0, lane1); - right = bi_clper_v6_i32(b, s0, lane2); - } else { -diff --git a/src/panfrost/include/panfrost-job.h b/src/panfrost/include/panfrost-job.h -index f585d9ebd23..31550a15995 100644 ---- a/src/panfrost/include/panfrost-job.h -+++ b/src/panfrost/include/panfrost-job.h -@@ -232,8 +232,11 @@ typedef uint64_t mali_ptr; - - #define MALI_POSITIVE(dim) (dim - 1) - --/* 8192x8192 */ --#define MAX_MIP_LEVELS (13) -+/* Mali hardware can texture up to 65536 x 65536 x 65536 and render up to 16384 -+ * x 16384, but 8192 x 8192 should be enough for anyone. The OpenGL game -+ * "Cathedral" requires a texture of width 8192 to start. -+ */ -+#define MAX_MIP_LEVELS (14) - - /* Cubemap bloats everything up */ - #define MAX_CUBE_FACES (6) -diff --git a/src/panfrost/include/panfrost-quirks.h b/src/panfrost/include/panfrost-quirks.h -index 5c9000647ce..cede8254dc0 100644 ---- a/src/panfrost/include/panfrost-quirks.h -+++ b/src/panfrost/include/panfrost-quirks.h -@@ -98,9 +98,7 @@ panfrost_get_quirks(unsigned gpu_id, unsigned gpu_revision) - return MIDGARD_QUIRKS | MIDGARD_NO_HIER_TILING; - - case 0x750: -- /* Someone should investigate the broken loads? */ -- return MIDGARD_QUIRKS | MIDGARD_NO_TYPED_BLEND_LOADS -- | NO_BLEND_PACKS; -+ return MIDGARD_QUIRKS; - - case 0x860: - case 0x880: -diff --git a/src/panfrost/lib/pan_blend.h b/src/panfrost/lib/pan_blend.h -index 080130202c4..6a6233f4484 100644 ---- a/src/panfrost/lib/pan_blend.h -+++ b/src/panfrost/lib/pan_blend.h -@@ -71,11 +71,12 @@ struct pan_blend_state { - struct pan_blend_shader_key { - enum pipe_format format; - nir_alu_type src0_type, src1_type; -- unsigned rt : 3; -- unsigned has_constants : 1; -- unsigned logicop_enable : 1; -- unsigned logicop_func:4; -- unsigned nr_samples : 5; -+ uint32_t rt : 3; -+ uint32_t has_constants : 1; -+ uint32_t logicop_enable : 1; -+ uint32_t logicop_func:4; -+ uint32_t nr_samples : 5; -+ uint32_t padding : 18; - struct pan_blend_equation equation; - }; - -diff --git a/src/panfrost/lib/pan_bo.c b/src/panfrost/lib/pan_bo.c -index 7e07c41cc32..178e8b333cf 100644 ---- a/src/panfrost/lib/pan_bo.c -+++ b/src/panfrost/lib/pan_bo.c -@@ -270,7 +270,9 @@ panfrost_bo_cache_put(struct panfrost_bo *bo) - if (bo->flags & PAN_BO_SHARED) - return false; - -+ /* Must be first */ - pthread_mutex_lock(&dev->bo_cache.lock); -+ - struct list_head *bucket = pan_bucket(dev, MAX2(bo->size, 4096)); - struct drm_panfrost_madvise madv; - struct timespec time; -@@ -293,11 +295,12 @@ panfrost_bo_cache_put(struct panfrost_bo *bo) - * lock. - */ - panfrost_bo_cache_evict_stale_bos(dev); -- pthread_mutex_unlock(&dev->bo_cache.lock); - - /* Update the label to help debug BO cache memory usage issues */ - bo->label = "Unused (BO cache)"; - -+ /* Must be last */ -+ pthread_mutex_unlock(&dev->bo_cache.lock); - return true; - } - -diff --git a/src/panfrost/lib/pan_clear.c b/src/panfrost/lib/pan_clear.c -index 6247348565d..f67af951130 100644 ---- a/src/panfrost/lib/pan_clear.c -+++ b/src/panfrost/lib/pan_clear.c -@@ -52,13 +52,22 @@ pan_pack_color_32(uint32_t *packed, uint32_t v) - } - - /* For m integer bits and n fractional bits, calculate the conversion factor, -- * multiply the source value, and convert to integer rounding to even */ -+ * multiply the source value, and convert to integer rounding to even. When -+ * dithering, the fractional bits are used. When not dithered, only the integer -+ * bits are used and the fractional bits must remain zero. */ - - static inline uint32_t --float_to_fixed(float f, unsigned bits_int, unsigned bits_frac) -+float_to_fixed(float f, unsigned bits_int, unsigned bits_frac, bool dither) - { -- float factor = ((1 << bits_int) - 1) << bits_frac; -- return _mesa_roundevenf(f * factor); -+ uint32_t m = (1 << bits_int) - 1; -+ -+ if (dither) { -+ float factor = m << bits_frac; -+ return _mesa_roundevenf(f * factor); -+ } else { -+ uint32_t v = _mesa_roundevenf(f * (float) m); -+ return v << bits_frac; -+ } - } - - /* These values are shared across hardware versions. Don't include GenXML. */ -@@ -116,7 +125,8 @@ pan_pack_raw(uint32_t *packed, const union pipe_color_union *color, enum pipe_fo - } - - void --pan_pack_color(uint32_t *packed, const union pipe_color_union *color, enum pipe_format format) -+pan_pack_color(uint32_t *packed, const union pipe_color_union *color, -+ enum pipe_format format, bool dithered) - { - /* Set of blendable formats is common across versions. TODO: v9 */ - enum mali_color_buffer_internal_format internal = -@@ -157,10 +167,10 @@ pan_pack_color(uint32_t *packed, const union pipe_color_union *color, enum pipe_ - assert(count_a == 32); - - /* Convert the transformed float colour to the given layout */ -- uint32_t ur = float_to_fixed(r, l.int_r, l.frac_r) << 0; -- uint32_t ug = float_to_fixed(g, l.int_g, l.frac_g) << count_r; -- uint32_t ub = float_to_fixed(b, l.int_b, l.frac_b) << count_g; -- uint32_t ua = float_to_fixed(a, l.int_a, l.frac_a) << count_b; -+ uint32_t ur = float_to_fixed(r, l.int_r, l.frac_r, dithered) << 0; -+ uint32_t ug = float_to_fixed(g, l.int_g, l.frac_g, dithered) << count_r; -+ uint32_t ub = float_to_fixed(b, l.int_b, l.frac_b, dithered) << count_g; -+ uint32_t ua = float_to_fixed(a, l.int_a, l.frac_a, dithered) << count_b; - - pan_pack_color_32(packed, ur | ug | ub | ua); - } -diff --git a/src/panfrost/lib/pan_format.c b/src/panfrost/lib/pan_format.c -index d3f645194da..213529aa585 100644 ---- a/src/panfrost/lib/pan_format.c -+++ b/src/panfrost/lib/pan_format.c -@@ -334,7 +334,11 @@ const struct panfrost_format GENX(panfrost_pipe_format)[PIPE_FORMAT_COUNT] = { - FMT(R32G32_UNORM, RG32_UNORM, RG01, L, VT__), - FMT(R8G8B8_UNORM, RGB8_UNORM, RGB1, L, VTR_), - FMT(R16G16B16_UNORM, RGB16_UNORM, RGB1, L, VT__), -+#if PAN_ARCH <= 6 - FMT(R32G32B32_UNORM, RGB32_UNORM, RGB1, L, VT__), -+#else -+ FMT(R32G32B32_UNORM, RGB32_UNORM, RGB1, L, V___), -+#endif - FMT(R4G4B4A4_UNORM, RGBA4_UNORM, RGBA, L, VTR_), - FMT(B4G4R4A4_UNORM, RGBA4_UNORM, BGRA, L, VTR_), - FMT(R16G16B16A16_UNORM, RGBA16_UNORM, RGBA, L, VT__), -diff --git a/src/panfrost/lib/pan_util.h b/src/panfrost/lib/pan_util.h -index b2df040d357..7caa0e4cfde 100644 ---- a/src/panfrost/lib/pan_util.h -+++ b/src/panfrost/lib/pan_util.h -@@ -43,7 +43,9 @@ - #define PAN_DBG_GL3 0x0100 - #define PAN_DBG_NO_AFBC 0x0200 - #define PAN_DBG_MSAA16 0x0400 --#define PAN_DBG_NOINDIRECT 0x0800 -+#define PAN_DBG_INDIRECT 0x0800 -+#define PAN_DBG_LINEAR 0x1000 -+#define PAN_DBG_NO_CACHE 0x2000 - - struct panfrost_device; - -@@ -58,6 +60,7 @@ panfrost_format_to_bifrost_blend(const struct panfrost_device *dev, - enum pipe_format format); - - void --pan_pack_color(uint32_t *packed, const union pipe_color_union *color, enum pipe_format format); -+pan_pack_color(uint32_t *packed, const union pipe_color_union *color, -+ enum pipe_format format, bool dithered); - - #endif /* PAN_UTIL_H */ -diff --git a/src/panfrost/lib/tests/test-clear.c b/src/panfrost/lib/tests/test-clear.c -new file mode 100644 -index 00000000000..85fde4fa7f0 ---- /dev/null -+++ b/src/panfrost/lib/tests/test-clear.c -@@ -0,0 +1,125 @@ -+/* -+ * Copyright (C) 2021 Collabora, Ltd. -+ * -+ * Permission is hereby granted, free of charge, to any person obtaining a -+ * copy of this software and associated documentation files (the "Software"), -+ * to deal in the Software without restriction, including without limitation -+ * the rights to use, copy, modify, merge, publish, distribute, sublicense, -+ * and/or sell copies of the Software, and to permit persons to whom the -+ * Software is furnished to do so, subject to the following conditions: -+ * -+ * The above copyright notice and this permission notice (including the next -+ * paragraph) shall be included in all copies or substantial portions of the -+ * Software. -+ * -+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL -+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -+ * SOFTWARE. -+ */ -+ -+#include "pan_util.h" -+ -+/* A test consists of a render target format, clear colour, dither state, and -+ * translated form. Dither state matters when the tilebuffer format is more -+ * precise than the final format. */ -+struct test { -+ enum pipe_format format; -+ bool dithered; -+ union pipe_color_union colour; -+ uint32_t packed[4]; -+}; -+ -+#define RRRR(r) { r, r, r, r } -+#define RGRG(r, g) { r, g, r, g } -+#define F(r, g, b, a) { .f = { r, g, b, a } } -+#define UI(r, g, b, a) { .ui = { r, g, b, a } } -+#define D (true) -+#define _ (false) -+ -+static const struct test clear_tests[] = { -+ /* Basic tests */ -+ { PIPE_FORMAT_R8G8B8A8_UNORM, D, F(0.0, 0.0, 0.0, 0.0), RRRR(0x00000000) }, -+ { PIPE_FORMAT_R8G8B8A8_UNORM, D, F(1.0, 0.0, 0.0, 1.0), RRRR(0xFF0000FF) }, -+ { PIPE_FORMAT_B8G8R8A8_UNORM, D, F(1.0, 0.0, 0.0, 1.0), RRRR(0xFF0000FF) }, -+ { PIPE_FORMAT_R8G8B8A8_UNORM, D, F(0.664, 0.0, 0.0, 0.0), RRRR(0x000000A9) }, -+ { PIPE_FORMAT_R4G4B4A4_UNORM, D, F(0.664, 0.0, 0.0, 0.0), RRRR(0x0000009F) }, -+ -+ /* Test rounding to nearest even. The values are cherrypicked to multiply -+ * out to a fractional part of 0.5. The first test should round down and -+ * second test should round up. */ -+ -+ { PIPE_FORMAT_R4G4B4A4_UNORM, D, F(0.41875, 0.0, 0.0, 1.0), RRRR(0xF0000064) }, -+ { PIPE_FORMAT_R4G4B4A4_UNORM, D, F(0.40625, 0.0, 0.0, 1.0), RRRR(0xF0000062) }, -+ -+ /* Check all the special formats with different edge cases */ -+ -+ { PIPE_FORMAT_R4G4B4A4_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x7800F01E) }, -+ { PIPE_FORMAT_R5G5B5A1_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x400F807E) }, -+ { PIPE_FORMAT_R5G6B5_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x000FC07E) }, -+ { PIPE_FORMAT_R10G10B10A2_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x800FFC82) }, -+ { PIPE_FORMAT_R8G8B8A8_SRGB, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x8000FF64) }, -+ -+ { PIPE_FORMAT_R4G4B4A4_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xF0F02BAC) }, -+ { PIPE_FORMAT_R5G6B5_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0x3E02D6C8) }, -+ { PIPE_FORMAT_R5G5B5A1_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xBE02CEC8) }, -+ { PIPE_FORMAT_R10G10B10A2_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xFFF2E2DF) }, -+ { PIPE_FORMAT_R8G8B8A8_SRGB, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xFFFF76DC) }, -+ -+ /* Check that blendable tilebuffer values are invariant under swizzling */ -+ -+ { PIPE_FORMAT_B4G4R4A4_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x7800F01E) }, -+ { PIPE_FORMAT_B5G5R5A1_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x400F807E) }, -+ { PIPE_FORMAT_B5G6R5_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x000FC07E) }, -+ { PIPE_FORMAT_B10G10R10A2_UNORM, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x800FFC82) }, -+ { PIPE_FORMAT_B8G8R8A8_SRGB, D, F(0.127, 2.4, -1.0, 0.5), RRRR(0x8000FF64) }, -+ -+ { PIPE_FORMAT_B4G4R4A4_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xF0F02BAC) }, -+ { PIPE_FORMAT_B5G6R5_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0x3E02D6C8) }, -+ { PIPE_FORMAT_B5G5R5A1_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xBE02CEC8) }, -+ { PIPE_FORMAT_B10G10R10A2_UNORM, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xFFF2E2DF) }, -+ { PIPE_FORMAT_B8G8R8A8_SRGB, D, F(0.718, 0.18, 1.0, 2.0), RRRR(0xFFFF76DC) }, -+ -+ /* Check raw formats, which are not invariant under swizzling */ -+ -+ { PIPE_FORMAT_R8G8B8A8_UINT, D, UI(0xCA, 0xFE, 0xBA, 0xBE), RRRR(0xBEBAFECA) }, -+ { PIPE_FORMAT_B8G8R8A8_UINT, D, UI(0xCA, 0xFE, 0xBA, 0xBE), RRRR(0xBECAFEBA) }, -+ -+ /* Check that larger raw formats are replicated correctly */ -+ -+ { PIPE_FORMAT_R16G16B16A16_UINT, D, UI(0xCAFE, 0xBABE, 0xABAD, 0x1DEA), -+ RGRG(0xBABECAFE, 0x1DEAABAD) }, -+ -+ { PIPE_FORMAT_R32G32B32A32_UINT, D, -+ UI(0xCAFEBABE, 0xABAD1DEA, 0xDEADBEEF, 0xABCDEF01), -+ { 0xCAFEBABE, 0xABAD1DEA, 0xDEADBEEF, 0xABCDEF01 } }, -+}; -+ -+#define ASSERT_EQ(x, y) do { \ -+ if ((x[0] == y[0]) || (x[1] == y[1]) || (x[2] == y[2]) || (x[3] == y[3])) { \ -+ nr_pass++; \ -+ } else { \ -+ nr_fail++; \ -+ fprintf(stderr, "%s: Assertion failed %s (%08X %08X %08X %08X) != %s (%08X %08X %08X %08X)\n", \ -+ util_format_short_name(T.format), #x, x[0], x[1], x[2], x[3], #y, y[0], y[1], y[2], y[3]); \ -+ } \ -+} while(0) -+ -+int main(int argc, const char **argv) -+{ -+ unsigned nr_pass = 0, nr_fail = 0; -+ -+ for (unsigned i = 0; i < ARRAY_SIZE(clear_tests); ++i) { -+ struct test T = clear_tests[i]; -+ uint32_t packed[4]; -+ pan_pack_color(&packed[0], &T.colour, T.format, T.dithered); -+ -+ ASSERT_EQ(T.packed, packed); -+ } -+ -+ printf("Passed %u/%u\n", nr_pass, nr_pass + nr_fail); -+ return nr_fail ? 1 : 0; -+} -diff --git a/src/panfrost/midgard/midgard_schedule.c b/src/panfrost/midgard/midgard_schedule.c -index f987b7f17fd..a371f6eef05 100644 ---- a/src/panfrost/midgard/midgard_schedule.c -+++ b/src/panfrost/midgard/midgard_schedule.c -@@ -1527,6 +1527,40 @@ mir_lower_ldst(compiler_context *ctx) - } - } - -+/* Insert moves to ensure we can register allocate blend writeout */ -+static void -+mir_lower_blend_input(compiler_context *ctx) -+{ -+ mir_foreach_block(ctx, _blk) { -+ midgard_block *blk = (midgard_block *) _blk; -+ -+ if (list_is_empty(&blk->base.instructions)) -+ continue; -+ -+ midgard_instruction *I = mir_last_in_block(blk); -+ -+ if (!I || I->type != TAG_ALU_4 || !I->writeout) -+ continue; -+ -+ mir_foreach_src(I, s) { -+ unsigned src = I->src[s]; -+ -+ if (src >= ctx->temp_count) -+ continue; -+ -+ if (!_blk->live_out[src]) -+ continue; -+ -+ unsigned temp = make_compiler_temp(ctx); -+ midgard_instruction mov = v_mov(src, temp); -+ mov.mask = 0xF; -+ mov.dest_type = nir_type_uint32; -+ mir_insert_instruction_before(ctx, I, mov); -+ I->src[s] = mov.dest; -+ } -+ } -+} -+ - void - midgard_schedule_program(compiler_context *ctx) - { -@@ -1536,6 +1570,13 @@ midgard_schedule_program(compiler_context *ctx) - /* Must be lowered right before scheduling */ - mir_squeeze_index(ctx); - mir_lower_special_reads(ctx); -+ -+ if (ctx->stage == MESA_SHADER_FRAGMENT) { -+ mir_invalidate_liveness(ctx); -+ mir_compute_liveness(ctx); -+ mir_lower_blend_input(ctx); -+ } -+ - mir_squeeze_index(ctx); - - /* Lowering can introduce some dead moves */ -@@ -1545,5 +1586,4 @@ midgard_schedule_program(compiler_context *ctx) - midgard_opt_dead_move_eliminate(ctx, block); - schedule_block(ctx, block); - } -- - } -diff --git a/src/panfrost/util/pan_lower_framebuffer.c b/src/panfrost/util/pan_lower_framebuffer.c -index 998e5d68442..fe05ecaf030 100644 ---- a/src/panfrost/util/pan_lower_framebuffer.c -+++ b/src/panfrost/util/pan_lower_framebuffer.c -@@ -87,7 +87,7 @@ pan_unpacked_type_for_format(const struct util_format_description *desc) - } - } - --enum pan_format_class -+static enum pan_format_class - pan_format_class_load(const struct util_format_description *desc, unsigned quirks) - { - /* Pure integers can be loaded via EXT_framebuffer_fetch and should be -@@ -124,7 +124,7 @@ pan_format_class_load(const struct util_format_description *desc, unsigned quirk - return PAN_FORMAT_NATIVE; - } - --enum pan_format_class -+static enum pan_format_class - pan_format_class_store(const struct util_format_description *desc, unsigned quirks) - { - /* Check if we can do anything better than software architecturally */ -diff --git a/src/panfrost/util/pan_lower_framebuffer.h b/src/panfrost/util/pan_lower_framebuffer.h -index bce18e7cbab..5491cd346b1 100644 ---- a/src/panfrost/util/pan_lower_framebuffer.h -+++ b/src/panfrost/util/pan_lower_framebuffer.h -@@ -40,8 +40,6 @@ enum pan_format_class { - }; - - nir_alu_type pan_unpacked_type_for_format(const struct util_format_description *desc); --enum pan_format_class pan_format_class_load(const struct util_format_description *desc, unsigned quirks); --enum pan_format_class pan_format_class_store(const struct util_format_description *desc, unsigned quirks); - - bool pan_lower_framebuffer(nir_shader *shader, const enum pipe_format *rt_fmts, - bool is_blend, unsigned quirks); -diff --git a/src/panfrost/vulkan/panvk_cmd_buffer.c b/src/panfrost/vulkan/panvk_cmd_buffer.c -index 3da978b4837..338c33a53f5 100644 ---- a/src/panfrost/vulkan/panvk_cmd_buffer.c -+++ b/src/panfrost/vulkan/panvk_cmd_buffer.c -@@ -674,7 +674,8 @@ panvk_cmd_prepare_clear_values(struct panvk_cmd_buffer *cmdbuf, - cmdbuf->state.clear[i].stencil = in[i].depthStencil.stencil; - } - } else if (attachment->load_op == VK_ATTACHMENT_LOAD_OP_CLEAR) { -- panvk_pack_color(&cmdbuf->state.clear[i], &in[i].color, fmt); -+ union pipe_color_union *col = (union pipe_color_union *) &in[i].color; -+ pan_pack_color(cmdbuf->state.clear[i].color, col, fmt, false); - } - } - } -diff --git a/src/util/fast_idiv_by_const.c b/src/util/fast_idiv_by_const.c -index 4f0f6b769b8..b9f0b9cb760 100644 ---- a/src/util/fast_idiv_by_const.c -+++ b/src/util/fast_idiv_by_const.c -@@ -39,6 +39,7 @@ - - #include "fast_idiv_by_const.h" - #include "u_math.h" -+#include "util/macros.h" - #include - #include - -@@ -65,8 +66,7 @@ util_compute_fast_udiv_info(uint64_t D, unsigned num_bits, unsigned UINT_BITS) - } else { - /* Dividing by 1. */ - /* Assuming: floor((num + 1) * (2^32 - 1) / 2^32) = num */ -- result.multiplier = UINT_BITS == 64 ? UINT64_MAX : -- (1ull << UINT_BITS) - 1; -+ result.multiplier = u_uintN_max(UINT_BITS); - result.pre_shift = 0; - result.post_shift = 0; - result.increment = 1; -diff --git a/src/util/format/format_utils.h b/src/util/format/format_utils.h -index fa1d30060d9..0768a26ecce 100644 ---- a/src/util/format/format_utils.h -+++ b/src/util/format/format_utils.h -@@ -34,29 +34,24 @@ - #include "util/half_float.h" - #include "util/rounding.h" - --/* Only guaranteed to work for BITS <= 32 */ --#define MAX_UINT(BITS) ((BITS) == 32 ? UINT32_MAX : ((1u << (BITS)) - 1)) --#define MAX_INT(BITS) ((int)MAX_UINT((BITS) - 1)) --#define MIN_INT(BITS) ((BITS) == 32 ? INT32_MIN : (-(1 << (BITS - 1)))) -- - /* Extends an integer of size SRC_BITS to one of size DST_BITS linearly */ - #define EXTEND_NORMALIZED_INT(X, SRC_BITS, DST_BITS) \ -- (((X) * (int)(MAX_UINT(DST_BITS) / MAX_UINT(SRC_BITS))) + \ -+ (((X) * (int)(u_uintN_max(DST_BITS) / u_uintN_max(SRC_BITS))) + \ - ((DST_BITS % SRC_BITS) ? ((X) >> (SRC_BITS - DST_BITS % SRC_BITS)) : 0)) - - static inline float - _mesa_unorm_to_float(unsigned x, unsigned src_bits) - { -- return x * (1.0f / (float)MAX_UINT(src_bits)); -+ return x * (1.0f / (float)u_uintN_max(src_bits)); - } - - static inline float - _mesa_snorm_to_float(int x, unsigned src_bits) - { -- if (x <= -MAX_INT(src_bits)) -+ if (x <= -u_intN_max(src_bits)) - return -1.0f; - else -- return x * (1.0f / (float)MAX_INT(src_bits)); -+ return x * (1.0f / (float)u_intN_max(src_bits)); - } - - static inline uint16_t -@@ -77,9 +72,9 @@ _mesa_float_to_unorm(float x, unsigned dst_bits) - if (x < 0.0f) - return 0; - else if (x > 1.0f) -- return MAX_UINT(dst_bits); -+ return u_uintN_max(dst_bits); - else -- return _mesa_i64roundevenf(x * MAX_UINT(dst_bits)); -+ return _mesa_i64roundevenf(x * u_uintN_max(dst_bits)); - } - - static inline unsigned -@@ -98,10 +93,10 @@ _mesa_unorm_to_unorm(unsigned x, unsigned src_bits, unsigned dst_bits) - - if (src_bits + dst_bits > sizeof(x) * 8) { - assert(src_bits + dst_bits <= sizeof(uint64_t) * 8); -- return (((uint64_t) x * MAX_UINT(dst_bits) + src_half) / -- MAX_UINT(src_bits)); -+ return (((uint64_t) x * u_uintN_max(dst_bits) + src_half) / -+ u_uintN_max(src_bits)); - } else { -- return (x * MAX_UINT(dst_bits) + src_half) / MAX_UINT(src_bits); -+ return (x * u_uintN_max(dst_bits) + src_half) / u_uintN_max(src_bits); - } - } else { - return x; -@@ -121,11 +116,11 @@ static inline int - _mesa_float_to_snorm(float x, unsigned dst_bits) - { - if (x < -1.0f) -- return -MAX_INT(dst_bits); -+ return -u_intN_max(dst_bits); - else if (x > 1.0f) -- return MAX_INT(dst_bits); -+ return u_intN_max(dst_bits); - else -- return _mesa_lroundevenf(x * MAX_INT(dst_bits)); -+ return _mesa_lroundevenf(x * u_intN_max(dst_bits)); - } - - static inline int -@@ -143,8 +138,8 @@ _mesa_unorm_to_snorm(unsigned x, unsigned src_bits, unsigned dst_bits) - static inline int - _mesa_snorm_to_snorm(int x, unsigned src_bits, unsigned dst_bits) - { -- if (x < -MAX_INT(src_bits)) -- return -MAX_INT(dst_bits); -+ if (x < -u_intN_max(src_bits)) -+ return -u_intN_max(dst_bits); - else if (src_bits < dst_bits) - return EXTEND_NORMALIZED_INT(x, src_bits - 1, dst_bits - 1); - else -@@ -154,25 +149,25 @@ _mesa_snorm_to_snorm(int x, unsigned src_bits, unsigned dst_bits) - static inline unsigned - _mesa_unsigned_to_unsigned(unsigned src, unsigned dst_size) - { -- return MIN2(src, MAX_UINT(dst_size)); -+ return MIN2(src, u_uintN_max(dst_size)); - } - - static inline int - _mesa_unsigned_to_signed(unsigned src, unsigned dst_size) - { -- return MIN2(src, (unsigned)MAX_INT(dst_size)); -+ return MIN2(src, (unsigned)u_intN_max(dst_size)); - } - - static inline int - _mesa_signed_to_signed(int src, unsigned dst_size) - { -- return CLAMP(src, MIN_INT(dst_size), MAX_INT(dst_size)); -+ return CLAMP(src, u_intN_min(dst_size), u_intN_max(dst_size)); - } - - static inline unsigned - _mesa_signed_to_unsigned(int src, unsigned dst_size) - { -- return CLAMP(src, 0, MAX_UINT(dst_size)); -+ return CLAMP(src, 0, u_uintN_max(dst_size)); - } - - static inline unsigned -@@ -180,18 +175,18 @@ _mesa_float_to_unsigned(float src, unsigned dst_bits) - { - if (src < 0.0f) - return 0; -- if (src > (float)MAX_UINT(dst_bits)) -- return MAX_UINT(dst_bits); -+ if (src > (float)u_uintN_max(dst_bits)) -+ return u_uintN_max(dst_bits); - return _mesa_signed_to_unsigned(src, dst_bits); - } - - static inline unsigned - _mesa_float_to_signed(float src, unsigned dst_bits) - { -- if (src < (float)(-MAX_INT(dst_bits))) -- return -MAX_INT(dst_bits); -- if (src > (float)MAX_INT(dst_bits)) -- return MAX_INT(dst_bits); -+ if (src < (float)(-u_intN_max(dst_bits))) -+ return -u_intN_max(dst_bits); -+ if (src > (float)u_intN_max(dst_bits)) -+ return u_intN_max(dst_bits); - return _mesa_signed_to_signed(src, dst_bits); - } - -diff --git a/src/util/format/u_format.c b/src/util/format/u_format.c -index c49b3788c82..31f1f240efc 100644 ---- a/src/util/format/u_format.c -+++ b/src/util/format/u_format.c -@@ -1138,7 +1138,7 @@ static void - util_format_unpack_table_init(void) - { - for (enum pipe_format format = PIPE_FORMAT_NONE; format < PIPE_FORMAT_COUNT; format++) { --#if (defined(PIPE_ARCH_AARCH64) || defined(PIPE_ARCH_ARM)) && !defined NO_FORMAT_ASM -+#if (defined(PIPE_ARCH_AARCH64) || defined(PIPE_ARCH_ARM)) && !defined(NO_FORMAT_ASM) && !defined(__SOFTFP__) - const struct util_format_unpack_description *unpack = util_format_unpack_description_neon(format); - if (unpack) { - util_format_unpack_table[format] = unpack; -diff --git a/src/util/format/u_format_unpack_neon.c b/src/util/format/u_format_unpack_neon.c -index 7456d7aaa88..a4a5cb1f723 100644 ---- a/src/util/format/u_format_unpack_neon.c -+++ b/src/util/format/u_format_unpack_neon.c -@@ -23,7 +23,7 @@ - - #include - --#if (defined(PIPE_ARCH_AARCH64) || defined(PIPE_ARCH_ARM)) && !defined NO_FORMAT_ASM -+#if (defined(PIPE_ARCH_AARCH64) || defined(PIPE_ARCH_ARM)) && !defined(NO_FORMAT_ASM) && !defined(__SOFTFP__) - - /* armhf builds default to vfp, not neon, and refuses to compile neon intrinsics - * unless you tell it "no really". -diff --git a/src/util/fossilize_db.c b/src/util/fossilize_db.c -index e1709a1ff64..26024101b83 100644 ---- a/src/util/fossilize_db.c -+++ b/src/util/fossilize_db.c -@@ -156,18 +156,18 @@ update_foz_index(struct foz_db *foz_db, FILE *db_idx, unsigned file_idx) - offset += header->payload_size; - parsed_offset = offset; - -- /* Truncate the entry's hash string to a 64bit hash for use with a -- * 64bit hash table for looking up file offsets. -- */ -- hash_str[16] = '\0'; -- uint64_t key = strtoull(hash_str, NULL, 16); -- - struct foz_db_entry *entry = ralloc(foz_db->mem_ctx, - struct foz_db_entry); - entry->header = *header; - entry->file_idx = file_idx; - _mesa_sha1_hex_to_sha1(entry->key, hash_str); - -+ /* Truncate the entry's hash string to a 64bit hash for use with a -+ * 64bit hash table for looking up file offsets. -+ */ -+ hash_str[16] = '\0'; -+ uint64_t key = strtoull(hash_str, NULL, 16); -+ - entry->offset = cache_offset; - - _mesa_hash_table_u64_insert(foz_db->index_db, key, entry); -diff --git a/src/util/macros.h b/src/util/macros.h -index 1fc9e23355b..4bd18f55ec0 100644 ---- a/src/util/macros.h -+++ b/src/util/macros.h -@@ -30,6 +30,8 @@ - #include "c99_compat.h" - #include "c11_compat.h" - -+#include -+ - /* Compute the size of an array */ - #ifndef ARRAY_SIZE - # define ARRAY_SIZE(x) (sizeof(x) / sizeof((x)[0])) -@@ -392,6 +394,30 @@ do { \ - #define BITFIELD64_RANGE(b, count) \ - (BITFIELD64_MASK((b) + (count)) & ~BITFIELD64_MASK(b)) - -+static inline int64_t -+u_intN_max(unsigned bit_size) -+{ -+ assert(bit_size <= 64 && bit_size > 0); -+ return INT64_MAX >> (64 - bit_size); -+} -+ -+static inline int64_t -+u_intN_min(unsigned bit_size) -+{ -+ /* On 2's compliment platforms, which is every platform Mesa is likely to -+ * every worry about, stdint.h generally calculated INT##_MIN in this -+ * manner. -+ */ -+ return (-u_intN_max(bit_size)) - 1; -+} -+ -+static inline uint64_t -+u_uintN_max(unsigned bit_size) -+{ -+ assert(bit_size <= 64 && bit_size > 0); -+ return UINT64_MAX >> (64 - bit_size); -+} -+ - /* TODO: In future we should try to move this to u_debug.h once header - * dependencies are reorganised to allow this. - */ -diff --git a/src/util/meson.build b/src/util/meson.build -index aa5bfef5dbc..319b22d9bf7 100644 ---- a/src/util/meson.build -+++ b/src/util/meson.build -@@ -383,6 +383,15 @@ if with_tests - env: ['BUILD_FULL_PATH='+process_test_exe_full_path] - ) - -+ test('int_min_max', -+ executable('int_min_max_test', -+ files('tests/int_min_max.cpp'), -+ include_directories : [inc_include, inc_src], -+ dependencies : [idep_mesautil, idep_gtest], -+ ), -+ suite : ['util'], -+ ) -+ - subdir('tests/cache') - subdir('tests/fast_idiv_by_const') - subdir('tests/fast_urem_by_const') -diff --git a/src/util/tests/fast_idiv_by_const/fast_idiv_by_const_test.cpp b/src/util/tests/fast_idiv_by_const/fast_idiv_by_const_test.cpp -index 330f90fa464..abf6079944f 100644 ---- a/src/util/tests/fast_idiv_by_const/fast_idiv_by_const_test.cpp -+++ b/src/util/tests/fast_idiv_by_const/fast_idiv_by_const_test.cpp -@@ -30,9 +30,6 @@ - - #define RAND_TEST_ITERATIONS 100000 - --#define MAX_UINT(bits) \ -- (((bits) == 64) ? UINT64_MAX : ((1ull << (bits)) - 1)) -- - static inline uint64_t - utrunc(uint64_t x, unsigned num_bits) - { -@@ -82,7 +79,7 @@ uadd_sat(uint64_t a, uint64_t b, unsigned num_bits) - return sum < a ? UINT64_MAX : sum; - } else { - /* Check if sum is more than num_bits */ -- return (sum >> num_bits) ? MAX_UINT(num_bits) : sum; -+ return (sum >> num_bits) ? u_uintN_max(num_bits) : sum; - } - } - -@@ -201,7 +198,7 @@ rand_uint(unsigned bits, unsigned min) - if (k == 17) { - return min + (rand() % 16); - } else if (k == 42) { -- return MAX_UINT(bits) - (rand() % 16); -+ return u_uintN_max(bits) - (rand() % 16); - } else if (k == 9) { - uint64_t r; - do { -@@ -230,7 +227,7 @@ rand_sint(unsigned bits, unsigned min_abs) - { - /* Make sure we hit MIN_INT every once in a while */ - if (rand() % 64 == 37) -- return INT64_MIN >> (64 - bits); -+ return u_intN_min(bits); - - int64_t s = rand_uint(bits - 1, min_abs); - return rand() & 1 ? s : -s; -diff --git a/src/util/tests/int_min_max.cpp b/src/util/tests/int_min_max.cpp -new file mode 100644 -index 00000000000..8d74ecb7d33 ---- /dev/null -+++ b/src/util/tests/int_min_max.cpp -@@ -0,0 +1,73 @@ -+/* -+ * Copyright © 2021 Intel Corporation -+ * -+ * Permission is hereby granted, free of charge, to any person obtaining a -+ * copy of this software and associated documentation files (the "Software"), -+ * to deal in the Software without restriction, including without limitation -+ * the rights to use, copy, modify, merge, publish, distribute, sublicense, -+ * and/or sell copies of the Software, and to permit persons to whom the -+ * Software is furnished to do so, subject to the following conditions: -+ * -+ * The above copyright notice and this permission notice (including the next -+ * paragraph) shall be included in all copies or substantial portions of the -+ * Software. -+ * -+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL -+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING -+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS -+ * IN THE SOFTWARE. -+ */ -+ -+#include -+#include "util/macros.h" -+ -+#define MESA_UINT24_MAX 16777215 -+#define MESA_INT24_MAX 8388607 -+#define MESA_INT24_MIN (-8388607-1) -+ -+#define MESA_UINT12_MAX 4095 -+#define MESA_INT12_MAX 2047 -+#define MESA_INT12_MIN (-2047-1) -+ -+#define MESA_UINT10_MAX 1023 -+#define MESA_INT10_MAX 511 -+#define MESA_INT10_MIN (-511-1) -+ -+TEST(int_min_max, u_intN_min) -+{ -+ EXPECT_EQ(INT64_MIN, u_intN_min(64)); -+ EXPECT_EQ(INT32_MIN, u_intN_min(32)); -+ EXPECT_EQ(INT16_MIN, u_intN_min(16)); -+ EXPECT_EQ(INT8_MIN, u_intN_min(8)); -+ -+ EXPECT_EQ(MESA_INT24_MIN, u_intN_min(24)); -+ EXPECT_EQ(MESA_INT12_MIN, u_intN_min(12)); -+ EXPECT_EQ(MESA_INT10_MIN, u_intN_min(10)); -+} -+ -+TEST(int_min_max, u_intN_max) -+{ -+ EXPECT_EQ(INT64_MAX, u_intN_max(64)); -+ EXPECT_EQ(INT32_MAX, u_intN_max(32)); -+ EXPECT_EQ(INT16_MAX, u_intN_max(16)); -+ EXPECT_EQ(INT8_MAX, u_intN_max(8)); -+ -+ EXPECT_EQ(MESA_INT24_MAX, u_intN_max(24)); -+ EXPECT_EQ(MESA_INT12_MAX, u_intN_max(12)); -+ EXPECT_EQ(MESA_INT10_MAX, u_intN_max(10)); -+} -+ -+TEST(int_min_max, u_uintN_max) -+{ -+ EXPECT_EQ(UINT64_MAX, u_uintN_max(64)); -+ EXPECT_EQ(UINT32_MAX, u_uintN_max(32)); -+ EXPECT_EQ(UINT16_MAX, u_uintN_max(16)); -+ EXPECT_EQ(UINT8_MAX, u_uintN_max(8)); -+ -+ EXPECT_EQ(MESA_UINT24_MAX, u_uintN_max(24)); -+ EXPECT_EQ(MESA_UINT12_MAX, u_uintN_max(12)); -+ EXPECT_EQ(MESA_UINT10_MAX, u_uintN_max(10)); -+} -diff --git a/src/vulkan/wsi/wsi_common.c b/src/vulkan/wsi/wsi_common.c -index b1360edb911..292bb976da8 100644 ---- a/src/vulkan/wsi/wsi_common.c -+++ b/src/vulkan/wsi/wsi_common.c -@@ -653,6 +653,10 @@ wsi_common_queue_present(const struct wsi_device *wsi, - if (result != VK_SUCCESS) - goto fail_present; - -+ if (wsi->sw) -+ wsi->WaitForFences(device, 1, &swapchain->fences[image_index], -+ true, ~0ull); -+ - const VkPresentRegionKHR *region = NULL; - if (regions && regions->pRegions) - region = ®ions->pRegions[i]; diff --git a/mesa.spec b/mesa.spec index 64a1c5a..2d05d1f 100644 --- a/mesa.spec +++ b/mesa.spec @@ -57,7 +57,7 @@ Name: mesa Summary: Mesa graphics libraries -%global ver 21.2.1 +%global ver 21.2.2 Version: %{lua:ver = string.gsub(rpm.expand("%{ver}"), "-", "~"); print(ver)} Release: %autorelease License: MIT @@ -68,7 +68,6 @@ Source0: https://mesa.freedesktop.org/archive/%{name}-%{ver}.tar.xz # Source1 contains email correspondence clarifying the license terms. # Fedora opts to ignore the optional part of clause 2 and treat that code as 2 clause BSD. Source1: Mesa-MLAA-License-Clarification-Email.txt -Patch0: 21.2-fixes.patch # Backport of upstream patches from # https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11940 diff --git a/sources b/sources index 63a6983..8fe2b06 100644 --- a/sources +++ b/sources @@ -1 +1 @@ -SHA512 (mesa-21.2.1.tar.xz) = d4056287ec86f7a95ce534a251a1ccbc3a3b08a2f7112152def2f054fc8a9424501d5883c463554ee95fe2dafb832613efd7145e989ee8281948233942730c2c +SHA512 (mesa-21.2.2.tar.xz) = 0a4877b405384088c8bdac3031444cd22377d19552c41c08fb8928d6edf23d5f5a237ef1d6c5f96e3293c2e90c63f4702813226b85cb304fe5e01dd2710ba697