namespace aco {
namespace {
+#define isel_err(...) _isel_err(ctx, __FILE__, __LINE__, __VA_ARGS__)
+
+static void _isel_err(isel_context *ctx, const char *file, unsigned line,
+ const nir_instr *instr, const char *msg)
+{
+ char *out;
+ size_t outsize;
+ FILE *memf = open_memstream(&out, &outsize);
+
+ fprintf(memf, "%s: ", msg);
+ nir_print_instr(instr, memf);
+ fclose(memf);
+
+ _aco_err(ctx->program, file, line, out);
+ free(out);
+}
+
class loop_info_RAII {
isel_context* ctx;
unsigned header_idx_old;
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), dst0, dst1);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
return;
}
aco_opcode op = dst.regClass() == s1 ? aco_opcode::s_cselect_b32 : aco_opcode::s_cselect_b64;
bld.sop2(op, Definition(dst), then, els, bld.scc(bool_to_scalar_condition(ctx, cond)));
} else {
- fprintf(stderr, "Unimplemented uniform bcsel bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented uniform bcsel bit size");
}
return;
}
void visit_alu_instr(isel_context *ctx, nir_alu_instr *instr)
{
if (!instr->dest.dest.is_ssa) {
- fprintf(stderr, "nir alu dst not in ssa: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "nir alu dst not in ssa");
abort();
}
Builder bld(ctx->program, ctx->block);
aco_opcode opcode = dst.size() == 1 ? aco_opcode::s_not_b32 : aco_opcode::s_not_b64;
bld.sop1(opcode, Definition(dst), bld.def(s1, scc), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
}
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
Temp src = get_alu_src(ctx, instr->src[0]);
bld.vop2(aco_opcode::v_max_i32, Definition(dst), src, bld.vsub32(bld.def(v1), Operand(0u), src));
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
upper = bld.vop2(aco_opcode::v_cndmask_b32, bld.def(v1), Operand(0u), neg, gtz);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s1) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_max_i32, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s1) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_max_u32, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s1) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_min_i32, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s1) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_min_u32, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s2) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_or_b64, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s2) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_and_b64, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s2) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_xor_b64, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s1) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_lshr_b32, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s2) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_lshl_b64, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s2) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_ashr_i64, dst, true);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (src.regClass() == s2) {
bld.sop1(aco_opcode::s_ff1_i32_b64, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
Temp carry = bld.vsub32(Definition(msb), Operand(31u), Operand(msb_rev), true).def(1).getTemp();
bld.vop2_e64(aco_opcode::v_cndmask_b32, Definition(dst), msb, Operand((uint32_t)-1), carry);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
bld.vop1(aco_opcode::v_bfrev_b32, Definition(dst), get_alu_src(ctx, instr->src[0]));
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
Temp dst1 = bld.vadd32(bld.def(v1), src01, src11, false, carry);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), dst0, dst1);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop2_e64(aco_opcode::v_cndmask_b32, Definition(dst), tmp, Operand((uint32_t) -1), carry);
}
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
carry = bld.vop2_e64(aco_opcode::v_cndmask_b32, bld.def(v1), Operand(0u), Operand(1u), carry);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), carry, Operand(0u));
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
Temp upper = bld.vsub32(bld.def(v1), src01, src11, false, borrow);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
borrow = bld.vop2_e64(aco_opcode::v_cndmask_b32, bld.def(v1), Operand(0u), Operand(1u), borrow);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), borrow, Operand(0u));
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == s1) {
emit_sop2_instruction(ctx, instr, aco_opcode::s_mul_i32, dst, false);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
as_vgpr(ctx, get_alu_src(ctx, instr->src[1])));
bld.pseudo(aco_opcode::p_as_uniform, Definition(dst), tmp);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
as_vgpr(ctx, get_alu_src(ctx, instr->src[1])));
bld.pseudo(aco_opcode::p_as_uniform, Definition(dst), tmp);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
bld.vop3(aco_opcode::v_mul_f64, Definition(dst), src0, src1);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
bld.vop3(aco_opcode::v_add_f64, Definition(dst), src0, src1);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
VOP3A_instruction* sub = static_cast<VOP3A_instruction*>(add);
sub->neg[1] = true;
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop3(aco_opcode::v_max_f64, Definition(dst), src0, src1);
}
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop3(aco_opcode::v_min_f64, Definition(dst), src0, src1);
}
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_max3_f32, dst, ctx->block->fp_mode.must_flush_denorms32);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_min3_f32, dst, ctx->block->fp_mode.must_flush_denorms32);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_med3_f32, dst, ctx->block->fp_mode.must_flush_denorms32);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.size() == 1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_max3_u32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.size() == 1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_min3_u32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.size() == 1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_med3_u32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.size() == 1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_max3_i32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.size() == 1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_min3_i32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.size() == 1) {
emit_vop3a_instruction(ctx, instr, aco_opcode::v_med3_i32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
/* Lowered at NIR level for precision reasons. */
emit_vop1_instruction(ctx, instr, aco_opcode::v_rsq_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
upper = bld.vop2(aco_opcode::v_xor_b32, bld.def(v1), Operand(0x80000000u), upper);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
upper = bld.vop2(aco_opcode::v_and_b32, bld.def(v1), Operand(0x7FFFFFFFu), upper);
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
VOP3A_instruction* vop3 = static_cast<VOP3A_instruction*>(add);
vop3->clamp = true;
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
emit_log2(ctx, bld, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
/* Lowered at NIR level for precision reasons. */
emit_vop1_instruction(ctx, instr, aco_opcode::v_rcp_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
emit_vop1_instruction(ctx, instr, aco_opcode::v_exp_f32, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
/* Lowered at NIR level for precision reasons. */
emit_vop1_instruction(ctx, instr, aco_opcode::v_sqrt_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
emit_vop1_instruction(ctx, instr, aco_opcode::v_fract_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
emit_floor_f64(ctx, bld, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop3(aco_opcode::v_add_f64, Definition(dst), trunc, add);
}
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
emit_trunc_f64(ctx, bld, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), dst0, dst1);
}
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
aco_opcode opcode = instr->op == nir_op_fsin ? aco_opcode::v_sin_f32 : aco_opcode::v_cos_f32;
bld.vop1(opcode, Definition(dst), tmp);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
bld.vop3(aco_opcode::v_ldexp_f64, Definition(dst), as_vgpr(ctx, src0), src1);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v2) {
bld.vop1(aco_opcode::v_frexp_mant_f64, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (instr->src[0].src.ssa->bit_size == 64) {
bld.vop1(aco_opcode::v_frexp_exp_i32_f64, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), Operand(0u), upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (instr->src[0].src.ssa->bit_size == 64) {
emit_vop1_instruction(ctx, instr, aco_opcode::v_cvt_f32_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop3(aco_opcode::v_add_f64, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
upper = bld.vop3(aco_opcode::v_ldexp_f64, bld.def(v2), upper, Operand(32u));
bld.vop3(aco_opcode::v_add_f64, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (instr->src[0].src.ssa->bit_size == 64) {
emit_vop1_instruction(ctx, instr, aco_opcode::v_cvt_i32_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (instr->src[0].src.ssa->bit_size == 64) {
emit_vop1_instruction(ctx, instr, aco_opcode::v_cvt_u32_f64, dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lower, upper);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
Temp src0 = bld.tmp(v1);
Temp src1 = bld.tmp(v1);
bld.pseudo(aco_opcode::p_split_vector, Definition(src0), Definition(src1), src);
- if (!ctx->block->fp_mode.care_about_round32 || ctx->block->fp_mode.round32 == fp_round_tz)
+ if (0 && (!ctx->block->fp_mode.care_about_round32 || ctx->block->fp_mode.round32 == fp_round_tz)) {
bld.vop3(aco_opcode::v_cvt_pkrtz_f16_f32, Definition(dst), src0, src1);
- else
- bld.vop3(aco_opcode::v_cvt_pk_u16_u32, Definition(dst),
- bld.vop1(aco_opcode::v_cvt_f32_f16, bld.def(v1), src0),
- bld.vop1(aco_opcode::v_cvt_f32_f16, bld.def(v1), src1));
+ } else {
+ src0 = bld.vop1(aco_opcode::v_cvt_f16_f32, bld.def(v1), src0);
+ src1 = bld.vop1(aco_opcode::v_cvt_f16_f32, bld.def(v1), src1);
+ if (ctx->program->chip_class >= GFX10) {
+ /* the high bits of v_cvt_f16_f32 isn't zero'd on GFX10 */
+ bld.vop3(aco_opcode::v_pack_b32_f16, Definition(dst), src0, src1);
+ } else {
+ bld.vop3(aco_opcode::v_cvt_pk_u16_u32, Definition(dst), src0, src1);
+ }
+ }
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
if (dst.regClass() == v1) {
bld.vop1(aco_opcode::v_cvt_f32_f16, Definition(dst), get_alu_src(ctx, instr->src[0]));
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop1(aco_opcode::v_cvt_f32_f16, Definition(dst),
bld.vop2(aco_opcode::v_lshrrev_b32, bld.def(v1), Operand(16u), as_vgpr(ctx, get_alu_src(ctx, instr->src[0]))));
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (dst.regClass() == v1) {
bld.vop3(aco_opcode::v_bfm_b32, Definition(dst), bits, offset);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.vop3(aco_opcode::v_bfi_b32, Definition(dst), bitmask, insert, base);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
} else if (src.regClass() == s2) {
bld.sop1(aco_opcode::s_bcnt1_i32_b64, Definition(dst), bld.def(s1, scc), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
emit_comparison(ctx, instr, dst, aco_opcode::v_cmp_eq_f16, aco_opcode::v_cmp_eq_f32, aco_opcode::v_cmp_eq_f64);
break;
}
- case nir_op_fne: {
+ case nir_op_fneu: {
emit_comparison(ctx, instr, dst, aco_opcode::v_cmp_neq_f16, aco_opcode::v_cmp_neq_f32, aco_opcode::v_cmp_neq_f64);
break;
}
break;
}
default:
- fprintf(stderr, "Unknown NIR ALU instr: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unknown NIR ALU instr");
}
}
ctx->shader->info.stage == MESA_SHADER_GEOMETRY) {
bool stored_to_temps = store_output_to_temps(ctx, instr);
if (!stored_to_temps) {
- fprintf(stderr, "Unimplemented output offset instruction:\n");
- nir_print_instr(instr->src[1].ssa->parent_instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(instr->src[1].ssa->parent_instr, "Unimplemented output offset instruction");
abort();
}
} else if (ctx->stage == vertex_es ||
nir_instr *off_instr = instr->src[0].ssa->parent_instr;
if (off_instr->type != nir_instr_type_load_const) {
- fprintf(stderr, "Unimplemented nir_intrinsic_load_input offset\n");
- nir_print_instr(off_instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(off_instr, "Unimplemented nir_intrinsic_load_input offset");
}
uint32_t offset = nir_instr_as_load_const(off_instr)->value[0].u32;
nir_instr *off_instr = instr->src[offset_idx].ssa->parent_instr;
if (off_instr->type != nir_instr_type_load_const ||
nir_instr_as_load_const(off_instr)->value[0].u32 != 0) {
- fprintf(stderr, "Unimplemented nir_intrinsic_load_input offset\n");
- nir_print_instr(off_instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(off_instr, "Unimplemented nir_intrinsic_load_input offset");
}
Temp prim_mask = get_arg(ctx, ctx->args->ac.prim_mask);
}
/* LOD */
+ assert(nir_src_as_uint(instr->src[1]) == 0);
Temp lod = bld.vop1(aco_opcode::v_mov_b32, bld.def(v1), Operand(0u));
/* Resource */
} else if (src.regClass() == s2) {
bld.sop1(aco_opcode::s_mov_b64, dst, src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
}
} else if (instr->src[0].ssa->bit_size == 64 && src.regClass() == v2) {
bld.vopc(aco_opcode::v_cmp_lg_u64, lanemask_tmp, Operand(0u), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
if (dst.size() != bld.lm.size()) {
/* Wave32 with ballot size set to 64 */
tmp = bld.vop2(aco_opcode::v_and_b32, bld.def(v1), Operand(1u), tmp);
emit_wqm(ctx, bld.vopc(aco_opcode::v_cmp_lg_u32, bld.def(bld.lm), Operand(0u), tmp), dst);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
}
break;
} else if (src.regClass() == s2) {
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), src);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lo, hi);
emit_split_vector(ctx, dst, 2);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
}
break;
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lo, hi);
emit_split_vector(ctx, dst, 2);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lo, hi);
emit_split_vector(ctx, dst, 2);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
bld.pseudo(aco_opcode::p_create_vector, Definition(dst), lo, hi);
emit_split_vector(ctx, dst, 2);
} else {
- fprintf(stderr, "Unimplemented NIR instr bit size: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented NIR instr bit size");
}
break;
}
break;
}
default:
- fprintf(stderr, "Unimplemented intrinsic instr: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unimplemented intrinsic instr");
abort();
break;
}
break;
default:
- fprintf(stderr, "Unknown NIR jump instr: ");
- nir_print_instr(&instr->instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(&instr->instr, "Unknown NIR jump instr");
abort();
}
visit_jump(ctx, nir_instr_as_jump(instr));
break;
default:
- fprintf(stderr, "Unknown NIR instr type: ");
- nir_print_instr(instr, stderr);
- fprintf(stderr, "\n");
+ isel_err(instr, "Unknown NIR instr type");
//abort();
}
}
cleanup_cfg(program);
}
+
+void select_trap_handler_shader(Program *program, struct nir_shader *shader,
+ ac_shader_config* config,
+ struct radv_shader_args *args)
+{
+ assert(args->options->chip_class == GFX8);
+
+ init_program(program, compute_cs, args->shader_info,
+ args->options->chip_class, args->options->family, config);
+
+ isel_context ctx = {};
+ ctx.program = program;
+ ctx.args = args;
+ ctx.options = args->options;
+ ctx.stage = program->stage;
+
+ ctx.block = ctx.program->create_and_insert_block();
+ ctx.block->loop_nest_depth = 0;
+ ctx.block->kind = block_kind_top_level;
+
+ program->workgroup_size = 1; /* XXX */
+
+ add_startpgm(&ctx);
+ append_logical_start(ctx.block);
+
+ Builder bld(ctx.program, ctx.block);
+
+ /* Load the buffer descriptor from TMA. */
+ bld.smem(aco_opcode::s_load_dwordx4, Definition(PhysReg{ttmp4}, s4),
+ Operand(PhysReg{tma}, s2), Operand(0u));
+
+ /* Store TTMP0-TTMP1. */
+ bld.smem(aco_opcode::s_buffer_store_dwordx2, Operand(PhysReg{ttmp4}, s4),
+ Operand(0u), Operand(PhysReg{ttmp0}, s2), memory_sync_info(), true);
+
+ uint32_t hw_regs_idx[] = {
+ 2, /* HW_REG_STATUS */
+ 3, /* HW_REG_TRAP_STS */
+ 4, /* HW_REG_HW_ID */
+ 7, /* HW_REG_IB_STS */
+ };
+
+ /* Store some hardware registers. */
+ for (unsigned i = 0; i < ARRAY_SIZE(hw_regs_idx); i++) {
+ /* "((size - 1) << 11) | register" */
+ bld.sopk(aco_opcode::s_getreg_b32, Definition(PhysReg{ttmp8}, s1),
+ ((20 - 1) << 11) | hw_regs_idx[i]);
+
+ bld.smem(aco_opcode::s_buffer_store_dword, Operand(PhysReg{ttmp4}, s4),
+ Operand(8u + i * 4), Operand(PhysReg{ttmp8}, s1), memory_sync_info(), true);
+ }
+
+ program->config->float_mode = program->blocks[0].fp_mode.val;
+
+ append_logical_end(ctx.block);
+ ctx.block->kind |= block_kind_uniform;
+ bld.sopp(aco_opcode::s_endpgm);
+
+ cleanup_cfg(program);
+}
}