diff options
| -rw-r--r-- | src/amd/common/ac_nir_to_llvm.c | 187 |
1 files changed, 92 insertions, 95 deletions
diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c index 09f47fe8d09..a052a7109d4 100644 --- a/src/amd/common/ac_nir_to_llvm.c +++ b/src/amd/common/ac_nir_to_llvm.c @@ -80,7 +80,6 @@ struct nir_to_llvm_context { unsigned max_workgroup_size; LLVMContextRef context; LLVMModuleRef module; - LLVMBuilderRef builder; LLVMValueRef main_function; LLVMValueRef descriptor_sets[AC_UD_MAX_SETS]; @@ -395,7 +394,7 @@ get_tcs_out_patch_stride(struct nir_to_llvm_context *ctx) static LLVMValueRef get_tcs_out_patch0_offset(struct nir_to_llvm_context *ctx) { - return LLVMBuildMul(ctx->builder, + return LLVMBuildMul(ctx->ac.builder, unpack_param(&ctx->ac, ctx->tcs_out_offsets, 0, 16), LLVMConstInt(ctx->ac.i32, 4, false), ""); } @@ -403,7 +402,7 @@ get_tcs_out_patch0_offset(struct nir_to_llvm_context *ctx) static LLVMValueRef get_tcs_out_patch0_patch_data_offset(struct nir_to_llvm_context *ctx) { - return LLVMBuildMul(ctx->builder, + return LLVMBuildMul(ctx->ac.builder, unpack_param(&ctx->ac, ctx->tcs_out_offsets, 16, 16), LLVMConstInt(ctx->ac.i32, 4, false), ""); } @@ -414,7 +413,7 @@ get_tcs_in_current_patch_offset(struct nir_to_llvm_context *ctx) LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx); LLVMValueRef rel_patch_id = get_rel_patch_id(ctx); - return LLVMBuildMul(ctx->builder, patch_stride, rel_patch_id, ""); + return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, ""); } static LLVMValueRef @@ -424,8 +423,8 @@ get_tcs_out_current_patch_offset(struct nir_to_llvm_context *ctx) LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx); LLVMValueRef rel_patch_id = get_rel_patch_id(ctx); - return LLVMBuildAdd(ctx->builder, patch0_offset, - LLVMBuildMul(ctx->builder, patch_stride, + return LLVMBuildAdd(ctx->ac.builder, patch0_offset, + LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, ""), ""); } @@ -438,8 +437,8 @@ get_tcs_out_current_patch_data_offset(struct nir_to_llvm_context *ctx) LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx); LLVMValueRef rel_patch_id = get_rel_patch_id(ctx); - return LLVMBuildAdd(ctx->builder, patch0_patch_data_offset, - LLVMBuildMul(ctx->builder, patch_stride, + return LLVMBuildAdd(ctx->ac.builder, patch0_patch_data_offset, + LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, ""), ""); } @@ -1021,7 +1020,7 @@ static void create_function(struct nir_to_llvm_context *ctx, } ctx->main_function = create_llvm_function( - ctx->context, ctx->module, ctx->builder, NULL, 0, &args, + ctx->context, ctx->module, ctx->ac.builder, NULL, 0, &args, ctx->max_workgroup_size, ctx->options->unsafe_math); set_llvm_calling_convention(ctx->main_function, stage); @@ -1046,7 +1045,7 @@ static void create_function(struct nir_to_llvm_context *ctx, ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr", LLVMPointerType(ctx->ac.i8, AC_CONST_ADDR_SPACE), NULL, 0, AC_FUNC_ATTR_READNONE); - ctx->ring_offsets = LLVMBuildBitCast(ctx->builder, ctx->ring_offsets, + ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets, ac_array_in_const_addr_space(ctx->ac.v4i32), ""); } } @@ -2372,8 +2371,8 @@ radv_load_resource(struct ac_shader_abi *abi, LLVMValueRef index, stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false); offset = LLVMConstInt(ctx->ac.i32, base_offset, false); - index = LLVMBuildMul(ctx->builder, index, stride, ""); - offset = LLVMBuildAdd(ctx->builder, offset, index, ""); + index = LLVMBuildMul(ctx->ac.builder, index, stride, ""); + offset = LLVMBuildAdd(ctx->ac.builder, offset, index, ""); desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset); desc_ptr = cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32); @@ -2771,15 +2770,15 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct nir_to_llvm_context *ctx, vertices_per_patch = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 9, 6); num_patches = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 0, 9); - total_vertices = LLVMBuildMul(ctx->builder, vertices_per_patch, + total_vertices = LLVMBuildMul(ctx->ac.builder, vertices_per_patch, num_patches, ""); constant16 = LLVMConstInt(ctx->ac.i32, 16, false); if (vertex_index) { - base_addr = LLVMBuildMul(ctx->builder, rel_patch_id, + base_addr = LLVMBuildMul(ctx->ac.builder, rel_patch_id, vertices_per_patch, ""); - base_addr = LLVMBuildAdd(ctx->builder, base_addr, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, vertex_index, ""); param_stride = total_vertices; @@ -2788,17 +2787,17 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct nir_to_llvm_context *ctx, param_stride = num_patches; } - base_addr = LLVMBuildAdd(ctx->builder, base_addr, - LLVMBuildMul(ctx->builder, param_index, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, + LLVMBuildMul(ctx->ac.builder, param_index, param_stride, ""), ""); - base_addr = LLVMBuildMul(ctx->builder, base_addr, constant16, ""); + base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, ""); if (!vertex_index) { LLVMValueRef patch_data_offset = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 16, 16); - base_addr = LLVMBuildAdd(ctx->builder, base_addr, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, patch_data_offset, ""); } return base_addr; @@ -2814,7 +2813,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_params(struct nir_to_llvm_context LLVMValueRef param_index; if (indir_index) - param_index = LLVMBuildAdd(ctx->builder, LLVMConstInt(ctx->ac.i32, param, false), + param_index = LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, param, false), indir_index, ""); else { if (const_index && !is_compact) @@ -2848,25 +2847,25 @@ get_dw_address(struct nir_to_llvm_context *ctx, { if (vertex_index) { - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, - LLVMBuildMul(ctx->builder, + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, + LLVMBuildMul(ctx->ac.builder, vertex_index, stride, ""), ""); } if (indir_index) - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, - LLVMBuildMul(ctx->builder, indir_index, + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, + LLVMBuildMul(ctx->ac.builder, indir_index, LLVMConstInt(ctx->ac.i32, 4, false), ""), ""); else if (const_index && !compact_const_index) - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, LLVMConstInt(ctx->ac.i32, const_index, false), ""); - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, LLVMConstInt(ctx->ac.i32, param * 4, false), ""); if (const_index && compact_const_index) - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, LLVMConstInt(ctx->ac.i32, const_index, false), ""); return dw_addr; } @@ -2907,7 +2906,7 @@ load_tcs_varyings(struct ac_shader_abi *abi, for (unsigned i = 0; i < num_components + component; i++) { value[i] = ac_lds_load(&ctx->ac, dw_addr); - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, ""); } result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component); @@ -2976,7 +2975,7 @@ store_tcs_output(struct ac_shader_abi *abi, if (store_lds || is_tess_factor) { LLVMValueRef dw_addr_chan = - LLVMBuildAdd(ctx->builder, dw_addr, + LLVMBuildAdd(ctx->ac.builder, dw_addr, LLVMConstInt(ctx->ac.i32, chan, false), ""); ac_lds_store(&ctx->ac, dw_addr_chan, value); } @@ -3021,7 +3020,7 @@ load_tes_input(struct ac_shader_abi *abi, is_compact, vertex_index, param_index); LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false); - buf_addr = LLVMBuildAdd(ctx->builder, buf_addr, comp_offset, ""); + buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, ""); result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL, buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false); @@ -3046,7 +3045,7 @@ load_gs_input(struct ac_shader_abi *abi, vtx_offset_param = vertex_index; assert(vtx_offset_param < 6); - vtx_offset = LLVMBuildMul(ctx->builder, ctx->gs_vtx_offset[vtx_offset_param], + vtx_offset = LLVMBuildMul(ctx->ac.builder, ctx->gs_vtx_offset[vtx_offset_param], LLVMConstInt(ctx->ac.i32, 4, false), ""); param = shader_io_get_unique_index(location); @@ -3069,7 +3068,7 @@ load_gs_input(struct ac_shader_abi *abi, vtx_offset, soffset, 0, 1, 0, true, false); - value[i] = LLVMBuildBitCast(ctx->builder, value[i], + value[i] = LLVMBuildBitCast(ctx->ac.builder, value[i], type, ""); } } @@ -4003,10 +4002,10 @@ static LLVMValueRef load_sample_position(struct ac_shader_abi *abi, LLVMValueRef result; LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false)); - ptr = LLVMBuildBitCast(ctx->builder, ptr, + ptr = LLVMBuildBitCast(ctx->ac.builder, ptr, ac_array_in_const_addr_space(ctx->ac.v2f32), ""); - sample_id = LLVMBuildAdd(ctx->builder, sample_id, ctx->sample_pos_offset, ""); + sample_id = LLVMBuildAdd(ctx->ac.builder, sample_id, ctx->sample_pos_offset, ""); result = ac_build_load_invariant(&ctx->ac, ptr, sample_id); return result; @@ -4165,7 +4164,7 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr assert(stream == 0); /* Write vertex attribute values to GSVS ring */ - gs_next_vertex = LLVMBuildLoad(ctx->builder, + gs_next_vertex = LLVMBuildLoad(ctx->ac.builder, ctx->gs_next_vertex, ""); @@ -4174,7 +4173,7 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr * have any effect, and GS threads have no externally observable * effects other than emitting vertices. */ - can_emit = LLVMBuildICmp(ctx->builder, LLVMIntULT, gs_next_vertex, + can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex, LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), ""); ac_build_kill_if_false(&ctx->ac, can_emit); @@ -4196,13 +4195,13 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr slot_inc = 2; } for (unsigned j = 0; j < length; j++) { - LLVMValueRef out_val = LLVMBuildLoad(ctx->builder, + LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], ""); LLVMValueRef voffset = LLVMConstInt(ctx->ac.i32, (slot * 4 + j) * ctx->gs_max_out_vertices, false); - voffset = LLVMBuildAdd(ctx->builder, voffset, gs_next_vertex, ""); - voffset = LLVMBuildMul(ctx->builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), ""); + voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, ""); + voffset = LLVMBuildMul(ctx->ac.builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), ""); - out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->ac.i32, ""); + out_val = LLVMBuildBitCast(ctx->ac.builder, out_val, ctx->ac.i32, ""); ac_build_buffer_store_dword(&ctx->ac, ctx->gsvs_ring, out_val, 1, @@ -4212,9 +4211,9 @@ visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addr idx += slot_inc; } - gs_next_vertex = LLVMBuildAdd(ctx->builder, gs_next_vertex, + gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex, ctx->ac.i32_1, ""); - LLVMBuildStore(ctx->builder, gs_next_vertex, ctx->gs_next_vertex); + LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex); ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (0 << 8), ctx->gs_wave_id); } @@ -4239,8 +4238,8 @@ load_tess_coord(struct ac_shader_abi *abi) }; if (ctx->tes_primitive_mode == GL_TRIANGLES) - coord[2] = LLVMBuildFSub(ctx->builder, ctx->ac.f32_1, - LLVMBuildFAdd(ctx->builder, coord[0], coord[1], ""), ""); + coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1, + LLVMBuildFAdd(ctx->ac.builder, coord[0], coord[1], ""), ""); return ac_build_gather_values(&ctx->ac, coord, 3); } @@ -4531,7 +4530,7 @@ static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi, LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md); - result = LLVMBuildLoad(ctx->builder, buffer_ptr, ""); + result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, ""); LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md); return result; @@ -4544,7 +4543,7 @@ static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md); - result = LLVMBuildLoad(ctx->builder, buffer_ptr, ""); + result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, ""); LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md); return result; @@ -4565,7 +4564,7 @@ static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi, unsigned offset = binding->offset; unsigned stride = binding->size; unsigned type_size; - LLVMBuilderRef builder = ctx->builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMTypeRef type; assert(base_index < layout->binding_count); @@ -5327,7 +5326,7 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx, for (unsigned i = 0; i < attrib_count; ++i, ++idx) { if (ctx->options->key.vs.instance_rate_inputs & (1u << (index + i))) { - buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.instance_id, + buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.instance_id, ctx->abi.start_instance, ""); if (ctx->options->key.vs.as_ls) { ctx->shader_info->vs.vgpr_comp_cnt = @@ -5337,7 +5336,7 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx, MAX2(1, ctx->shader_info->vs.vgpr_comp_cnt); } } else - buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.vertex_id, + buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id, ctx->abi.base_vertex, ""); t_offset = LLVMConstInt(ctx->ac.i32, index + i, false); @@ -5353,7 +5352,7 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx, for (unsigned chan = 0; chan < 4; chan++) { LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false); ctx->inputs[radeon_llvm_reg_index_soa(idx, chan)] = - ac_to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->builder, + ac_to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->ac.builder, input, llvm_chan, "")); } } @@ -5383,12 +5382,12 @@ static void interp_fs_input(struct nir_to_llvm_context *ctx, * to NaN. */ if (interp) { - interp_param = LLVMBuildBitCast(ctx->builder, interp_param, + interp_param = LLVMBuildBitCast(ctx->ac.builder, interp_param, ctx->ac.v2f32, ""); - i = LLVMBuildExtractElement(ctx->builder, interp_param, + i = LLVMBuildExtractElement(ctx->ac.builder, interp_param, ctx->ac.i32_0, ""); - j = LLVMBuildExtractElement(ctx->builder, interp_param, + j = LLVMBuildExtractElement(ctx->ac.builder, interp_param, ctx->ac.i32_1, ""); } @@ -5468,9 +5467,9 @@ prepare_interp_optimize(struct nir_to_llvm_context *ctx, } if (uses_center && uses_centroid) { - LLVMValueRef sel = LLVMBuildICmp(ctx->builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, ""); - ctx->persp_centroid = LLVMBuildSelect(ctx->builder, sel, ctx->persp_center, ctx->persp_centroid, ""); - ctx->linear_centroid = LLVMBuildSelect(ctx->builder, sel, ctx->linear_center, ctx->linear_centroid, ""); + LLVMValueRef sel = LLVMBuildICmp(ctx->ac.builder, LLVMIntSLT, ctx->abi.prim_mask, ctx->ac.i32_0, ""); + ctx->persp_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->persp_center, ctx->persp_centroid, ""); + ctx->linear_centroid = LLVMBuildSelect(ctx->ac.builder, sel, ctx->linear_center, ctx->linear_centroid, ""); } } @@ -5882,7 +5881,7 @@ radv_load_output(struct nir_to_llvm_context *ctx, unsigned index, unsigned chan) LLVMValueRef output = ctx->nir->outputs[radeon_llvm_reg_index_soa(index, chan)]; - return LLVMBuildLoad(ctx->builder, output, ""); + return LLVMBuildLoad(ctx->ac.builder, output, ""); } static void @@ -5905,7 +5904,7 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx, si_build_alloca_undef(&ctx->ac, ctx->ac.f32, ""); } - LLVMBuildStore(ctx->builder, ac_to_float(&ctx->ac, ctx->abi.view_index), *tmp_out); + LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, ctx->abi.view_index), *tmp_out); ctx->output_mask |= 1ull << VARYING_SLOT_LAYER; } @@ -5987,10 +5986,10 @@ handle_vs_outputs_post(struct nir_to_llvm_context *ctx, */ LLVMValueRef v = viewport_index_value; v = ac_to_integer(&ctx->ac, v); - v = LLVMBuildShl(ctx->builder, v, + v = LLVMBuildShl(ctx->ac.builder, v, LLVMConstInt(ctx->ac.i32, 16, false), ""); - v = LLVMBuildOr(ctx->builder, v, + v = LLVMBuildOr(ctx->ac.builder, v, ac_to_integer(&ctx->ac, pos_args[1].out[2]), ""); pos_args[1].out[2] = ac_to_float(&ctx->ac, v); @@ -6108,18 +6107,18 @@ handle_es_outputs_post(struct nir_to_llvm_context *ctx, param_index = shader_io_get_unique_index(i); if (lds_base) { - dw_addr = LLVMBuildAdd(ctx->builder, lds_base, + dw_addr = LLVMBuildAdd(ctx->ac.builder, lds_base, LLVMConstInt(ctx->ac.i32, param_index * 4, false), ""); } for (j = 0; j < length; j++) { - LLVMValueRef out_val = LLVMBuildLoad(ctx->builder, out_ptr[j], ""); - out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->ac.i32, ""); + LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[j], ""); + out_val = LLVMBuildBitCast(ctx->ac.builder, out_val, ctx->ac.i32, ""); if (ctx->ac.chip_class >= GFX9) { ac_lds_store(&ctx->ac, dw_addr, - LLVMBuildLoad(ctx->builder, out_ptr[j], "")); - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->ac.i32_1, ""); + LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "")); + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, ""); } else { ac_build_buffer_store_dword(&ctx->ac, ctx->esgs_ring, @@ -6137,7 +6136,7 @@ handle_ls_outputs_post(struct nir_to_llvm_context *ctx) { LLVMValueRef vertex_id = ctx->rel_auto_id; LLVMValueRef vertex_dw_stride = unpack_param(&ctx->ac, ctx->ls_out_layout, 13, 8); - LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->builder, vertex_id, + LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id, vertex_dw_stride, ""); for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) { @@ -6153,13 +6152,13 @@ handle_ls_outputs_post(struct nir_to_llvm_context *ctx) mark_tess_output(ctx, false, param); if (length > 4) mark_tess_output(ctx, false, param + 1); - LLVMValueRef dw_addr = LLVMBuildAdd(ctx->builder, base_dw_addr, + LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr, LLVMConstInt(ctx->ac.i32, param * 4, false), ""); for (unsigned j = 0; j < length; j++) { ac_lds_store(&ctx->ac, dw_addr, - LLVMBuildLoad(ctx->builder, out_ptr[j], "")); - dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->ac.i32_1, ""); + LLVMBuildLoad(ctx->ac.builder, out_ptr[j], "")); + dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr, ctx->ac.i32_1, ""); } } } @@ -6182,7 +6181,7 @@ ac_build_insert_new_block(struct nir_to_llvm_context *ctx, const char *name) LLVMBasicBlockRef new_block; /* get current basic block */ - current_block = LLVMGetInsertBlock(ctx->builder); + current_block = LLVMGetInsertBlock(ctx->ac.builder); /* chqeck if there's another block after this one */ next_block = LLVMGetNextBasicBlock(current_block); @@ -6203,7 +6202,7 @@ ac_nir_build_if(struct ac_build_if_state *ifthen, struct nir_to_llvm_context *ctx, LLVMValueRef condition) { - LLVMBasicBlockRef block = LLVMGetInsertBlock(ctx->builder); + LLVMBasicBlockRef block = LLVMGetInsertBlock(ctx->ac.builder); memset(ifthen, 0, sizeof *ifthen); ifthen->ctx = ctx; @@ -6220,7 +6219,7 @@ ac_nir_build_if(struct ac_build_if_state *ifthen, "if-true-block"); /* successive code goes into the true block */ - LLVMPositionBuilderAtEnd(ctx->builder, ifthen->true_block); + LLVMPositionBuilderAtEnd(ctx->ac.builder, ifthen->true_block); } /** @@ -6229,7 +6228,7 @@ ac_nir_build_if(struct ac_build_if_state *ifthen, static void ac_nir_build_endif(struct ac_build_if_state *ifthen) { - LLVMBuilderRef builder = ifthen->ctx->builder; + LLVMBuilderRef builder = ifthen->ctx->ac.builder; /* Insert branch to the merge block from current block */ LLVMBuildBr(builder, ifthen->merge_block); @@ -6289,7 +6288,7 @@ write_tess_factors(struct nir_to_llvm_context *ctx) } ac_nir_build_if(&if_ctx, ctx, - LLVMBuildICmp(ctx->builder, LLVMIntEQ, + LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, invocation_id, ctx->ac.i32_0, "")); tess_inner_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER); @@ -6298,9 +6297,9 @@ write_tess_factors(struct nir_to_llvm_context *ctx) mark_tess_output(ctx, true, tess_inner_index); mark_tess_output(ctx, true, tess_outer_index); lds_base = get_tcs_out_current_patch_data_offset(ctx); - lds_inner = LLVMBuildAdd(ctx->builder, lds_base, + lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base, LLVMConstInt(ctx->ac.i32, tess_inner_index * 4, false), ""); - lds_outer = LLVMBuildAdd(ctx->builder, lds_base, + lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base, LLVMConstInt(ctx->ac.i32, tess_outer_index * 4, false), ""); for (i = 0; i < 4; i++) { @@ -6311,20 +6310,20 @@ write_tess_factors(struct nir_to_llvm_context *ctx) // LINES reverseal if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) { outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer); - lds_outer = LLVMBuildAdd(ctx->builder, lds_outer, + lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer, ctx->ac.i32_1, ""); outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer); } else { for (i = 0; i < outer_comps; i++) { outer[i] = out[i] = ac_lds_load(&ctx->ac, lds_outer); - lds_outer = LLVMBuildAdd(ctx->builder, lds_outer, + lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_outer, ctx->ac.i32_1, ""); } for (i = 0; i < inner_comps; i++) { inner[i] = out[outer_comps+i] = ac_lds_load(&ctx->ac, lds_inner); - lds_inner = LLVMBuildAdd(ctx->builder, lds_inner, + lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_inner, ctx->ac.i32_1, ""); } } @@ -6339,13 +6338,13 @@ write_tess_factors(struct nir_to_llvm_context *ctx) buffer = ctx->hs_ring_tess_factor; tf_base = ctx->tess_factor_offset; - byteoffset = LLVMBuildMul(ctx->builder, rel_patch_id, + byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id, LLVMConstInt(ctx->ac.i32, 4 * stride, false), ""); unsigned tf_offset = 0; if (ctx->options->chip_class <= VI) { ac_nir_build_if(&inner_if_ctx, ctx, - LLVMBuildICmp(ctx->builder, LLVMIntEQ, + LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, rel_patch_id, ctx->ac.i32_0, "")); /* Store the dynamic HS control word. */ @@ -6556,7 +6555,7 @@ static void ac_llvm_finalize_module(struct nir_to_llvm_context * ctx) LLVMRunFunctionPassManager(passmgr, ctx->main_function); LLVMFinalizeFunctionPassManager(passmgr); - LLVMDisposeBuilder(ctx->builder); + LLVMDisposeBuilder(ctx->ac.builder); LLVMDisposePassManager(passmgr); } @@ -6609,12 +6608,12 @@ ac_setup_rings(struct nir_to_llvm_context *ctx) ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_ESGS_GS, false)); ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_GS, false)); - ctx->gsvs_ring = LLVMBuildBitCast(ctx->builder, ctx->gsvs_ring, ctx->ac.v4i32, ""); + ctx->gsvs_ring = LLVMBuildBitCast(ctx->ac.builder, ctx->gsvs_ring, ctx->ac.v4i32, ""); - ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->ac.i32, 2, false), ""); - tmp = LLVMBuildExtractElement(ctx->builder, ctx->gsvs_ring, ctx->ac.i32_1, ""); - tmp = LLVMBuildOr(ctx->builder, tmp, ctx->gsvs_ring_stride, ""); - ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, tmp, ctx->ac.i32_1, ""); + ctx->gsvs_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->ac.i32, 2, false), ""); + tmp = LLVMBuildExtractElement(ctx->ac.builder, ctx->gsvs_ring, ctx->ac.i32_1, ""); + tmp = LLVMBuildOr(ctx->ac.builder, tmp, ctx->gsvs_ring_stride, ""); + ctx->gsvs_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->gsvs_ring, tmp, ctx->ac.i32_1, ""); } if (ctx->stage == MESA_SHADER_TESS_CTRL || @@ -6751,8 +6750,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm, options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH : AC_FLOAT_MODE_DEFAULT; - ctx.builder = ac_create_builder(ctx.context, float_mode); - ctx.ac.builder = ctx.builder; + ctx.ac.builder = ac_create_builder(ctx.context, float_mode); memset(shader_info, 0, sizeof(*shader_info)); @@ -6883,7 +6881,7 @@ LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm, } } - LLVMBuildRetVoid(ctx.builder); + LLVMBuildRetVoid(ctx.ac.builder); if (options->dump_preoptir) ac_dump_module(ctx.module); @@ -7105,7 +7103,7 @@ static void ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx) { LLVMValueRef vtx_offset = - LLVMBuildMul(ctx->builder, ctx->abi.vertex_id, + LLVMBuildMul(ctx->ac.builder, ctx->abi.vertex_id, LLVMConstInt(ctx->ac.i32, 4, false), ""); int idx = 0; @@ -7135,7 +7133,7 @@ ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx) vtx_offset, soffset, 0, 1, 1, true, false); - LLVMBuildStore(ctx->builder, + LLVMBuildStore(ctx->ac.builder, ac_to_float(&ctx->ac, value), ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]); } idx += slot_inc; @@ -7168,8 +7166,7 @@ void ac_create_gs_copy_shader(LLVMTargetMachineRef tm, options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH : AC_FLOAT_MODE_DEFAULT; - ctx.builder = ac_create_builder(ctx.context, float_mode); - ctx.ac.builder = ctx.builder; + ctx.ac.builder = ac_create_builder(ctx.context, float_mode); ctx.stage = MESA_SHADER_VERTEX; create_function(&ctx, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX); @@ -7196,7 +7193,7 @@ void ac_create_gs_copy_shader(LLVMTargetMachineRef tm, ctx.nir = NULL; - LLVMBuildRetVoid(ctx.builder); + LLVMBuildRetVoid(ctx.ac.builder); ac_llvm_finalize_module(&ctx); |