diff options
| -rw-r--r-- | src/amd/common/ac_nir_to_llvm.c | 518 | ||||
| -rw-r--r-- | src/amd/common/ac_nir_to_llvm.h | 6 |
2 files changed, 290 insertions, 234 deletions
diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c index b086402434..e21e17de95 100644 --- a/src/amd/common/ac_nir_to_llvm.c +++ b/src/amd/common/ac_nir_to_llvm.c @@ -54,11 +54,35 @@ enum desc_type { DESC_BUFFER, }; +struct nir_to_llvm_context; + +struct ac_nir_context { + struct ac_llvm_context ac; + struct ac_shader_abi *abi; + + gl_shader_stage stage; + + struct hash_table *defs; + struct hash_table *phis; + struct hash_table *vars; + + LLVMBasicBlockRef continue_block; + LLVMBasicBlockRef break_block; + + LLVMValueRef outputs[RADEON_LLVM_MAX_OUTPUTS * 4]; + + int num_locals; + LLVMValueRef *locals; + + struct nir_to_llvm_context *nctx; /* TODO get rid of this */ +}; + struct nir_to_llvm_context { struct ac_llvm_context ac; const struct ac_nir_compiler_options *options; struct ac_shader_variant_info *shader_info; struct ac_shader_abi abi; + struct ac_nir_context *nir; unsigned max_workgroup_size; LLVMContextRef context; @@ -68,7 +92,6 @@ struct nir_to_llvm_context { struct hash_table *defs; struct hash_table *phis; - struct hash_table *vars; LLVMValueRef descriptor_sets[AC_UD_MAX_SETS]; LLVMValueRef ring_offsets; @@ -118,9 +141,6 @@ struct nir_to_llvm_context { LLVMValueRef sample_coverage; LLVMValueRef frag_pos[4]; - LLVMBasicBlockRef continue_block; - LLVMBasicBlockRef break_block; - LLVMTypeRef i1; LLVMTypeRef i8; LLVMTypeRef i16; @@ -151,12 +171,9 @@ struct nir_to_llvm_context { LLVMValueRef lds; LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4]; - LLVMValueRef outputs[RADEON_LLVM_MAX_OUTPUTS * 4]; uint64_t input_mask; uint64_t output_mask; - int num_locals; - LLVMValueRef *locals; uint8_t num_output_clips; uint8_t num_output_culls; @@ -1011,32 +1028,32 @@ build_store_values_extended(struct nir_to_llvm_context *ctx, } } -static LLVMTypeRef get_def_type(struct nir_to_llvm_context *ctx, +static LLVMTypeRef get_def_type(struct ac_nir_context *ctx, const nir_ssa_def *def) { - LLVMTypeRef type = LLVMIntTypeInContext(ctx->context, def->bit_size); + LLVMTypeRef type = LLVMIntTypeInContext(ctx->ac.context, def->bit_size); if (def->num_components > 1) { type = LLVMVectorType(type, def->num_components); } return type; } -static LLVMValueRef get_src(struct nir_to_llvm_context *ctx, nir_src src) +static LLVMValueRef get_src(struct ac_nir_context *nir, nir_src src) { assert(src.is_ssa); - struct hash_entry *entry = _mesa_hash_table_search(ctx->defs, src.ssa); + struct hash_entry *entry = _mesa_hash_table_search(nir->defs, src.ssa); return (LLVMValueRef)entry->data; } -static LLVMBasicBlockRef get_block(struct nir_to_llvm_context *ctx, +static LLVMBasicBlockRef get_block(struct ac_nir_context *nir, const struct nir_block *b) { - struct hash_entry *entry = _mesa_hash_table_search(ctx->defs, b); + struct hash_entry *entry = _mesa_hash_table_search(nir->defs, b); return (LLVMBasicBlockRef)entry->data; } -static LLVMValueRef get_alu_src(struct nir_to_llvm_context *ctx, +static LLVMValueRef get_alu_src(struct ac_nir_context *ctx, nir_alu_src src, unsigned num_components) { @@ -1057,20 +1074,20 @@ static LLVMValueRef get_alu_src(struct nir_to_llvm_context *ctx, if (need_swizzle || num_components != src_components) { LLVMValueRef masks[] = { - LLVMConstInt(ctx->i32, src.swizzle[0], false), - LLVMConstInt(ctx->i32, src.swizzle[1], false), - LLVMConstInt(ctx->i32, src.swizzle[2], false), - LLVMConstInt(ctx->i32, src.swizzle[3], false)}; + LLVMConstInt(ctx->ac.i32, src.swizzle[0], false), + LLVMConstInt(ctx->ac.i32, src.swizzle[1], false), + LLVMConstInt(ctx->ac.i32, src.swizzle[2], false), + LLVMConstInt(ctx->ac.i32, src.swizzle[3], false)}; if (src_components > 1 && num_components == 1) { - value = LLVMBuildExtractElement(ctx->builder, value, + value = LLVMBuildExtractElement(ctx->ac.builder, value, masks[0], ""); } else if (src_components == 1 && num_components > 1) { LLVMValueRef values[] = {value, value, value, value}; value = ac_build_gather_values(&ctx->ac, values, num_components); } else { LLVMValueRef swizzle = LLVMConstVector(masks, num_components); - value = LLVMBuildShuffleVector(ctx->builder, value, value, + value = LLVMBuildShuffleVector(ctx->ac.builder, value, value, swizzle, ""); } } @@ -1488,7 +1505,7 @@ static LLVMValueRef emit_ddxy_interp( return ac_build_gather_values(&ctx->ac, result, 4); } -static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *instr) +static void visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr) { LLVMValueRef src[4], result = NULL; unsigned num_components = instr->dest.dest.ssa.num_components; @@ -1522,38 +1539,38 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst break; case nir_op_fneg: src[0] = to_float(&ctx->ac, src[0]); - result = LLVMBuildFNeg(ctx->builder, src[0], ""); + result = LLVMBuildFNeg(ctx->ac.builder, src[0], ""); break; case nir_op_ineg: - result = LLVMBuildNeg(ctx->builder, src[0], ""); + result = LLVMBuildNeg(ctx->ac.builder, src[0], ""); break; case nir_op_inot: - result = LLVMBuildNot(ctx->builder, src[0], ""); + result = LLVMBuildNot(ctx->ac.builder, src[0], ""); break; case nir_op_iadd: - result = LLVMBuildAdd(ctx->builder, src[0], src[1], ""); + result = LLVMBuildAdd(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_fadd: src[0] = to_float(&ctx->ac, src[0]); src[1] = to_float(&ctx->ac, src[1]); - result = LLVMBuildFAdd(ctx->builder, src[0], src[1], ""); + result = LLVMBuildFAdd(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_fsub: src[0] = to_float(&ctx->ac, src[0]); src[1] = to_float(&ctx->ac, src[1]); - result = LLVMBuildFSub(ctx->builder, src[0], src[1], ""); + result = LLVMBuildFSub(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_isub: - result = LLVMBuildSub(ctx->builder, src[0], src[1], ""); + result = LLVMBuildSub(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_imul: - result = LLVMBuildMul(ctx->builder, src[0], src[1], ""); + result = LLVMBuildMul(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_imod: - result = LLVMBuildSRem(ctx->builder, src[0], src[1], ""); + result = LLVMBuildSRem(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_umod: - result = LLVMBuildURem(ctx->builder, src[0], src[1], ""); + result = LLVMBuildURem(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_fmod: src[0] = to_float(&ctx->ac, src[0]); @@ -1561,27 +1578,27 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst result = ac_build_fdiv(&ctx->ac, src[0], src[1]); result = emit_intrin_1f_param(&ctx->ac, "llvm.floor", to_float_type(&ctx->ac, def_type), result); - result = LLVMBuildFMul(ctx->builder, src[1] , result, ""); - result = LLVMBuildFSub(ctx->builder, src[0], result, ""); + result = LLVMBuildFMul(ctx->ac.builder, src[1] , result, ""); + result = LLVMBuildFSub(ctx->ac.builder, src[0], result, ""); break; case nir_op_frem: src[0] = to_float(&ctx->ac, src[0]); src[1] = to_float(&ctx->ac, src[1]); - result = LLVMBuildFRem(ctx->builder, src[0], src[1], ""); + result = LLVMBuildFRem(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_irem: - result = LLVMBuildSRem(ctx->builder, src[0], src[1], ""); + result = LLVMBuildSRem(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_idiv: - result = LLVMBuildSDiv(ctx->builder, src[0], src[1], ""); + result = LLVMBuildSDiv(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_udiv: - result = LLVMBuildUDiv(ctx->builder, src[0], src[1], ""); + result = LLVMBuildUDiv(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_fmul: src[0] = to_float(&ctx->ac, src[0]); src[1] = to_float(&ctx->ac, src[1]); - result = LLVMBuildFMul(ctx->builder, src[0], src[1], ""); + result = LLVMBuildFMul(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_fdiv: src[0] = to_float(&ctx->ac, src[0]); @@ -1590,32 +1607,32 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst break; case nir_op_frcp: src[0] = to_float(&ctx->ac, src[0]); - result = ac_build_fdiv(&ctx->ac, ctx->f32one, src[0]); + result = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, src[0]); break; case nir_op_iand: - result = LLVMBuildAnd(ctx->builder, src[0], src[1], ""); + result = LLVMBuildAnd(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_ior: - result = LLVMBuildOr(ctx->builder, src[0], src[1], ""); + result = LLVMBuildOr(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_ixor: - result = LLVMBuildXor(ctx->builder, src[0], src[1], ""); + result = LLVMBuildXor(ctx->ac.builder, src[0], src[1], ""); break; case nir_op_ishl: - result = LLVMBuildShl(ctx->builder, src[0], - LLVMBuildZExt(ctx->builder, src[1], + result = LLVMBuildShl(ctx->ac.builder, src[0], + LLVMBuildZExt(ctx->ac.builder, src[1], LLVMTypeOf(src[0]), ""), ""); break; case nir_op_ishr: - result = LLVMBuildAShr(ctx->builder, src[0], - LLVMBuildZExt(ctx->builder, src[1], + result = LLVMBuildAShr(ctx->ac.builder, src[0], + LLVMBuildZExt(ctx->ac.builder, src[1], LLVMTypeOf(src[0]), ""), ""); break; case nir_op_ushr: - result = LLVMBuildLShr(ctx->builder, src[0], - LLVMBuildZExt(ctx->builder, src[1], + result = LLVMBuildLShr(ctx->ac.builder, src[0], + LLVMBuildZExt(ctx->ac.builder, src[1], LLVMTypeOf(src[0]), ""), ""); break; @@ -1717,7 +1734,7 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst case nir_op_frsq: result = emit_intrin_1f_param(&ctx->ac, "llvm.sqrt", to_float_type(&ctx->ac, def_type), src[0]); - result = ac_build_fdiv(&ctx->ac, ctx->f32one, result); + result = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, result); break; case nir_op_fpow: result = emit_intrin_2f_param(&ctx->ac, "llvm.pow", @@ -1753,10 +1770,10 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst result = emit_bitfield_insert(&ctx->ac, src[0], src[1], src[2], src[3]); break; case nir_op_bitfield_reverse: - result = ac_build_intrinsic(&ctx->ac, "llvm.bitreverse.i32", ctx->i32, src, 1, AC_FUNC_ATTR_READNONE); + result = ac_build_intrinsic(&ctx->ac, "llvm.bitreverse.i32", ctx->ac.i32, src, 1, AC_FUNC_ATTR_READNONE); break; case nir_op_bit_count: - result = ac_build_intrinsic(&ctx->ac, "llvm.ctpop.i32", ctx->i32, src, 1, AC_FUNC_ATTR_READNONE); + result = ac_build_intrinsic(&ctx->ac, "llvm.ctpop.i32", ctx->ac.i32, src, 1, AC_FUNC_ATTR_READNONE); break; case nir_op_vec2: case nir_op_vec3: @@ -1768,40 +1785,40 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst case nir_op_f2i32: case nir_op_f2i64: src[0] = to_float(&ctx->ac, src[0]); - result = LLVMBuildFPToSI(ctx->builder, src[0], def_type, ""); + result = LLVMBuildFPToSI(ctx->ac.builder, src[0], def_type, ""); break; case nir_op_f2u32: case nir_op_f2u64: src[0] = to_float(&ctx->ac, src[0]); - result = LLVMBuildFPToUI(ctx->builder, src[0], def_type, ""); + result = LLVMBuildFPToUI(ctx->ac.builder, src[0], def_type, ""); break; case nir_op_i2f32: case nir_op_i2f64: - result = LLVMBuildSIToFP(ctx->builder, src[0], to_float_type(&ctx->ac, def_type), ""); + result = LLVMBuildSIToFP(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), ""); break; case nir_op_u2f32: case nir_op_u2f64: - result = LLVMBuildUIToFP(ctx->builder, src[0], to_float_type(&ctx->ac, def_type), ""); + result = LLVMBuildUIToFP(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), ""); break; case nir_op_f2f64: - result = LLVMBuildFPExt(ctx->builder, src[0], to_float_type(&ctx->ac, def_type), ""); + result = LLVMBuildFPExt(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), ""); break; case nir_op_f2f32: - result = LLVMBuildFPTrunc(ctx->builder, src[0], to_float_type(&ctx->ac, def_type), ""); + result = LLVMBuildFPTrunc(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), ""); break; case nir_op_u2u32: case nir_op_u2u64: if (get_elem_bits(&ctx->ac, LLVMTypeOf(src[0])) < get_elem_bits(&ctx->ac, def_type)) - result = LLVMBuildZExt(ctx->builder, src[0], def_type, ""); + result = LLVMBuildZExt(ctx->ac.builder, src[0], def_type, ""); else - result = LLVMBuildTrunc(ctx->builder, src[0], def_type, ""); + result = LLVMBuildTrunc(ctx->ac.builder, src[0], def_type, ""); break; case nir_op_i2i32: case nir_op_i2i64: if (get_elem_bits(&ctx->ac, LLVMTypeOf(src[0])) < get_elem_bits(&ctx->ac, def_type)) - result = LLVMBuildSExt(ctx->builder, src[0], def_type, ""); + result = LLVMBuildSExt(ctx->ac.builder, src[0], def_type, ""); else - result = LLVMBuildTrunc(ctx->builder, src[0], def_type, ""); + result = LLVMBuildTrunc(ctx->ac.builder, src[0], def_type, ""); break; case nir_op_bcsel: result = emit_bcsel(&ctx->ac, src[0], src[1], src[2]); @@ -1834,7 +1851,7 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst result = emit_i2b(&ctx->ac, src[0]); break; case nir_op_fquantize2f16: - result = emit_f2f16(ctx, src[0]); + result = emit_f2f16(ctx->nctx, src[0]); break; case nir_op_umul_high: result = emit_umul_high(&ctx->ac, src[0], src[1]); @@ -1854,36 +1871,36 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst case nir_op_fddy_fine: case nir_op_fddx_coarse: case nir_op_fddy_coarse: - result = emit_ddxy(ctx, instr->op, src[0]); + result = emit_ddxy(ctx->nctx, instr->op, src[0]); break; case nir_op_unpack_64_2x32_split_x: { assert(instr->src[0].src.ssa->num_components == 1); - LLVMValueRef tmp = LLVMBuildBitCast(ctx->builder, src[0], - LLVMVectorType(ctx->i32, 2), + LLVMValueRef tmp = LLVMBuildBitCast(ctx->ac.builder, src[0], + LLVMVectorType(ctx->ac.i32, 2), ""); - result = LLVMBuildExtractElement(ctx->builder, tmp, - ctx->i32zero, ""); + result = LLVMBuildExtractElement(ctx->ac.builder, tmp, + ctx->ac.i32_0, ""); break; } case nir_op_unpack_64_2x32_split_y: { assert(instr->src[0].src.ssa->num_components == 1); - LLVMValueRef tmp = LLVMBuildBitCast(ctx->builder, src[0], - LLVMVectorType(ctx->i32, 2), + LLVMValueRef tmp = LLVMBuildBitCast(ctx->ac.builder, src[0], + LLVMVectorType(ctx->ac.i32, 2), ""); - result = LLVMBuildExtractElement(ctx->builder, tmp, - ctx->i32one, ""); + result = LLVMBuildExtractElement(ctx->ac.builder, tmp, + ctx->ac.i32_0, ""); break; } case nir_op_pack_64_2x32_split: { - LLVMValueRef tmp = LLVMGetUndef(LLVMVectorType(ctx->i32, 2)); - tmp = LLVMBuildInsertElement(ctx->builder, tmp, - src[0], ctx->i32zero, ""); - tmp = LLVMBuildInsertElement(ctx->builder, tmp, - src[1], ctx->i32one, ""); - result = LLVMBuildBitCast(ctx->builder, tmp, ctx->i64, ""); + LLVMValueRef tmp = LLVMGetUndef(LLVMVectorType(ctx->ac.i32, 2)); + tmp = LLVMBuildInsertElement(ctx->ac.builder, tmp, + src[0], ctx->ac.i32_0, ""); + tmp = LLVMBuildInsertElement(ctx->ac.builder, tmp, + src[1], ctx->ac.i32_1, ""); + result = LLVMBuildBitCast(ctx->ac.builder, tmp, ctx->ac.i64, ""); break; } @@ -1902,12 +1919,12 @@ static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *inst } } -static void visit_load_const(struct nir_to_llvm_context *ctx, +static void visit_load_const(struct ac_nir_context *ctx, const nir_load_const_instr *instr) { LLVMValueRef values[4], value = NULL; LLVMTypeRef element_type = - LLVMIntTypeInContext(ctx->context, instr->def.bit_size); + LLVMIntTypeInContext(ctx->ac.context, instr->def.bit_size); for (unsigned i = 0; i < instr->def.num_components; ++i) { switch (instr->def.bit_size) { @@ -2168,7 +2185,7 @@ static LLVMValueRef build_tex_intrinsic(struct nir_to_llvm_context *ctx, static LLVMValueRef visit_vulkan_resource_index(struct nir_to_llvm_context *ctx, nir_intrinsic_instr *instr) { - LLVMValueRef index = get_src(ctx, instr->src[0]); + LLVMValueRef index = get_src(ctx->nir, instr->src[0]); unsigned desc_set = nir_intrinsic_desc_set(instr); unsigned binding = nir_intrinsic_binding(instr); LLVMValueRef desc_ptr = ctx->descriptor_sets[desc_set]; @@ -2204,10 +2221,10 @@ static LLVMValueRef visit_load_push_constant(struct nir_to_llvm_context *ctx, LLVMValueRef ptr, addr; addr = LLVMConstInt(ctx->i32, nir_intrinsic_base(instr), 0); - addr = LLVMBuildAdd(ctx->builder, addr, get_src(ctx, instr->src[0]), ""); + addr = LLVMBuildAdd(ctx->builder, addr, get_src(ctx->nir, instr->src[0]), ""); ptr = ac_build_gep0(&ctx->ac, ctx->push_constants, addr); - ptr = cast_ptr(ctx, ptr, get_def_type(ctx, &instr->dest.ssa)); + ptr = cast_ptr(ctx, ptr, get_def_type(ctx->nir, &instr->dest.ssa)); return LLVMBuildLoad(ctx->builder, ptr, ""); } @@ -2215,7 +2232,7 @@ static LLVMValueRef visit_load_push_constant(struct nir_to_llvm_context *ctx, static LLVMValueRef visit_get_buffer_size(struct nir_to_llvm_context *ctx, const nir_intrinsic_instr *instr) { - LLVMValueRef desc = get_src(ctx, instr->src[0]); + LLVMValueRef desc = get_src(ctx->nir, instr->src[0]); return get_buffer_size(ctx, desc, false); } @@ -2223,7 +2240,7 @@ static void visit_store_ssbo(struct nir_to_llvm_context *ctx, nir_intrinsic_instr *instr) { const char *store_name; - LLVMValueRef src_data = get_src(ctx, instr->src[0]); + LLVMValueRef src_data = get_src(ctx->nir, instr->src[0]); LLVMTypeRef data_type = ctx->f32; int elem_size_mult = get_elem_bits(&ctx->ac, LLVMTypeOf(src_data)) / 32; int components_32bit = elem_size_mult * instr->num_components; @@ -2234,7 +2251,7 @@ static void visit_store_ssbo(struct nir_to_llvm_context *ctx, if (ctx->stage == MESA_SHADER_FRAGMENT) ctx->shader_info->fs.writes_memory = true; - params[1] = get_src(ctx, instr->src[1]); + params[1] = get_src(ctx->nir, instr->src[1]); params[2] = LLVMConstInt(ctx->i32, 0, false); /* vindex */ params[4] = ctx->i1false; /* glc */ params[5] = ctx->i1false; /* slc */ @@ -2246,7 +2263,7 @@ static void visit_store_ssbo(struct nir_to_llvm_context *ctx, base_data = trim_vector(ctx, base_data, instr->num_components); base_data = LLVMBuildBitCast(ctx->builder, base_data, data_type, ""); - base_offset = get_src(ctx, instr->src[2]); /* voffset */ + base_offset = get_src(ctx->nir, instr->src[2]); /* voffset */ while (writemask) { int start, count; LLVMValueRef data; @@ -2315,12 +2332,12 @@ static LLVMValueRef visit_atomic_ssbo(struct nir_to_llvm_context *ctx, ctx->shader_info->fs.writes_memory = true; if (instr->intrinsic == nir_intrinsic_ssbo_atomic_comp_swap) { - params[arg_count++] = llvm_extract_elem(ctx, get_src(ctx, instr->src[3]), 0); + params[arg_count++] = llvm_extract_elem(ctx, get_src(ctx->nir, instr->src[3]), 0); } - params[arg_count++] = llvm_extract_elem(ctx, get_src(ctx, instr->src[2]), 0); - params[arg_count++] = get_src(ctx, instr->src[0]); + params[arg_count++] = llvm_extract_elem(ctx, get_src(ctx->nir, instr->src[2]), 0); + params[arg_count++] = get_src(ctx->nir, instr->src[0]); params[arg_count++] = LLVMConstInt(ctx->i32, 0, false); /* vindex */ - params[arg_count++] = get_src(ctx, instr->src[1]); /* voffset */ + params[arg_count++] = get_src(ctx->nir, instr->src[1]); /* voffset */ params[arg_count++] = ctx->i1false; /* slc */ switch (instr->intrinsic) { @@ -2375,7 +2392,7 @@ static LLVMValueRef visit_load_buffer(struct nir_to_llvm_context *ctx, const char *load_name; LLVMTypeRef data_type = ctx->f32; LLVMValueRef offset = LLVMConstInt(ctx->i32, i * 4, false); - offset = LLVMBuildAdd(ctx->builder, get_src(ctx, instr->src[1]), offset, ""); + offset = LLVMBuildAdd(ctx->builder, get_src(ctx->nir, instr->src[1]), offset, ""); if (load_components == 3) data_type = LLVMVectorType(ctx->f32, 4); @@ -2392,7 +2409,7 @@ static LLVMValueRef visit_load_buffer(struct nir_to_llvm_context *ctx, unreachable("unhandled number of components"); LLVMValueRef params[] = { - get_src(ctx, instr->src[0]), + get_src(ctx->nir, instr->src[0]), LLVMConstInt(ctx->i32, 0, false), offset, ctx->i1false, @@ -2418,15 +2435,15 @@ static LLVMValueRef visit_load_buffer(struct nir_to_llvm_context *ctx, } return LLVMBuildBitCast(ctx->builder, ret, - get_def_type(ctx, &instr->dest.ssa), ""); + get_def_type(ctx->nir, &instr->dest.ssa), ""); } static LLVMValueRef visit_load_ubo_buffer(struct nir_to_llvm_context *ctx, const nir_intrinsic_instr *instr) { LLVMValueRef results[8], ret; - LLVMValueRef rsrc = get_src(ctx, instr->src[0]); - LLVMValueRef offset = get_src(ctx, instr->src[1]); + LLVMValueRef rsrc = get_src(ctx->nir, instr->src[0]); + LLVMValueRef offset = get_src(ctx->nir, instr->src[1]); int num_components = instr->num_components; if (instr->dest.ssa.bit_size == 64) @@ -2447,7 +2464,7 @@ static LLVMValueRef visit_load_ubo_buffer(struct nir_to_llvm_context *ctx, ret = ac_build_gather_values(&ctx->ac, results, instr->num_components); return LLVMBuildBitCast(ctx->builder, ret, - get_def_type(ctx, &instr->dest.ssa), ""); + get_def_type(ctx->nir, &instr->dest.ssa), ""); } static void @@ -2469,7 +2486,7 @@ radv_get_deref_offset(struct nir_to_llvm_context *ctx, nir_deref_var *deref, if (vertex_index_ref) { LLVMValueRef vtx = LLVMConstInt(ctx->i32, deref_array->base_offset, false); if (deref_array->deref_array_type == nir_deref_array_type_indirect) { - vtx = LLVMBuildAdd(ctx->builder, vtx, get_src(ctx, deref_array->indirect), ""); + vtx = LLVMBuildAdd(ctx->builder, vtx, get_src(ctx->nir, deref_array->indirect), ""); } *vertex_index_ref = vtx; } @@ -2500,7 +2517,7 @@ radv_get_deref_offset(struct nir_to_llvm_context *ctx, nir_deref_var *deref, continue; assert(deref_array->deref_array_type == nir_deref_array_type_indirect); - index = get_src(ctx, deref_array->indirect); + index = get_src(ctx->nir, deref_array->indirect); stride = LLVMConstInt(ctx->i32, size, 0); local_offset = LLVMBuildMul(ctx->builder, stride, index, ""); @@ -2703,7 +2720,7 @@ load_tcs_input(struct nir_to_llvm_context *ctx, ctx->i32one, ""); } result = ac_build_gather_values(&ctx->ac, value, instr->num_components); - result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), ""); + result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), ""); return result; } @@ -2740,7 +2757,7 @@ load_tcs_output(struct nir_to_llvm_context *ctx, ctx->i32one, ""); } result = ac_build_gather_values(&ctx->ac, value, instr->num_components); - result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), ""); + result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), ""); return result; } @@ -2841,7 +2858,7 @@ load_tes_input(struct nir_to_llvm_context *ctx, result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, instr->num_components, NULL, buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false); result = trim_vector(ctx, result, instr->num_components); - result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), ""); + result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), ""); return result; } @@ -2887,7 +2904,7 @@ load_gs_input(struct nir_to_llvm_context *ctx, } static LLVMValueRef -build_gep_for_deref(struct nir_to_llvm_context *ctx, +build_gep_for_deref(struct ac_nir_context *ctx, nir_deref_var *deref) { struct hash_entry *entry = _mesa_hash_table_search(ctx->vars, deref->var); @@ -2899,10 +2916,10 @@ build_gep_for_deref(struct nir_to_llvm_context *ctx, switch (tail->deref_type) { case nir_deref_type_array: { nir_deref_array *array = nir_deref_as_array(tail); - offset = LLVMConstInt(ctx->i32, array->base_offset, 0); + offset = LLVMConstInt(ctx->ac.i32, array->base_offset, 0); if (array->deref_array_type == nir_deref_array_type_indirect) { - offset = LLVMBuildAdd(ctx->builder, offset, + offset = LLVMBuildAdd(ctx->ac.builder, offset, get_src(ctx, array->indirect), ""); @@ -2912,7 +2929,7 @@ build_gep_for_deref(struct nir_to_llvm_context *ctx, case nir_deref_type_struct: { nir_deref_struct *deref_struct = nir_deref_as_struct(tail); - offset = LLVMConstInt(ctx->i32, + offset = LLVMConstInt(ctx->ac.i32, deref_struct->index, 0); break; } @@ -2975,23 +2992,23 @@ static LLVMValueRef visit_load_var(struct nir_to_llvm_context *ctx, instr->variables[0]->var->type, false); count -= chan / 4; LLVMValueRef tmp_vec = ac_build_gather_values_extended( - &ctx->ac, ctx->locals + idx + chan, count, + &ctx->ac, ctx->nir->locals + idx + chan, count, 4, true); values[chan] = LLVMBuildExtractElement(ctx->builder, tmp_vec, indir_index, ""); } else { - values[chan] = LLVMBuildLoad(ctx->builder, ctx->locals[idx + chan + const_index * 4], ""); + values[chan] = LLVMBuildLoad(ctx->builder, ctx->nir->locals[idx + chan + const_index * 4], ""); } } break; case nir_var_shared: { - LLVMValueRef address = build_gep_for_deref(ctx, + LLVMValueRef address = build_gep_for_deref(ctx->nir, instr->variables[0]); LLVMValueRef val = LLVMBuildLoad(ctx->builder, address, ""); return LLVMBuildBitCast(ctx->builder, val, - get_def_type(ctx, &instr->dest.ssa), + get_def_type(ctx->nir, &instr->dest.ssa), ""); } case nir_var_shader_out: @@ -3003,7 +3020,7 @@ static LLVMValueRef visit_load_var(struct nir_to_llvm_context *ctx, instr->variables[0]->var->type, false); count -= chan / 4; LLVMValueRef tmp_vec = ac_build_gather_values_extended( - &ctx->ac, ctx->outputs + idx + chan, count, + &ctx->ac, ctx->nir->outputs + idx + chan, count, 4, true); values[chan] = LLVMBuildExtractElement(ctx->builder, @@ -3011,7 +3028,7 @@ static LLVMValueRef visit_load_var(struct nir_to_llvm_context *ctx, indir_index, ""); } else { values[chan] = LLVMBuildLoad(ctx->builder, - ctx->outputs[idx + chan + const_index * 4], + ctx->nir->outputs[idx + chan + const_index * 4], ""); } } @@ -3020,7 +3037,7 @@ static LLVMValueRef visit_load_var(struct nir_to_llvm_context *ctx, unreachable("unhandle variable mode"); } ret = ac_build_gather_values(&ctx->ac, values, ve); - return LLVMBuildBitCast(ctx->builder, ret, get_def_type(ctx, &instr->dest.ssa), ""); + return LLVMBuildBitCast(ctx->builder, ret, get_def_type(ctx->nir, &instr->dest.ssa), ""); } static void @@ -3029,7 +3046,7 @@ visit_store_var(struct nir_to_llvm_context *ctx, { LLVMValueRef temp_ptr, value; int idx = instr->variables[0]->var->data.driver_location; - LLVMValueRef src = to_float(&ctx->ac, get_src(ctx, instr->src[0])); + LLVMValueRef src = to_float(&ctx->ac, get_src(ctx->nir, instr->src[0])); int writemask = instr->const_index[0]; LLVMValueRef indir_index; unsigned const_index; @@ -3072,7 +3089,7 @@ visit_store_var(struct nir_to_llvm_context *ctx, instr->variables[0]->var->type, false); count -= chan / 4; LLVMValueRef tmp_vec = ac_build_gather_values_extended( - &ctx->ac, ctx->outputs + idx + chan, count, + &ctx->ac, ctx->nir->outputs + idx + chan, count, stride, true); if (get_llvm_num_components(tmp_vec) > 1) { @@ -3080,11 +3097,11 @@ visit_store_var(struct nir_to_llvm_context *ctx, value, indir_index, ""); } else tmp_vec = value; - build_store_values_extended(ctx, ctx->outputs + idx + chan, + build_store_values_extended(ctx, ctx->nir->outputs + idx + chan, count, stride, tmp_vec); } else { - temp_ptr = ctx->outputs[idx + chan + const_index * stride]; + temp_ptr = ctx->nir->outputs[idx + chan + const_index * stride]; LLVMBuildStore(ctx->builder, value, temp_ptr); } @@ -3101,15 +3118,15 @@ visit_store_var(struct nir_to_llvm_context *ctx, instr->variables[0]->var->type, false); count -= chan / 4; LLVMValueRef tmp_vec = ac_build_gather_values_extended( - &ctx->ac, ctx->locals + idx + chan, count, + &ctx->ac, ctx->nir->locals + idx + chan, count, 4, true); tmp_vec = LLVMBuildInsertElement(ctx->builder, tmp_vec, value, indir_index, ""); - build_store_values_extended(ctx, ctx->locals + idx + chan, + build_store_values_extended(ctx, ctx->nir->locals + idx + chan, count, 4, tmp_vec); } else { - temp_ptr = ctx->locals[idx + chan + const_index * 4]; + temp_ptr = ctx->nir->locals[idx + chan + const_index * 4]; LLVMBuildStore(ctx->builder, value, temp_ptr); } @@ -3117,9 +3134,9 @@ visit_store_var(struct nir_to_llvm_context *ctx, break; case nir_var_shared: { int writemask = instr->const_index[0]; - LLVMValueRef address = build_gep_for_deref(ctx, + LLVMValueRef address = build_gep_for_deref(ctx->nir, instr->variables[0]); - LLVMValueRef val = get_src(ctx, instr->src[0]); + LLVMValueRef val = get_src(ctx->nir, instr->src[0]); unsigned components = glsl_get_vector_elements( nir_deref_tail(&instr->variables[0]->deref)->type); @@ -3261,14 +3278,14 @@ static LLVMValueRef get_image_coords(struct nir_to_llvm_context *ctx, if(instr->variables[0]->deref.child) type = instr->variables[0]->deref.child->type; - LLVMValueRef src0 = get_src(ctx, instr->src[0]); + LLVMValueRef src0 = get_src(ctx->nir, instr->src[0]); LLVMValueRef coords[4]; LLVMValueRef masks[] = { LLVMConstInt(ctx->i32, 0, false), LLVMConstInt(ctx->i32, 1, false), LLVMConstInt(ctx->i32, 2, false), LLVMConstInt(ctx->i32, 3, false), }; LLVMValueRef res; - LLVMValueRef sample_index = llvm_extract_elem(ctx, get_src(ctx, instr->src[1]), 0); + LLVMValueRef sample_index = llvm_extract_elem(ctx, get_src(ctx->nir, instr->src[1]), 0); int count; enum glsl_sampler_dim dim = glsl_get_sampler_dim(type); @@ -3346,7 +3363,7 @@ static LLVMValueRef visit_image_load(struct nir_to_llvm_context *ctx, type = glsl_without_array(type); if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) { params[0] = get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER); - params[1] = LLVMBuildExtractElement(ctx->builder, get_src(ctx, instr->src[0]), + params[1] = LLVMBuildExtractElement(ctx->builder, get_src(ctx->nir, instr->src[0]), LLVMConstInt(ctx->i32, 0, false), ""); /* vindex */ params[2] = LLVMConstInt(ctx->i32, 0, false); /* voffset */ params[3] = ctx->i1false; /* glc */ @@ -3406,9 +3423,9 @@ static void visit_image_store(struct nir_to_llvm_context *ctx, ctx->shader_info->fs.writes_memory = true; if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) { - params[0] = to_float(&ctx->ac, get_src(ctx, instr->src[2])); /* data */ + params[0] = to_float(&ctx->ac, get_src(ctx->nir, instr->src[2])); /* data */ params[1] = get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER); - params[2] = LLVMBuildExtractElement(ctx->builder, get_src(ctx, instr->src[0]), + params[2] = LLVMBuildExtractElement(ctx->builder, get_src(ctx->nir, instr->src[0]), LLVMConstInt(ctx->i32, 0, false), ""); /* vindex */ params[3] = LLVMConstInt(ctx->i32, 0, false); /* voffset */ params[4] = glc; /* glc */ @@ -3421,7 +3438,7 @@ static void visit_image_store(struct nir_to_llvm_context *ctx, LLVMValueRef da = is_da ? ctx->i1true : ctx->i1false; LLVMValueRef slc = ctx->i1false; - params[0] = to_float(&ctx->ac, get_src(ctx, instr->src[2])); + params[0] = to_float(&ctx->ac, get_src(ctx->nir, instr->src[2])); params[1] = get_image_coords(ctx, instr); /* coords */ params[2] = get_sampler_desc(ctx, instr->variables[0], DESC_IMAGE); params[3] = LLVMConstInt(ctx->i32, 15, false); /* dmask */ @@ -3495,12 +3512,12 @@ static LLVMValueRef visit_image_atomic(struct nir_to_llvm_context *ctx, } if (instr->intrinsic == nir_intrinsic_image_atomic_comp_swap) - params[param_count++] = get_src(ctx, instr->src[3]); - params[param_count++] = get_src(ctx, instr->src[2]); + params[param_count++] = get_src(ctx->nir, instr->src[3]); + params[param_count++] = get_src(ctx->nir, instr->src[2]); if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) { params[param_count++] = get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER); - params[param_count++] = LLVMBuildExtractElement(ctx->builder, get_src(ctx, instr->src[0]), + params[param_count++] = LLVMBuildExtractElement(ctx->builder, get_src(ctx->nir, instr->src[0]), LLVMConstInt(ctx->i32, 0, false), ""); /* vindex */ params[param_count++] = ctx->i32zero; /* voffset */ params[param_count++] = ctx->i1false; /* slc */ @@ -3601,7 +3618,7 @@ static void emit_discard_if(struct nir_to_llvm_context *ctx, ctx->shader_info->fs.can_discard = true; cond = LLVMBuildICmp(ctx->builder, LLVMIntNE, - get_src(ctx, instr->src[0]), + get_src(ctx->nir, instr->src[0]), ctx->i32zero, ""); cond = LLVMBuildSelect(ctx->builder, cond, @@ -3625,11 +3642,11 @@ static LLVMValueRef visit_var_atomic(struct nir_to_llvm_context *ctx, const nir_intrinsic_instr *instr) { LLVMValueRef ptr, result; - LLVMValueRef src = get_src(ctx, instr->src[0]); - ptr = build_gep_for_deref(ctx, instr->variables[0]); + LLVMValueRef src = get_src(ctx->nir, instr->src[0]); + ptr = build_gep_for_deref(ctx->nir, instr->variables[0]); if (instr->intrinsic == nir_intrinsic_var_atomic_comp_swap) { - LLVMValueRef src1 = get_src(ctx, instr->src[1]); + LLVMValueRef src1 = get_src(ctx->nir, instr->src[1]); result = LLVMBuildAtomicCmpXchg(ctx->builder, ptr, src, src1, LLVMAtomicOrderingSequentiallyConsistent, @@ -3749,7 +3766,7 @@ static LLVMValueRef visit_interp(struct nir_to_llvm_context *ctx, case nir_intrinsic_interp_var_at_sample: case nir_intrinsic_interp_var_at_offset: location = INTERP_CENTER; - src0 = get_src(ctx, instr->src[0]); + src0 = get_src(ctx->nir, instr->src[0]); break; default: break; @@ -3866,7 +3883,7 @@ visit_emit_vertex(struct nir_to_llvm_context *ctx, /* loop num outputs */ idx = 0; for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) { - LLVMValueRef *out_ptr = &ctx->outputs[i * 4]; + LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4]; int length = 4; int slot = idx; int slot_inc = 1; @@ -3928,7 +3945,7 @@ visit_load_tess_coord(struct nir_to_llvm_context *ctx, LLVMValueRef result = ac_build_gather_values(&ctx->ac, coord, instr->num_components); return LLVMBuildBitCast(ctx->builder, result, - get_def_type(ctx, &instr->dest.ssa), ""); + get_def_type(ctx->nir, &instr->dest.ssa), ""); } static void visit_intrinsic(struct nir_to_llvm_context *ctx, @@ -4109,7 +4126,7 @@ static void visit_intrinsic(struct nir_to_llvm_context *ctx, break; } if (result) { - _mesa_hash_table_insert(ctx->defs, &instr->dest.ssa, result); + _mesa_hash_table_insert(ctx->nir->defs, &instr->dest.ssa, result); } } @@ -4163,7 +4180,7 @@ static LLVMValueRef get_sampler_desc(struct nir_to_llvm_context *ctx, assert(child->deref_array_type != nir_deref_array_type_wildcard); offset += child->base_offset * stride; if (child->deref_array_type == nir_deref_array_type_indirect) { - index = get_src(ctx, child->indirect); + index = get_src(ctx->nir, child->indirect); } constant_index = child->base_offset; @@ -4312,39 +4329,39 @@ static void visit_tex(struct nir_to_llvm_context *ctx, nir_tex_instr *instr) for (unsigned i = 0; i < instr->num_srcs; i++) { switch (instr->src[i].src_type) { case nir_tex_src_coord: - coord = get_src(ctx, instr->src[i].src); + coord = get_src(ctx->nir, instr->src[i].src); break; case nir_tex_src_projector: break; case nir_tex_src_comparator: - comparator = get_src(ctx, instr->src[i].src); + comparator = get_src(ctx->nir, instr->src[i].src); break; case nir_tex_src_offset: - offsets = get_src(ctx, instr->src[i].src); + offsets = get_src(ctx->nir, instr->src[i].src); const_src = i; break; case nir_tex_src_bias: - bias = get_src(ctx, instr->src[i].src); + bias = get_src(ctx->nir, instr->src[i].src); break; case nir_tex_src_lod: { nir_const_value *val = nir_src_as_const_value(instr->src[i].src); if (val && val->i32[0] == 0) lod_is_zero = true; - lod = get_src(ctx, instr->src[i].src); + lod = get_src(ctx->nir, instr->src[i].src); break; } case nir_tex_src_ms_index: - sample_index = get_src(ctx, instr->src[i].src); + sample_index = get_src(ctx->nir, instr->src[i].src); break; case nir_tex_src_ms_mcs: break; case nir_tex_src_ddx: - ddx = get_src(ctx, instr->src[i].src); + ddx = get_src(ctx->nir, instr->src[i].src); num_deriv_comp = instr->src[i].src.ssa->num_components; break; case nir_tex_src_ddy: - ddy = get_src(ctx, instr->src[i].src); + ddy = get_src(ctx->nir, instr->src[i].src); break; case nir_tex_src_texture_offset: case nir_tex_src_sampler_offset: @@ -4576,21 +4593,21 @@ write_result: if (result) { assert(instr->dest.is_ssa); result = to_integer(&ctx->ac, result); - _mesa_hash_table_insert(ctx->defs, &instr->dest.ssa, result); + _mesa_hash_table_insert(ctx->nir->defs, &instr->dest.ssa, result); } } -static void visit_phi(struct nir_to_llvm_context *ctx, nir_phi_instr *instr) +static void visit_phi(struct ac_nir_context *ctx, nir_phi_instr *instr) { LLVMTypeRef type = get_def_type(ctx, &instr->dest.ssa); - LLVMValueRef result = LLVMBuildPhi(ctx->builder, type, ""); + LLVMValueRef result = LLVMBuildPhi(ctx->ac.builder, type, ""); _mesa_hash_table_insert(ctx->defs, &instr->dest.ssa, result); _mesa_hash_table_insert(ctx->phis, instr, result); } -static void visit_post_phi(struct nir_to_llvm_context *ctx, +static void visit_post_phi(struct ac_nir_context *ctx, nir_phi_instr *instr, LLVMValueRef llvm_phi) { @@ -4602,7 +4619,7 @@ static void visit_post_phi(struct nir_to_llvm_context *ctx, } } -static void phi_post_pass(struct nir_to_llvm_context *ctx) +static void phi_post_pass(struct ac_nir_context *ctx) { struct hash_entry *entry; hash_table_foreach(ctx->phis, entry) { @@ -4612,31 +4629,31 @@ static void phi_post_pass(struct nir_to_llvm_context *ctx) } -static void visit_ssa_undef(struct nir_to_llvm_context *ctx, +static void visit_ssa_undef(struct ac_nir_context *ctx, const nir_ssa_undef_instr *instr) { unsigned num_components = instr->def.num_components; LLVMValueRef undef; if (num_components == 1) - undef = LLVMGetUndef(ctx->i32); + undef = LLVMGetUndef(ctx->ac.i32); else { - undef = LLVMGetUndef(LLVMVectorType(ctx->i32, num_components)); + undef = LLVMGetUndef(LLVMVectorType(ctx->ac.i32, num_components)); } _mesa_hash_table_insert(ctx->defs, &instr->def, undef); } -static void visit_jump(struct nir_to_llvm_context *ctx, +static void visit_jump(struct ac_nir_context *ctx, const nir_jump_instr *instr) { switch (instr->type) { case nir_jump_break: - LLVMBuildBr(ctx->builder, ctx->break_block); - LLVMClearInsertionPosition(ctx->builder); + LLVMBuildBr(ctx->ac.builder, ctx->break_block); + LLVMClearInsertionPosition(ctx->ac.builder); break; case nir_jump_continue: - LLVMBuildBr(ctx->builder, ctx->continue_block); - LLVMClearInsertionPosition(ctx->builder); + LLVMBuildBr(ctx->ac.builder, ctx->continue_block); + LLVMClearInsertionPosition(ctx->ac.builder); break; default: fprintf(stderr, "Unknown NIR jump instr: "); @@ -4646,12 +4663,12 @@ static void visit_jump(struct nir_to_llvm_context *ctx, } } -static void visit_cf_list(struct nir_to_llvm_context *ctx, +static void visit_cf_list(struct ac_nir_context *ctx, struct exec_list *list); -static void visit_block(struct nir_to_llvm_context *ctx, nir_block *block) +static void visit_block(struct ac_nir_context *ctx, nir_block *block) { - LLVMBasicBlockRef llvm_block = LLVMGetInsertBlock(ctx->builder); + LLVMBasicBlockRef llvm_block = LLVMGetInsertBlock(ctx->ac.builder); nir_foreach_instr(instr, block) { switch (instr->type) { @@ -4662,10 +4679,10 @@ static void visit_block(struct nir_to_llvm_context *ctx, nir_block *block) visit_load_const(ctx, nir_instr_as_load_const(instr)); break; case nir_instr_type_intrinsic: - visit_intrinsic(ctx, nir_instr_as_intrinsic(instr)); + visit_intrinsic(ctx->nctx, nir_instr_as_intrinsic(instr)); break; case nir_instr_type_tex: - visit_tex(ctx, nir_instr_as_tex(instr)); + visit_tex(ctx->nctx, nir_instr_as_tex(instr)); break; case nir_instr_type_phi: visit_phi(ctx, nir_instr_as_phi(instr)); @@ -4687,61 +4704,63 @@ static void visit_block(struct nir_to_llvm_context *ctx, nir_block *block) _mesa_hash_table_insert(ctx->defs, block, llvm_block); } -static void visit_if(struct nir_to_llvm_context *ctx, nir_if *if_stmt) +static void visit_if(struct ac_nir_context *ctx, nir_if *if_stmt) { LLVMValueRef value = get_src(ctx, if_stmt->condition); + LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx->ac.builder)); LLVMBasicBlockRef merge_block = - LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, ""); + LLVMAppendBasicBlockInContext(ctx->ac.context, fn, ""); LLVMBasicBlockRef if_block = - LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, ""); + LLVMAppendBasicBlockInContext(ctx->ac.context, fn, ""); LLVMBasicBlockRef else_block = merge_block; if (!exec_list_is_empty(&if_stmt->else_list)) else_block = LLVMAppendBasicBlockInContext( - ctx->context, ctx->main_function, ""); + ctx->ac.context, fn, ""); - LLVMValueRef cond = LLVMBuildICmp(ctx->builder, LLVMIntNE, value, - LLVMConstInt(ctx->i32, 0, false), ""); - LLVMBuildCondBr(ctx->builder, cond, if_block, else_block); + LLVMValueRef cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE, value, + LLVMConstInt(ctx->ac.i32, 0, false), ""); + LLVMBuildCondBr(ctx->ac.builder, cond, if_block, else_block); - LLVMPositionBuilderAtEnd(ctx->builder, if_block); + LLVMPositionBuilderAtEnd(ctx->ac.builder, if_block); visit_cf_list(ctx, &if_stmt->then_list); - if (LLVMGetInsertBlock(ctx->builder)) - LLVMBuildBr(ctx->builder, merge_block); + if (LLVMGetInsertBlock(ctx->ac.builder)) + LLVMBuildBr(ctx->ac.builder, merge_block); if (!exec_list_is_empty(&if_stmt->else_list)) { - LLVMPositionBuilderAtEnd(ctx->builder, else_block); + LLVMPositionBuilderAtEnd(ctx->ac.builder, else_block); visit_cf_list(ctx, &if_stmt->else_list); - if (LLVMGetInsertBlock(ctx->builder)) - LLVMBuildBr(ctx->builder, merge_block); + if (LLVMGetInsertBlock(ctx->ac.builder)) + LLVMBuildBr(ctx->ac.builder, merge_block); } - LLVMPositionBuilderAtEnd(ctx->builder, merge_block); + LLVMPositionBuilderAtEnd(ctx->ac.builder, merge_block); } -static void visit_loop(struct nir_to_llvm_context *ctx, nir_loop *loop) +static void visit_loop(struct ac_nir_context *ctx, nir_loop *loop) { + LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx->ac.builder)); LLVMBasicBlockRef continue_parent = ctx->continue_block; LLVMBasicBlockRef break_parent = ctx->break_block; ctx->continue_block = - LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, ""); + LLVMAppendBasicBlockInContext(ctx->ac.context, fn, ""); ctx->break_block = - LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, ""); + LLVMAppendBasicBlockInContext(ctx->ac.context, fn, ""); - LLVMBuildBr(ctx->builder, ctx->continue_block); - LLVMPositionBuilderAtEnd(ctx->builder, ctx->continue_block); + LLVMBuildBr(ctx->ac.builder, ctx->continue_block); + LLVMPositionBuilderAtEnd(ctx->ac.builder, ctx->continue_block); visit_cf_list(ctx, &loop->body); - if (LLVMGetInsertBlock(ctx->builder)) - LLVMBuildBr(ctx->builder, ctx->continue_block); - LLVMPositionBuilderAtEnd(ctx->builder, ctx->break_block); + if (LLVMGetInsertBlock(ctx->ac.builder)) + LLVMBuildBr(ctx->ac.builder, ctx->continue_block); + LLVMPositionBuilderAtEnd(ctx->ac.builder, ctx->break_block); ctx->continue_block = continue_parent; ctx->break_block = break_parent; } -static void visit_cf_list(struct nir_to_llvm_context *ctx, +static void visit_cf_list(struct ac_nir_context *ctx, struct exec_list *list) { foreach_list_typed(nir_cf_node, node, node, list) @@ -5018,7 +5037,7 @@ handle_shader_output_decl(struct nir_to_llvm_context *ctx, for (unsigned i = 0; i < attrib_count; ++i) { for (unsigned chan = 0; chan < 4; chan++) { - ctx->outputs[radeon_llvm_reg_index_soa(idx + i, chan)] = + ctx->nir->outputs[radeon_llvm_reg_index_soa(idx + i, chan)] = si_build_alloca_undef(ctx, ctx->f32, ""); } } @@ -5088,7 +5107,7 @@ glsl_to_llvm_type(struct nir_to_llvm_context *ctx, } static void -setup_locals(struct nir_to_llvm_context *ctx, +setup_locals(struct ac_nir_context *ctx, struct nir_function *func) { int i, j; @@ -5105,19 +5124,19 @@ setup_locals(struct nir_to_llvm_context *ctx, for (i = 0; i < ctx->num_locals; i++) { for (j = 0; j < 4; j++) { ctx->locals[i * 4 + j] = - si_build_alloca_undef(ctx, ctx->f32, "temp"); + si_build_alloca_undef(ctx->nctx, ctx->ac.f32, "temp"); } } } static void -setup_shared(struct nir_to_llvm_context *ctx, +setup_shared(struct ac_nir_context *ctx, struct nir_shader *nir) { nir_foreach_variable(variable, &nir->shared) { LLVMValueRef shared = LLVMAddGlobalInAddressSpace( - ctx->module, glsl_to_llvm_type(ctx, variable->type), + ctx->ac.module, glsl_to_llvm_type(ctx->nctx, variable->type), variable->name ? variable->name : "", LOCAL_ADDR_SPACE); _mesa_hash_table_insert(ctx->vars, variable, shared); @@ -5322,7 +5341,7 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx, i = VARYING_SLOT_CLIP_DIST0; for (j = 0; j < ctx->num_output_clips + ctx->num_output_culls; j++) slots[j] = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(i, j)], "")); + ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], "")); for (i = ctx->num_output_clips + ctx->num_output_culls; i < 8; i++) slots[i] = LLVMGetUndef(ctx->f32); @@ -5345,26 +5364,26 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx, if (ctx->output_mask & (1ull << VARYING_SLOT_POS)) { for (unsigned j = 0; j < 4; j++) pos_values[j] = LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_POS, j)], ""); + ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_POS, j)], ""); } si_llvm_init_export_args(ctx, pos_values, V_008DFC_SQ_EXP_POS, &pos_args[0]); if (ctx->output_mask & (1ull << VARYING_SLOT_PSIZ)) { outinfo->writes_pointsize = true; psize_value = LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_PSIZ, 0)], ""); + ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_PSIZ, 0)], ""); } if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) { outinfo->writes_layer = true; layer_value = LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)], ""); + ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)], ""); } if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) { outinfo->writes_viewport_index = true; viewport_index_value = LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_VIEWPORT, 0)], ""); + ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_VIEWPORT, 0)], ""); } uint32_t mask = ((outinfo->writes_pointsize == true ? 1 : 0) | @@ -5412,7 +5431,7 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx, for (unsigned j = 0; j < 4; j++) values[j] = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(i, j)], "")); + ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], "")); if (i == VARYING_SLOT_LAYER) { target = V_008DFC_SQ_EXP_PARAM + param_count; @@ -5468,7 +5487,7 @@ handle_es_outputs_post(struct nir_to_llvm_context *ctx, int j; uint64_t max_output_written = 0; for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) { - LLVMValueRef *out_ptr = &ctx->outputs[i * 4]; + LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4]; int param_index; int length = 4; @@ -5506,7 +5525,7 @@ handle_ls_outputs_post(struct nir_to_llvm_context *ctx) vertex_dw_stride, ""); for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) { - LLVMValueRef *out_ptr = &ctx->outputs[i * 4]; + LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4]; int length = 4; if (!(ctx->output_mask & (1ull << i))) @@ -5841,20 +5860,20 @@ handle_fs_outputs_post(struct nir_to_llvm_context *ctx) if (i == FRAG_RESULT_DEPTH) { ctx->shader_info->fs.writes_z = true; depth = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(i, 0)], "")); + ctx->nir->outputs[radeon_llvm_reg_index_soa(i, 0)], "")); } else if (i == FRAG_RESULT_STENCIL) { ctx->shader_info->fs.writes_stencil = true; stencil = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(i, 0)], "")); + ctx->nir->outputs[radeon_llvm_reg_index_soa(i, 0)], "")); } else if (i == FRAG_RESULT_SAMPLE_MASK) { ctx->shader_info->fs.writes_sample_mask = true; samplemask = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(i, 0)], "")); + ctx->nir->outputs[radeon_llvm_reg_index_soa(i, 0)], "")); } else { bool last = false; for (unsigned j = 0; j < 4; j++) values[j] = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder, - ctx->outputs[radeon_llvm_reg_index_soa(i, j)], "")); + ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], "")); if (!ctx->shader_info->fs.writes_z && !ctx->shader_info->fs.writes_stencil && !ctx->shader_info->fs.writes_sample_mask) last = ctx->output_mask <= ((1ull << (i + 1)) - 1); @@ -6028,6 +6047,52 @@ ac_nir_get_max_workgroup_size(enum chip_class chip_class, return max_workgroup_size; } +void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi, + struct nir_shader *nir, struct nir_to_llvm_context *nctx) +{ + struct ac_nir_context ctx = {}; + struct nir_function *func; + + ctx.ac = *ac; + ctx.abi = abi; + + ctx.nctx = nctx; + if (nctx) + nctx->nir = &ctx; + + ctx.stage = nir->stage; + + nir_foreach_variable(variable, &nir->outputs) + handle_shader_output_decl(nctx, variable); + + ctx.defs = _mesa_hash_table_create(NULL, _mesa_hash_pointer, + _mesa_key_pointer_equal); + ctx.phis = _mesa_hash_table_create(NULL, _mesa_hash_pointer, + _mesa_key_pointer_equal); + ctx.vars = _mesa_hash_table_create(NULL, _mesa_hash_pointer, + _mesa_key_pointer_equal); + + func = (struct nir_function *)exec_list_get_head(&nir->functions); + + setup_locals(&ctx, func); + + if (nir->stage == MESA_SHADER_COMPUTE) + setup_shared(&ctx, nir); + + visit_cf_list(&ctx, &func->impl->body); + phi_post_pass(&ctx); + + handle_shader_outputs_post(nctx); + + free(ctx.locals); + ralloc_free(ctx.defs); + ralloc_free(ctx.phis); + ralloc_free(ctx.vars); + + if (nctx) + nctx->nir = NULL; +} + static LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm, struct nir_shader *nir, @@ -6035,7 +6100,6 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm, const struct ac_nir_compiler_options *options) { struct nir_to_llvm_context ctx = {0}; - struct nir_function *func; unsigned i; ctx.options = options; ctx.shader_info = shader_info; @@ -6083,45 +6147,22 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm, ac_setup_rings(&ctx); + ctx.num_output_clips = nir->info.clip_distance_array_size; + ctx.num_output_culls = nir->info.cull_distance_array_size; + nir_foreach_variable(variable, &nir->inputs) handle_shader_input_decl(&ctx, variable); if (nir->stage == MESA_SHADER_FRAGMENT) handle_fs_inputs_pre(&ctx, nir); - ctx.num_output_clips = nir->info.clip_distance_array_size; - ctx.num_output_culls = nir->info.cull_distance_array_size; - - nir_foreach_variable(variable, &nir->outputs) - handle_shader_output_decl(&ctx, variable); - - ctx.defs = _mesa_hash_table_create(NULL, _mesa_hash_pointer, - _mesa_key_pointer_equal); - ctx.phis = _mesa_hash_table_create(NULL, _mesa_hash_pointer, - _mesa_key_pointer_equal); - ctx.vars = _mesa_hash_table_create(NULL, _mesa_hash_pointer, - _mesa_key_pointer_equal); - - func = (struct nir_function *)exec_list_get_head(&nir->functions); - - setup_locals(&ctx, func); + ac_nir_translate(&ctx.ac, &ctx.abi, nir, &ctx); - if (nir->stage == MESA_SHADER_COMPUTE) - setup_shared(&ctx, nir); - - visit_cf_list(&ctx, &func->impl->body); - phi_post_pass(&ctx); - - handle_shader_outputs_post(&ctx); LLVMBuildRetVoid(ctx.builder); ac_llvm_finalize_module(&ctx); ac_nir_eliminate_const_vs_outputs(&ctx); - free(ctx.locals); - ralloc_free(ctx.defs); - ralloc_free(ctx.phis); - ralloc_free(ctx.vars); if (nir->stage == MESA_SHADER_GEOMETRY) { unsigned addclip = ctx.num_output_clips + ctx.num_output_culls > 4; @@ -6355,7 +6396,7 @@ ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx) AC_FUNC_ATTR_LEGACY); LLVMBuildStore(ctx->builder, - to_float(&ctx->ac, value), ctx->outputs[radeon_llvm_reg_index_soa(i, j)]); + to_float(&ctx->ac, value), ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]); } idx += slot_inc; } @@ -6395,11 +6436,20 @@ void ac_create_gs_copy_shader(LLVMTargetMachineRef tm, ctx.num_output_clips = geom_shader->info.clip_distance_array_size; ctx.num_output_culls = geom_shader->info.cull_distance_array_size; + struct ac_nir_context nir_ctx = {}; + nir_ctx.ac = ctx.ac; + nir_ctx.abi = &ctx.abi; + + nir_ctx.nctx = &ctx; + ctx.nir = &nir_ctx; + nir_foreach_variable(variable, &geom_shader->outputs) handle_shader_output_decl(&ctx, variable); ac_gs_copy_shader_emit(&ctx); + ctx.nir = NULL; + LLVMBuildRetVoid(ctx.builder); ac_llvm_finalize_module(&ctx); diff --git a/src/amd/common/ac_nir_to_llvm.h b/src/amd/common/ac_nir_to_llvm.h index 54d54891ac..791c694c78 100644 --- a/src/amd/common/ac_nir_to_llvm.h +++ b/src/amd/common/ac_nir_to_llvm.h @@ -36,6 +36,8 @@ struct ac_shader_config; struct nir_shader; struct radv_pipeline_layout; +struct ac_llvm_context; +struct ac_shader_abi; struct ac_vs_variant_key { uint32_t instance_rate_inputs; @@ -221,4 +223,8 @@ void ac_create_gs_copy_shader(LLVMTargetMachineRef tm, const struct ac_nir_compiler_options *options, bool dump_shader); +struct nir_to_llvm_context; +void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi, + struct nir_shader *nir, struct nir_to_llvm_context *nctx); + #endif /* AC_NIR_TO_LLVM_H */ |