+ } else if (instr->format == Format::PSEUDO_BARRIER) {
+ Pseudo_barrier_instruction* barrier = static_cast<Pseudo_barrier_instruction*>(instr.get());
+
+ /* Anything larger than a workgroup isn't possible. Anything
+ * smaller requires no instructions and this pseudo instruction
+ * would only be included to control optimizations. */
+ bool emit_s_barrier = barrier->exec_scope == scope_workgroup &&
+ program->workgroup_size > program->wave_size;
+
+ bld.insert(std::move(instr));
+ if (emit_s_barrier)
+ bld.sopp(aco_opcode::s_barrier);
+ } else if (instr->opcode == aco_opcode::p_cvt_f16_f32_rtne) {
+ float_mode new_mode = block->fp_mode;
+ new_mode.round16_64 = fp_round_ne;
+ bool set_round = new_mode.round != block->fp_mode.round;
+
+ emit_set_mode(bld, new_mode, set_round, false);
+
+ instr->opcode = aco_opcode::v_cvt_f16_f32;
+ ctx.instructions.emplace_back(std::move(instr));
+
+ emit_set_mode(bld, block->fp_mode, set_round, false);