aboutsummaryrefslogtreecommitdiffstats
path: root/src/amd/common/ac_nir_to_llvm.c
diff options
context:
space:
mode:
Diffstat (limited to 'src/amd/common/ac_nir_to_llvm.c')
-rw-r--r--src/amd/common/ac_nir_to_llvm.c187
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);
generated by cgit v1.2.3 (git 2.39.1) at 2024-11-30 21:36:23 +0000