From: Marek Olšák <marek.ol...@amd.com>
---
src/gallium/drivers/radeonsi/si_shader.c | 116 ++++++++++-----------
src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c | 4 +-
.../drivers/radeonsi/si_shader_tgsi_setup.c | 46 ++++----
3 files changed, 79 insertions(+), 87 deletions(-)
diff --git a/src/gallium/drivers/radeonsi/si_shader.c
b/src/gallium/drivers/radeonsi/si_shader.c
index 0200172..29d3dd4 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -296,21 +296,21 @@ get_tcs_out_current_patch_data_offset(struct
si_shader_context *ctx)
return LLVMBuildAdd(gallivm->builder, patch0_patch_data_offset,
LLVMBuildMul(gallivm->builder, patch_stride,
rel_patch_id, ""),
"");
}
static LLVMValueRef get_instance_index_for_fetch(
struct si_shader_context *ctx,
unsigned param_start_instance, unsigned divisor)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef result = LLVMGetParam(ctx->main_fn,
ctx->param_instance_id);
/* The division must be done before START_INSTANCE is added. */
if (divisor > 1)
result = LLVMBuildUDiv(gallivm->builder, result,
LLVMConstInt(ctx->i32, divisor, 0), "");
return LLVMBuildAdd(gallivm->builder, result,
@@ -331,22 +331,21 @@ static LLVMValueRef extract_double_to_float(struct
si_shader_context *ctx,
LLVMValueRef value = LLVMBuildExtractElement(builder, dvec2, index, "");
return LLVMBuildFPTrunc(builder, value, ctx->f32, "");
}
static void declare_input_vs(
struct si_shader_context *ctx,
unsigned input_index,
const struct tgsi_full_declaration *decl,
LLVMValueRef out[4])
{
- struct lp_build_context *base = &ctx->bld_base.base;
- struct gallivm_state *gallivm = base->gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
unsigned chan;
unsigned fix_fetch;
unsigned num_fetches;
unsigned fetch_stride;
LLVMValueRef t_list_ptr;
LLVMValueRef t_offset;
LLVMValueRef t_list;
LLVMValueRef vertex_index;
@@ -567,21 +566,21 @@ static LLVMValueRef get_primitive_id(struct
lp_build_tgsi_context *bld_base,
}
/**
* Return the value of tgsi_ind_register for indexing.
* This is the indirect index with the constant offset added to it.
*/
static LLVMValueRef get_indirect_index(struct si_shader_context *ctx,
const struct tgsi_ind_register *ind,
int rel_index)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef result;
result = ctx->addrs[ind->Index][ind->Swizzle];
result = LLVMBuildLoad(gallivm->builder, result, "");
result = LLVMBuildAdd(gallivm->builder, result,
LLVMConstInt(ctx->i32, rel_index, 0), "");
return result;
}
/**
@@ -607,21 +606,21 @@ static LLVMValueRef get_bounded_indirect_index(struct
si_shader_context *ctx,
/**
* Calculate a dword address given an input or output register and a stride.
*/
static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
const struct tgsi_full_dst_register *dst,
const struct tgsi_full_src_register *src,
LLVMValueRef vertex_dw_stride,
LLVMValueRef base_addr)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct tgsi_shader_info *info = &ctx->shader->selector->info;
ubyte *name, *index, *array_first;
int first, param;
struct tgsi_full_dst_register reg;
/* Set the register description. The address computation is the same
* for sources and destinations. */
if (src) {
reg.Register.File = src->Register.File;
reg.Register.Index = src->Register.Index;
@@ -706,21 +705,21 @@ static LLVMValueRef get_dw_address(struct
si_shader_context *ctx,
* - per patch attribute 0 of patch 1
* ...
*
* Note that every attribute has 4 components.
*/
static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
LLVMValueRef rel_patch_id,
LLVMValueRef vertex_index,
LLVMValueRef param_index)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef base_addr, vertices_per_patch, num_patches, total_vertices;
LLVMValueRef param_stride, constant16;
vertices_per_patch = unpack_param(ctx, SI_PARAM_TCS_OFFCHIP_LAYOUT, 9,
6);
num_patches = unpack_param(ctx, SI_PARAM_TCS_OFFCHIP_LAYOUT, 0, 9);
total_vertices = LLVMBuildMul(gallivm->builder, vertices_per_patch,
num_patches, "");
constant16 = LLVMConstInt(ctx->i32, 16, 0);
if (vertex_index) {
@@ -750,21 +749,21 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct
si_shader_context *ctx,
patch_data_offset, "");
}
return base_addr;
}
static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
struct si_shader_context *ctx,
const struct tgsi_full_dst_register
*dst,
const struct tgsi_full_src_register
*src)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct tgsi_shader_info *info = &ctx->shader->selector->info;
ubyte *name, *index, *array_first;
struct tgsi_full_src_register reg;
LLVMValueRef vertex_index = NULL;
LLVMValueRef param_index = NULL;
unsigned param_index_base, param_base;
reg = src ? *src : tgsi_full_src_register_from_dst(dst);
if (reg.Register.Dimension) {
@@ -814,21 +813,21 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
return get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx),
vertex_index, param_index);
}
static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
enum tgsi_opcode_type type, unsigned swizzle,
LLVMValueRef buffer, LLVMValueRef offset,
LLVMValueRef base, bool readonly_memory)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef value, value2;
LLVMTypeRef llvm_type = tgsi2llvmtype(bld_base, type);
LLVMTypeRef vec_type = LLVMVectorType(llvm_type, 4);
if (swizzle == ~0) {
value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base,
offset,
0, 1, 0, readonly_memory);
return LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
}
@@ -856,30 +855,30 @@ static LLVMValueRef buffer_load(struct
lp_build_tgsi_context *bld_base,
*
* \param type output value type
* \param swizzle offset (typically 0..3); it can be ~0, which loads a
vec4
* \param dw_addr address in dwords
*/
static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
enum tgsi_opcode_type type, unsigned swizzle,
LLVMValueRef dw_addr)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef value;
if (swizzle == ~0) {
LLVMValueRef values[TGSI_NUM_CHANNELS];
for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; chan++)
values[chan] = lds_load(bld_base, type, chan, dw_addr);
- return lp_build_gather_values(bld_base->base.gallivm, values,
+ return lp_build_gather_values(gallivm, values,
TGSI_NUM_CHANNELS);
}
dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
LLVMConstInt(ctx->i32, swizzle, 0));
value = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
if (tgsi_type_is_64bit(type)) {
LLVMValueRef value2;
dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
@@ -897,21 +896,21 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context
*bld_base,
*
* \param swizzle offset (typically 0..3)
* \param dw_addr address in dwords
* \param value value to store
*/
static void lds_store(struct lp_build_tgsi_context *bld_base,
unsigned swizzle, LLVMValueRef dw_addr,
LLVMValueRef value)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
LLVMConstInt(ctx->i32, swizzle, 0));
value = LLVMBuildBitCast(gallivm->builder, value, ctx->i32, "");
ac_build_indexed_store(&ctx->ac, ctx->lds,
dw_addr, value);
}
static LLVMValueRef fetch_input_tcs(
@@ -967,21 +966,21 @@ static LLVMValueRef fetch_input_tes(
return buffer_load(bld_base, type, swizzle, buffer, base, addr, true);
}
static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_instruction *inst,
const struct tgsi_opcode_info *info,
LLVMValueRef dst[4])
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_dst_register *reg = &inst->Dst[0];
const struct tgsi_shader_info *sh_info = &ctx->shader->selector->info;
unsigned chan_index;
LLVMValueRef dw_addr, stride;
LLVMValueRef rw_buffers, buffer, base, buf_addr;
LLVMValueRef values[4];
bool skip_lds_store;
bool is_tess_factor = false;
/* Only handle per-patch and per-vertex outputs here.
@@ -1038,59 +1037,58 @@ static void store_output_tcs(struct
lp_build_tgsi_context *bld_base,
values[chan_index] = value;
if (inst->Dst[0].Register.WriteMask != 0xF && !is_tess_factor) {
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 1,
buf_addr, base,
4 * chan_index, 1, 0, true,
false);
}
}
if (inst->Dst[0].Register.WriteMask == 0xF && !is_tess_factor) {
- LLVMValueRef value =
lp_build_gather_values(bld_base->base.gallivm,
+ LLVMValueRef value = lp_build_gather_values(gallivm,
values, 4);
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4,
buf_addr,
base, 0, 1, 0, true, false);
}
}
static LLVMValueRef fetch_input_gs(
struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_src_register *reg,
enum tgsi_opcode_type type,
unsigned swizzle)
{
- struct lp_build_context *base = &bld_base->base;
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
struct lp_build_context *uint = &ctx->bld_base.uint_bld;
- struct gallivm_state *gallivm = base->gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef vtx_offset, soffset;
unsigned vtx_offset_param;
struct tgsi_shader_info *info = &shader->selector->info;
unsigned semantic_name = info->input_semantic_name[reg->Register.Index];
unsigned semantic_index =
info->input_semantic_index[reg->Register.Index];
unsigned param;
LLVMValueRef value;
if (swizzle != ~0 && semantic_name == TGSI_SEMANTIC_PRIMID)
return get_primitive_id(bld_base, swizzle);
if (!reg->Register.Dimension)
return NULL;
if (swizzle == ~0) {
LLVMValueRef values[TGSI_NUM_CHANNELS];
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
values[chan] = fetch_input_gs(bld_base, reg, type,
chan);
}
- return lp_build_gather_values(bld_base->base.gallivm, values,
+ return lp_build_gather_values(gallivm, values,
TGSI_NUM_CHANNELS);
}
/* Get the vertex offset parameter */
vtx_offset_param = reg->Dimension.Index;
if (vtx_offset_param < 2) {
vtx_offset_param += SI_PARAM_VTX0_OFFSET;
} else {
assert(vtx_offset_param < 6);
vtx_offset_param += SI_PARAM_VTX2_OFFSET - 2;
@@ -1167,23 +1165,21 @@ static void interp_fs_input(struct si_shader_context
*ctx,
unsigned input_index,
unsigned semantic_name,
unsigned semantic_index,
unsigned num_interp_inputs,
unsigned colors_read_mask,
LLVMValueRef interp_param,
LLVMValueRef prim_mask,
LLVMValueRef face,
LLVMValueRef result[4])
{
- struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
- struct lp_build_context *base = &bld_base->base;
- struct gallivm_state *gallivm = base->gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef attr_number;
LLVMValueRef i, j;
unsigned chan;
/* fs.constant returns the param from the middle vertex, so it's not
* really useful for flat shading. It's meant to be used for custom
* interpolation (but the intrinsic can't fetch from the other two
* vertices).
*
@@ -1640,21 +1636,21 @@ static LLVMValueRef fetch_constant(
LLVMValueRef addr, bufp;
LLVMValueRef result;
if (swizzle == LP_CHAN_ALL) {
unsigned chan;
LLVMValueRef values[4];
for (chan = 0; chan < TGSI_NUM_CHANNELS; ++chan)
values[chan] = fetch_constant(bld_base, reg, type,
chan);
- return lp_build_gather_values(bld_base->base.gallivm, values,
4);
+ return lp_build_gather_values(&ctx->gallivm, values, 4);
}
buf = reg->Register.Dimension ? reg->Dimension.Index : 0;
idx = reg->Register.Index * 4 + swizzle;
if (reg->Register.Dimension && reg->Dimension.Indirect) {
LLVMValueRef ptr = LLVMGetParam(ctx->main_fn,
SI_PARAM_CONST_BUFFERS);
LLVMValueRef index;
index = get_bounded_indirect_index(ctx, ®->DimIndirect,
reg->Dimension.Index,
@@ -1713,21 +1709,21 @@ static LLVMValueRef
si_llvm_pack_two_int32_as_int16(struct si_shader_context *ct
}
/* Initialize arguments for the shader export intrinsic */
static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
LLVMValueRef *values,
unsigned target,
struct ac_export_args *args)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct lp_build_context *base = &bld_base->base;
- LLVMBuilderRef builder = base->gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
LLVMValueRef val[4];
unsigned spi_shader_col_format = V_028714_SPI_SHADER_32_ABGR;
unsigned chan;
bool is_int8, is_int10;
/* Default is 0xf. Adjusted below depending on the format. */
args->enabled_channels = 0xf; /* writemask */
/* Specify whether the EXEC mask represents the valid mask */
args->valid_mask = 0;
@@ -1783,21 +1779,21 @@ static void si_llvm_init_export_args(struct
lp_build_tgsi_context *bld_base,
for (chan = 0; chan < 2; chan++) {
LLVMValueRef pack_args[2] = {
values[2 * chan],
values[2 * chan + 1]
};
LLVMValueRef packed;
packed = ac_build_cvt_pkrtz_f16(&ctx->ac, pack_args);
args->out[chan] =
- LLVMBuildBitCast(base->gallivm->builder,
+ LLVMBuildBitCast(ctx->gallivm.builder,
packed, ctx->f32, "");
}
break;
case V_028714_SPI_SHADER_UNORM16_ABGR:
for (chan = 0; chan < 4; chan++) {
val[chan] = ac_build_clamp(&ctx->ac, values[chan]);
val[chan] = LLVMBuildFMul(builder, val[chan],
LLVMConstReal(ctx->f32,
65535), "");
val[chan] = LLVMBuildFAdd(builder, val[chan],
@@ -1922,21 +1918,21 @@ static void si_alpha_test(struct lp_build_tgsi_context
*bld_base,
} else {
ac_build_kill(&ctx->ac, NULL);
}
}
static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context
*bld_base,
LLVMValueRef alpha,
unsigned samplemask_param)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef coverage;
/* alpha = alpha * popcount(coverage) / SI_NUM_SMOOTH_AA_SAMPLES */
coverage = LLVMGetParam(ctx->main_fn,
samplemask_param);
coverage = bitcast(bld_base, TGSI_TYPE_SIGNED, coverage);
coverage = lp_build_intrinsic(gallivm->builder, "llvm.ctpop.i32",
ctx->i32,
&coverage, 1, LP_FUNC_ATTR_READNONE);
@@ -2300,29 +2296,29 @@ handle_semantic:
pos_args[1].out[1] = base->zero; /* Y */
pos_args[1].out[2] = base->zero; /* Z */
pos_args[1].out[3] = base->zero; /* W */
if (shader->selector->info.writes_psize)
pos_args[1].out[0] = psize_value;
if (shader->selector->info.writes_edgeflag) {
/* The output is a float, but the hw expects an integer
* with the first bit containing the edge flag. */
- edgeflag_value = LLVMBuildFPToUI(base->gallivm->builder,
+ edgeflag_value = LLVMBuildFPToUI(ctx->gallivm.builder,
edgeflag_value,
ctx->i32, "");
edgeflag_value = lp_build_min(&bld_base->int_bld,
edgeflag_value,
ctx->i32_1);
/* The LLVM intrinsic expects a float. */
- pos_args[1].out[1] =
LLVMBuildBitCast(base->gallivm->builder,
+ pos_args[1].out[1] =
LLVMBuildBitCast(ctx->gallivm.builder,
edgeflag_value,
ctx->f32, "");
}
if (shader->selector->info.writes_layer)
pos_args[1].out[2] = layer_value;
if (shader->selector->info.writes_viewport_index)
pos_args[1].out[3] = viewport_index_value;
}
@@ -2347,21 +2343,21 @@ handle_semantic:
}
}
/**
* Forward all outputs from the vertex shader to the TES. This is only used
* for the fixed function TCS.
*/
static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef invocation_id, rw_buffers, buffer, buffer_offset;
LLVMValueRef lds_vertex_stride, lds_vertex_offset, lds_base;
uint64_t inputs;
invocation_id = unpack_param(ctx, SI_PARAM_REL_IDS, 8, 5);
rw_buffers = LLVMGetParam(ctx->main_fn, SI_PARAM_RW_BUFFERS);
buffer = ac_build_indexed_load_const(&ctx->ac, rw_buffers,
LLVMConstInt(ctx->i32, SI_HS_RING_TESS_OFFCHIP, 0));
@@ -2393,21 +2389,21 @@ static void si_copy_tcs_inputs(struct
lp_build_tgsi_context *bld_base)
buffer_offset, 0, 1, 0, true,
false);
}
}
static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
LLVMValueRef rel_patch_id,
LLVMValueRef invocation_id,
LLVMValueRef
tcs_out_current_patch_data_offset)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_shader *shader = ctx->shader;
unsigned tess_inner_index, tess_outer_index;
LLVMValueRef lds_base, lds_inner, lds_outer, byteoffset, buffer;
LLVMValueRef out[6], vec0, vec1, rw_buffers, tf_base, inner[4],
outer[4];
unsigned stride, outer_comps, inner_comps, i;
struct lp_build_if_state if_ctx, inner_if_ctx;
si_llvm_emit_barrier(NULL, bld_base, NULL);
/* Do this only for invocation 0, because the tess levels are per-patch,
@@ -2563,21 +2559,21 @@ static void si_llvm_emit_tcs_epilogue(struct
lp_build_tgsi_context *bld_base)
LLVMValueRef rel_patch_id, invocation_id, tf_lds_offset;
LLVMValueRef offchip_soffset, offchip_layout;
si_copy_tcs_inputs(bld_base);
rel_patch_id = get_rel_patch_id(ctx);
invocation_id = unpack_param(ctx, SI_PARAM_REL_IDS, 8, 5);
tf_lds_offset = get_tcs_out_current_patch_data_offset(ctx);
/* Return epilog parameters from this function. */
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
LLVMValueRef ret = ctx->return_value;
LLVMValueRef rw_buffers, rw0, rw1, tf_soffset;
unsigned vgpr;
/* RW_BUFFERS pointer */
rw_buffers = LLVMGetParam(ctx->main_fn,
SI_PARAM_RW_BUFFERS);
rw_buffers = LLVMBuildPtrToInt(builder, rw_buffers, ctx->i64, "");
rw_buffers = LLVMBuildBitCast(builder, rw_buffers, ctx->v2i32, "");
rw0 = LLVMBuildExtractElement(builder, rw_buffers,
@@ -2610,21 +2606,21 @@ static void si_llvm_emit_tcs_epilogue(struct
lp_build_tgsi_context *bld_base)
ret = LLVMBuildInsertValue(builder, ret, invocation_id, vgpr++, "");
ret = LLVMBuildInsertValue(builder, ret, tf_lds_offset, vgpr++, "");
ctx->return_value = ret;
}
static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
struct tgsi_shader_info *info = &shader->selector->info;
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
unsigned i, chan;
LLVMValueRef vertex_id = LLVMGetParam(ctx->main_fn,
ctx->param_rel_auto_id);
LLVMValueRef vertex_dw_stride =
unpack_param(ctx, SI_PARAM_LS_OUT_LAYOUT, 13, 8);
LLVMValueRef base_dw_addr = LLVMBuildMul(gallivm->builder, vertex_id,
vertex_dw_stride, "");
/* Write outputs to LDS. The next shader (TCS aka HS) will read
* its inputs from it. */
@@ -2639,21 +2635,21 @@ static void si_llvm_emit_ls_epilogue(struct
lp_build_tgsi_context *bld_base)
for (chan = 0; chan < 4; chan++) {
lds_store(bld_base, chan, dw_addr,
LLVMBuildLoad(gallivm->builder,
out_ptr[chan], ""));
}
}
}
static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_shader *es = ctx->shader;
struct tgsi_shader_info *info = &es->selector->info;
LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
ctx->param_es2gs_offset);
unsigned chan;
int i;
for (i = 0; i < info->num_outputs; i++) {
LLVMValueRef *out_ptr = ctx->outputs[i];
int param_index;
@@ -2682,21 +2678,21 @@ static void si_llvm_emit_gs_epilogue(struct
lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE,
LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
}
static void si_llvm_emit_vs_epilogue(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct tgsi_shader_info *info = &ctx->shader->selector->info;
struct si_shader_output_values *outputs = NULL;
int i,j;
assert(!ctx->shader->is_gs_copy_shader);
outputs = MALLOC((info->num_outputs + 1) * sizeof(outputs[0]));
/* Vertex color clamping.
*
@@ -2815,21 +2811,21 @@ static void si_export_mrt_z(struct
lp_build_tgsi_context *bld_base,
args.out[2] = base->undef; /* B, sample mask */
args.out[3] = base->undef; /* A, alpha to mask */
if (format == V_028710_SPI_SHADER_UINT16_ABGR) {
assert(!depth);
args.compr = 1; /* COMPR flag */
if (stencil) {
/* Stencil should be in X[23:16]. */
stencil = bitcast(bld_base, TGSI_TYPE_UNSIGNED,
stencil);
- stencil = LLVMBuildShl(base->gallivm->builder, stencil,
+ stencil = LLVMBuildShl(ctx->gallivm.builder, stencil,
LLVMConstInt(ctx->i32, 16, 0),
"");
args.out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT,
stencil);
mask |= 0x3;
}
if (samplemask) {
/* SampleMask should be in Y[15:0]. */
args.out[1] = samplemask;
mask |= 0xc;
}
} else {
@@ -2963,23 +2959,22 @@ static void si_export_null(struct lp_build_tgsi_context
*bld_base)
* vN+1 = Stencil
* vN+2 = SampleMask
* vN+3 = SampleMaskIn (used for OpenGL smoothing)
*
* The alpha-ref SGPR is returned via its original location.
*/
static void si_llvm_return_fs_outputs(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
- struct lp_build_context *base = &bld_base->base;
struct tgsi_shader_info *info = &shader->selector->info;
- LLVMBuilderRef builder = base->gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
unsigned i, j, first_vgpr, vgpr;
LLVMValueRef color[8][4] = {};
LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
LLVMValueRef ret;
/* Read the output values. */
for (i = 0; i < info->num_outputs; i++) {
unsigned semantic_name = info->output_semantic_name[i];
unsigned semantic_index = info->output_semantic_index[i];
@@ -3049,21 +3044,21 @@ static void si_llvm_return_fs_outputs(struct
lp_build_tgsi_context *bld_base)
/**
* Given a v8i32 resource descriptor for a buffer, extract the size of the
* buffer in number of elements and return it as an i32.
*/
static LLVMValueRef get_buffer_size(
struct lp_build_tgsi_context *bld_base,
LLVMValueRef descriptor)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef size =
LLVMBuildExtractElement(builder, descriptor,
LLVMConstInt(ctx->i32, 2, 0), "");
if (ctx->screen->b.chip_class == VI) {
/* On VI, the descriptor contains the size in bytes,
* but TXQ must return the size in elements.
* The stride is always non-zero for resources using TXQ.
*/
@@ -3296,21 +3291,21 @@ image_fetch_rsrc(
if (dcc_off && target != TGSI_TEXTURE_BUFFER)
*rsrc = force_dcc_off(ctx, *rsrc);
}
static LLVMValueRef image_fetch_coords(
struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_instruction *inst,
unsigned src)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
unsigned target = inst->Memory.Texture;
unsigned num_coords = tgsi_util_get_texture_coord_dim(target);
LLVMValueRef coords[4];
LLVMValueRef tmp;
int chan;
for (chan = 0; chan < num_coords; ++chan) {
tmp = lp_build_emit_fetch(bld_base, inst, src, chan);
tmp = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
@@ -3409,21 +3404,21 @@ static void buffer_append_args(
i1true : i1false; /* glc */
}
emit_data->args[emit_data->arg_count++] = i1false; /* slc */
}
static void load_fetch_args(
struct lp_build_tgsi_context * bld_base,
struct lp_build_emit_data * emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_instruction * inst = emit_data->inst;
unsigned target = inst->Memory.Texture;
LLVMValueRef rsrc;
emit_data->dst_type = ctx->v4f32;
if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef offset;
LLVMValueRef tmp;
@@ -3611,21 +3606,21 @@ static bool is_oneway_access_only(const struct
tgsi_full_instruction *inst,
}
return false;
}
static void load_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
const struct tgsi_shader_info *info = &ctx->shader->selector->info;
char intrinsic_name[64];
bool readonly_memory = false;
if (inst->Src[0].Register.File == TGSI_FILE_MEMORY) {
load_emit_memory(ctx, emit_data);
return;
}
@@ -3664,21 +3659,21 @@ static void load_emit(
emit_data->args, emit_data->arg_count,
get_load_intr_attribs(readonly_memory));
}
}
static void store_fetch_args(
struct lp_build_tgsi_context * bld_base,
struct lp_build_emit_data * emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
struct tgsi_full_src_register memory;
LLVMValueRef chans[4];
LLVMValueRef data;
LLVMValueRef rsrc;
unsigned chan;
emit_data->dst_type = LLVMVoidTypeInContext(gallivm->context);
@@ -3828,21 +3823,21 @@ static void store_emit_memory(
LLVMBuildStore(builder, data, derived_ptr);
}
}
static void store_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
const struct tgsi_shader_info *info = &ctx->shader->selector->info;
unsigned target = inst->Memory.Texture;
char intrinsic_name[64];
bool writeonly_memory = false;
if (inst->Dst[0].Register.File == TGSI_FILE_MEMORY) {
store_emit_memory(ctx, emit_data);
return;
@@ -3881,21 +3876,21 @@ static void store_emit(
emit_data->args, emit_data->arg_count,
get_store_intr_attribs(writeonly_memory));
}
}
static void atomic_fetch_args(
struct lp_build_tgsi_context * bld_base,
struct lp_build_emit_data * emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
LLVMValueRef data1, data2;
LLVMValueRef rsrc;
LLVMValueRef tmp;
emit_data->dst_type = ctx->f32;
tmp = lp_build_emit_fetch(bld_base, inst, 2, 0);
data1 = LLVMBuildBitCast(builder, tmp, ctx->i32, "");
@@ -4009,21 +4004,21 @@ static void atomic_emit_memory(struct si_shader_context
*ctx,
}
emit_data->output[emit_data->chan] = LLVMBuildBitCast(builder, result,
emit_data->dst_type, "");
}
static void atomic_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
char intrinsic_name[40];
LLVMValueRef tmp;
if (inst->Src[0].Register.File == TGSI_FILE_MEMORY) {
atomic_emit_memory(ctx, emit_data);
return;
}
@@ -4146,21 +4141,21 @@ static void resq_fetch_args(
0xf);
}
}
static void resq_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
const struct tgsi_full_instruction *inst = emit_data->inst;
LLVMValueRef out;
if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
out = LLVMBuildExtractElement(builder, emit_data->args[0],
LLVMConstInt(ctx->i32, 2, 0), "");
} else if (inst->Memory.Texture == TGSI_TEXTURE_BUFFER) {
out = get_buffer_size(bld_base, emit_data->args[0]);
} else {
@@ -4347,21 +4342,21 @@ static void txq_emit(const struct lp_build_tgsi_action
*action,
LLVMValueRef result = ac_build_image_opcode(&ctx->ac, &args);
emit_data->output[emit_data->chan] = fix_resinfo(ctx, target, result);
}
static void tex_fetch_args(
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_instruction *inst = emit_data->inst;
unsigned opcode = inst->Instruction.Opcode;
unsigned target = inst->Texture.Texture;
LLVMValueRef coords[5], derivs[6];
LLVMValueRef address[16];
unsigned num_coords = tgsi_util_get_texture_coord_dim(target);
int ref_pos = tgsi_util_get_shadow_ref_src_index(target);
unsigned count = 0;
unsigned chan;
unsigned num_deriv_channels = 0;
@@ -4873,21 +4868,21 @@ static void build_tex_intrinsic(const struct
lp_build_tgsi_action *action,
emit_data->output[emit_data->chan] =
ac_build_image_opcode(&ctx->ac, &args);
}
static void si_llvm_emit_txqs(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef res, samples;
LLVMValueRef res_ptr, samp_ptr, fmask_ptr = NULL;
tex_fetch_ptrs(bld_base, emit_data, &res_ptr, &samp_ptr, &fmask_ptr);
/* Read the samples from the descriptor directly. */
res = LLVMBuildBitCast(builder, res_ptr, ctx->v8i32, "");
samples = LLVMBuildExtractElement(
@@ -4902,21 +4897,21 @@ static void si_llvm_emit_txqs(
emit_data->output[emit_data->chan] = samples;
}
static void si_llvm_emit_ddxy(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
unsigned opcode = emit_data->info->opcode;
LLVMValueRef val;
int idx;
unsigned mask;
if (opcode == TGSI_OPCODE_DDX_FINE)
mask = AC_TID_MASK_LEFT;
else if (opcode == TGSI_OPCODE_DDY_FINE)
mask = AC_TID_MASK_TOP;
else
@@ -4934,40 +4929,40 @@ static void si_llvm_emit_ddxy(
/*
* this takes an I,J coordinate pair,
* and works out the X and Y derivatives.
* it returns DDX(I), DDX(J), DDY(I), DDY(J).
*/
static LLVMValueRef si_llvm_emit_ddxy_interp(
struct lp_build_tgsi_context *bld_base,
LLVMValueRef interp_ij)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef result[4], a;
unsigned i;
for (i = 0; i < 2; i++) {
a = LLVMBuildExtractElement(gallivm->builder, interp_ij,
LLVMConstInt(ctx->i32, i, 0), "");
result[i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDX,
a);
result[2+i] = lp_build_emit_llvm_unary(bld_base,
TGSI_OPCODE_DDY, a);
}
return lp_build_gather_values(gallivm, result, 4);
}
static void interp_fetch_args(
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_instruction *inst = emit_data->inst;
if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_OFFSET) {
/* offset is in second src, first two channels */
emit_data->args[0] = lp_build_emit_fetch(bld_base,
emit_data->inst, 1,
TGSI_CHAN_X);
emit_data->args[1] = lp_build_emit_fetch(bld_base,
emit_data->inst, 1,
TGSI_CHAN_Y);
@@ -4998,21 +4993,21 @@ static void interp_fetch_args(
emit_data->arg_count = 2;
}
}
static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef interp_param;
const struct tgsi_full_instruction *inst = emit_data->inst;
int input_index = inst->Src[0].Register.Index;
int chan;
int i;
LLVMValueRef attr_number;
LLVMValueRef params = LLVMGetParam(ctx->main_fn, SI_PARAM_PRIM_MASK);
int interp_param_idx;
unsigned interp = shader->selector->info.input_interpolate[input_index];
unsigned location;
@@ -5063,21 +5058,21 @@ static void build_interp_intrinsic(const struct
lp_build_tgsi_action *action,
ctx->f32, "");
temp1 = LLVMBuildFMul(gallivm->builder, ddx_el,
emit_data->args[0], "");
temp1 = LLVMBuildFAdd(gallivm->builder, temp1,
interp_el, "");
temp2 = LLVMBuildFMul(gallivm->builder, ddy_el,
emit_data->args[1], "");
ij_out[i] = LLVMBuildFAdd(gallivm->builder, temp2,
temp1, "");
}
- interp_param = lp_build_gather_values(bld_base->base.gallivm,
ij_out, 2);
+ interp_param = lp_build_gather_values(gallivm, ij_out, 2);
}
for (chan = 0; chan < 4; chan++) {
LLVMValueRef llvm_chan;
unsigned schan;
schan = tgsi_util_get_full_src_register_swizzle(&inst->Src[0],
chan);
llvm_chan = LLVMConstInt(ctx->i32, schan, 0);
if (interp_param) {
@@ -5194,21 +5189,21 @@ static unsigned si_llvm_get_stream(struct
lp_build_tgsi_context *bld_base,
/* Emit one vertex from the geometry shader */
static void si_llvm_emit_vertex(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct lp_build_context *uint = &bld_base->uint_bld;
struct si_shader *shader = ctx->shader;
struct tgsi_shader_info *info = &shader->selector->info;
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct lp_build_if_state if_state;
LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
SI_PARAM_GS2VS_OFFSET);
LLVMValueRef gs_next_vertex;
LLVMValueRef can_emit, kill;
unsigned chan, offset;
int i;
unsigned stream;
stream = si_llvm_get_stream(bld_base, emit_data);
@@ -5294,21 +5289,21 @@ static void si_llvm_emit_primitive(
stream = si_llvm_get_stream(bld_base, emit_data);
ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS |
(stream << 8),
LLVMGetParam(ctx->main_fn, SI_PARAM_GS_WAVE_ID));
}
static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
/* SI only (thanks to a hw bug workaround):
* The real barrier instruction isn’t needed, because an entire patch
* always fits into a single wave.
*/
if (HAVE_LLVM >= 0x0309 &&
ctx->screen->b.chip_class == SI &&
ctx->type == PIPE_SHADER_TESS_CTRL) {
emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
return;
@@ -5453,21 +5448,21 @@ static unsigned si_get_max_workgroup_size(struct
si_shader *shader)
* compile it for the maximum possible group size.
*/
max_work_group_size = SI_MAX_VARIABLE_THREADS_PER_BLOCK;
}
return max_work_group_size;
}
static void create_function(struct si_shader_context *ctx)
{
struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_shader *shader = ctx->shader;
LLVMTypeRef params[SI_NUM_PARAMS + SI_MAX_ATTRIBS], v3i32;
LLVMTypeRef returns[16+32*4];
unsigned i, last_sgpr, num_params, num_return_sgprs;
unsigned num_returns = 0;
unsigned num_prolog_vgprs = 0;
v3i32 = LLVMVectorType(ctx->i32, 3);
params[SI_PARAM_RW_BUFFERS] = const_array(ctx->v16i8,
SI_NUM_RW_BUFFERS);
@@ -5718,21 +5713,21 @@ static void create_function(struct si_shader_context
*ctx)
ctx->type == PIPE_SHADER_TESS_CTRL)
declare_tess_lds(ctx);
}
/**
* Load ESGS and GSVS ring buffer resource descriptors and save the variables
* for later use.
*/
static void preload_ring_buffers(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
SI_PARAM_RW_BUFFERS);
if ((ctx->type == PIPE_SHADER_VERTEX &&
ctx->shader->key.as_es) ||
(ctx->type == PIPE_SHADER_TESS_EVAL &&
ctx->shader->key.as_es) ||
ctx->type == PIPE_SHADER_GEOMETRY) {
@@ -5820,22 +5815,21 @@ static void preload_ring_buffers(struct
si_shader_context *ctx)
ctx->gsvs_ring[stream] = ring;
}
}
}
static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
LLVMValueRef param_rw_buffers,
unsigned param_pos_fixed_pt)
{
- struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef slot, desc, offset, row, bit, address[2];
/* Use the fixed-point gl_FragCoord input.
* Since the stipple pattern is 32x32 and it repeats, just get 5 bits
* per coordinate to get the repeating effect.
*/
address[0] = unpack_param(ctx, param_pos_fixed_pt, 0, 5);
address[1] = unpack_param(ctx, param_pos_fixed_pt, 16, 5);
@@ -6436,28 +6430,28 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
LLVMBuildBr(builder, end_bb);
}
LLVMPositionBuilderAtEnd(builder, end_bb);
LLVMBuildRetVoid(gallivm->builder);
/* Dump LLVM IR before any optimization passes */
if (sscreen->b.debug_flags & DBG_PREOPT_IR &&
r600_can_dump_shader(&sscreen->b, PIPE_SHADER_GEOMETRY))
- ac_dump_module(bld_base->base.gallivm->module);
+ ac_dump_module(ctx.gallivm.module);
si_llvm_finalize_module(&ctx,
r600_extra_shader_checks(&sscreen->b, PIPE_SHADER_GEOMETRY));
r = si_compile_llvm(sscreen, &ctx.shader->binary,
&ctx.shader->config, ctx.tm,
- bld_base->base.gallivm->module,
+ ctx.gallivm.module,
debug, PIPE_SHADER_GEOMETRY,
"GS Copy Shader");
if (!r) {
if (r600_can_dump_shader(&sscreen->b, PIPE_SHADER_GEOMETRY))
fprintf(stderr, "GS Copy Shader:\n");
si_shader_dump(sscreen, ctx.shader, debug,
PIPE_SHADER_GEOMETRY, stderr, true);
r = si_shader_binary_upload(sscreen, ctx.shader);
}
@@ -6859,21 +6853,21 @@ static bool si_compile_tgsi_main(struct
si_shader_context *ctx,
return false;
}
create_function(ctx);
preload_ring_buffers(ctx);
if (ctx->type == PIPE_SHADER_GEOMETRY) {
int i;
for (i = 0; i < 4; i++) {
ctx->gs_next_vertex[i] =
- lp_build_alloca(bld_base->base.gallivm,
+ lp_build_alloca(&ctx->gallivm,
ctx->i32, "");
}
}
if (!lp_build_tgsi_llvm(bld_base, sel->tokens)) {
fprintf(stderr, "Failed to translate shader from TGSI to
LLVM\n");
return false;
}
si_llvm_build_ret(ctx, ctx->return_value);
@@ -7339,41 +7333,39 @@ static void si_build_wrapper_function(struct
si_shader_context *ctx,
}
int si_compile_tgsi_shader(struct si_screen *sscreen,
LLVMTargetMachineRef tm,
struct si_shader *shader,
bool is_monolithic,
struct pipe_debug_callback *debug)
{
struct si_shader_selector *sel = shader->selector;
struct si_shader_context ctx;
- struct lp_build_tgsi_context *bld_base;
LLVMModuleRef mod;
int r = -1;
/* Dump TGSI code before doing TGSI->LLVM conversion in case the
* conversion fails. */
if (r600_can_dump_shader(&sscreen->b, sel->info.processor) &&
!(sscreen->b.debug_flags & DBG_NO_TGSI)) {
tgsi_dump(sel->tokens, 0);
si_dump_streamout(&sel->so);
}
si_init_shader_ctx(&ctx, sscreen, shader, tm);
ctx.separate_prolog = !is_monolithic;
memset(shader->info.vs_output_param_offset, EXP_PARAM_UNDEFINED,
sizeof(shader->info.vs_output_param_offset));
shader->info.uses_instanceid = sel->info.uses_instanceid;
- bld_base = &ctx.bld_base;
ctx.load_system_value = declare_system_value;
if (!si_compile_tgsi_main(&ctx, shader)) {
si_llvm_dispose(&ctx);
return -1;
}
if (is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
LLVMValueRef parts[3];
bool need_prolog;
@@ -7452,21 +7444,21 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
parts[0] = ctx.main_fn;
}
si_get_ps_epilog_key(shader, &epilog_key);
si_build_ps_epilog_function(&ctx, &epilog_key);
parts[need_prolog ? 2 : 1] = ctx.main_fn;
si_build_wrapper_function(&ctx, parts, need_prolog ? 3 : 2,
need_prolog ? 1 : 0);
}
- mod = bld_base->base.gallivm->module;
+ mod = ctx.gallivm.module;
/* Dump LLVM IR before any optimization passes */
if (sscreen->b.debug_flags & DBG_PREOPT_IR &&
r600_can_dump_shader(&sscreen->b, ctx.type))
ac_dump_module(mod);
si_llvm_finalize_module(&ctx,
r600_extra_shader_checks(&sscreen->b,
ctx.type));
/* Post-optimization transformations and analysis. */
diff --git a/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
b/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
index d7ec9ec..1e2d75d 100644
--- a/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
+++ b/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
@@ -494,21 +494,21 @@ static void emit_bfi(const struct lp_build_tgsi_action
*action,
lp_build_const_int32(gallivm, 32), "");
emit_data->output[emit_data->chan] =
LLVMBuildSelect(builder, cond, emit_data->args[1], bfi_sm5, "");
}
static void emit_bfe(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef bfe_sm5;
LLVMValueRef cond;
bfe_sm5 = ac_build_bfe(&ctx->ac, emit_data->args[0],
emit_data->args[1], emit_data->args[2],
emit_data->info->opcode == TGSI_OPCODE_IBFE);
/* Correct for GLSL semantics. */
cond = LLVMBuildICmp(builder, LLVMIntUGE, emit_data->args[2],
@@ -690,21 +690,21 @@ static void emit_up2h(const struct lp_build_tgsi_action
*action,
}
}
static void emit_fdiv(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
emit_data->output[emit_data->chan] =
- LLVMBuildFDiv(bld_base->base.gallivm->builder,
+ LLVMBuildFDiv(ctx->gallivm.builder,
emit_data->args[0], emit_data->args[1], "");
/* Use v_rcp_f32 instead of precise division. */
if (HAVE_LLVM >= 0x0309 &&
!LLVMIsConstant(emit_data->output[emit_data->chan]))
LLVMSetMetadata(emit_data->output[emit_data->chan],
ctx->fpmath_md_kind, ctx->fpmath_md_2p5_ulp);
}
/* 1/sqrt is translated to rsq for f32 if fp32 denormals are not enabled in
diff --git a/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
b/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
index 341c18d..3e38f0d 100644
--- a/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
+++ b/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
@@ -418,21 +418,21 @@ get_array_range(struct lp_build_tgsi_context *bld_base,
range.First = 0;
range.Last = bld_base->info->file_max[File];
return range;
}
static LLVMValueRef
emit_array_index(struct si_shader_context *ctx,
const struct tgsi_ind_register *reg,
unsigned offset)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
if (!reg) {
return LLVMConstInt(ctx->i32, offset, 0);
}
LLVMValueRef addr = LLVMBuildLoad(gallivm->builder,
ctx->addrs[reg->Index][reg->Swizzle], "");
return LLVMBuildAdd(gallivm->builder, addr, LLVMConstInt(ctx->i32,
offset, 0), "");
}
/**
* For indirect registers, construct a pointer directly to the requested
@@ -443,21 +443,21 @@ emit_array_index(struct si_shader_context *ctx,
*/
static LLVMValueRef
get_pointer_into_array(struct si_shader_context *ctx,
unsigned file,
unsigned swizzle,
unsigned reg_index,
const struct tgsi_ind_register *reg_indirect)
{
unsigned array_id;
struct tgsi_array_info *array;
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef idxs[2];
LLVMValueRef index;
LLVMValueRef alloca;
if (file != TGSI_FILE_TEMPORARY)
return NULL;
array_id = get_temp_array_id(&ctx->bld_base, reg_index, reg_indirect);
if (!array_id)
@@ -526,21 +526,21 @@ si_llvm_emit_fetch_64bit(struct lp_build_tgsi_context
*bld_base,
}
static LLVMValueRef
emit_array_fetch(struct lp_build_tgsi_context *bld_base,
unsigned File, enum tgsi_opcode_type type,
struct tgsi_declaration_range range,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
unsigned i, size = range.Last - range.First + 1;
LLVMTypeRef vec = LLVMVectorType(tgsi2llvmtype(bld_base, type), size);
LLVMValueRef result = LLVMGetUndef(vec);
struct tgsi_full_src_register tmp_reg = {};
tmp_reg.Register.File = File;
for (i = 0; i < size; ++i) {
tmp_reg.Register.Index = i + range.First;
@@ -553,21 +553,21 @@ emit_array_fetch(struct lp_build_tgsi_context *bld_base,
static LLVMValueRef
load_value_from_array(struct lp_build_tgsi_context *bld_base,
unsigned file,
enum tgsi_opcode_type type,
unsigned swizzle,
unsigned reg_index,
const struct tgsi_ind_register *reg_indirect)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef ptr;
ptr = get_pointer_into_array(ctx, file, swizzle, reg_index,
reg_indirect);
if (ptr) {
LLVMValueRef val = LLVMBuildLoad(builder, ptr, "");
if (tgsi_type_is_64bit(type)) {
LLVMValueRef ptr_hi, val_hi;
ptr_hi = LLVMBuildGEP(builder, ptr, &ctx->i32_1, 1, "");
val_hi = LLVMBuildLoad(builder, ptr_hi, "");
@@ -588,21 +588,21 @@ load_value_from_array(struct lp_build_tgsi_context
*bld_base,
static void
store_value_to_array(struct lp_build_tgsi_context *bld_base,
LLVMValueRef value,
unsigned file,
unsigned chan_index,
unsigned reg_index,
const struct tgsi_ind_register *reg_indirect)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef ptr;
ptr = get_pointer_into_array(ctx, file, chan_index, reg_index,
reg_indirect);
if (ptr) {
LLVMBuildStore(builder, value, ptr);
} else {
unsigned i, size;
struct tgsi_declaration_range range = get_array_range(bld_base,
file, reg_index, reg_indirect);
LLVMValueRef index = emit_array_index(ctx, reg_indirect,
reg_index - range.First);
@@ -657,30 +657,30 @@ get_output_ptr(struct lp_build_tgsi_context *bld_base,
unsigned index,
assert(index <= ctx->bld_base.info->file_max[TGSI_FILE_OUTPUT]);
return ctx->outputs[index][chan];
}
LLVMValueRef si_llvm_emit_fetch(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_src_register *reg,
enum tgsi_opcode_type type,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
LLVMValueRef result = NULL, ptr, ptr2;
if (swizzle == ~0) {
LLVMValueRef values[TGSI_NUM_CHANNELS];
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
values[chan] = si_llvm_emit_fetch(bld_base, reg, type,
chan);
}
- return lp_build_gather_values(bld_base->base.gallivm, values,
+ return lp_build_gather_values(&ctx->gallivm, values,
TGSI_NUM_CHANNELS);
}
if (reg->Register.Indirect) {
LLVMValueRef load = load_value_from_array(bld_base,
reg->Register.File, type,
swizzle, reg->Register.Index, ®->Indirect);
return bitcast(bld_base, type, load);
}
switch(reg->Register.File) {
@@ -755,35 +755,35 @@ LLVMValueRef si_llvm_emit_fetch(struct
lp_build_tgsi_context *bld_base,
return bitcast(bld_base, type, result);
}
static LLVMValueRef fetch_system_value(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_src_register *reg,
enum tgsi_opcode_type type,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef cval = ctx->system_values[reg->Register.Index];
if (LLVMGetTypeKind(LLVMTypeOf(cval)) == LLVMVectorTypeKind) {
cval = LLVMBuildExtractElement(gallivm->builder, cval,
LLVMConstInt(ctx->i32, swizzle,
0), "");
}
return bitcast(bld_base, type, cval);
}
static void emit_declaration(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_declaration *decl)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
unsigned first, last, i;
switch(decl->Declaration.File) {
case TGSI_FILE_ADDRESS:
{
unsigned idx;
for (idx = decl->Range.First; idx <= decl->Range.Last; idx++) {
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
ctx->addrs[idx][chan] = lp_build_alloca_undef(
&ctx->gallivm,
@@ -846,40 +846,40 @@ static void emit_declaration(struct lp_build_tgsi_context
*bld_base,
ctx->temps_count =
bld_base->info->file_max[TGSI_FILE_TEMPORARY] + 1;
ctx->temps = MALLOC(TGSI_NUM_CHANNELS *
ctx->temps_count * sizeof(LLVMValueRef));
}
if (!array_alloca) {
for (i = 0; i < decl_size; ++i) {
#ifdef DEBUG
snprintf(name, sizeof(name), "TEMP%d.%c",
first + i / 4, "xyzw"[i % 4]);
#endif
ctx->temps[first * TGSI_NUM_CHANNELS + i] =
-
lp_build_alloca_undef(bld_base->base.gallivm,
+ lp_build_alloca_undef(&ctx->gallivm,
ctx->f32,
name);
}
} else {
LLVMValueRef idxs[2] = {
ctx->i32_0,
NULL
};
unsigned j = 0;
if (writemask != TGSI_WRITEMASK_XYZW &&
!ctx->undef_alloca) {
/* Create a dummy alloca. We use it so that we
* have a pointer that is safe to load from if
* a shader ever reads from a channel that
* it never writes to.
*/
ctx->undef_alloca = lp_build_alloca_undef(
- bld_base->base.gallivm,
+ &ctx->gallivm,
ctx->f32, "undef");
}
for (i = 0; i < decl_size; ++i) {
LLVMValueRef ptr;
if (writemask & (1 << (i % 4))) {
#ifdef DEBUG
snprintf(name, sizeof(name),
"TEMP%d.%c",
first + i / 4, "xyzw"[i % 4]);
#endif
@@ -953,23 +953,23 @@ static void emit_declaration(struct lp_build_tgsi_context
*bld_base,
break;
}
}
void si_llvm_emit_store(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_instruction *inst,
const struct tgsi_opcode_info *info,
LLVMValueRef dst[4])
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_dst_register *reg = &inst->Dst[0];
- LLVMBuilderRef builder = ctx->bld_base.base.gallivm->builder;
+ LLVMBuilderRef builder = ctx->gallivm.builder;
LLVMValueRef temp_ptr, temp_ptr2 = NULL;
unsigned chan, chan_index;
bool is_vec_store = false;
enum tgsi_opcode_type dtype =
tgsi_opcode_infer_dst_type(inst->Instruction.Opcode);
if (dst[0]) {
LLVMTypeKind k = LLVMGetTypeKind(LLVMTypeOf(dst[0]));
is_vec_store = (k == LLVMVectorTypeKind);
}
@@ -1084,112 +1084,112 @@ static void emit_default_branch(LLVMBuilderRef
builder, LLVMBasicBlockRef target
{
if (!LLVMGetBasicBlockTerminator(LLVMGetInsertBlock(builder)))
LLVMBuildBr(builder, target);
}
static void bgnloop_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = push_flow(ctx);
flow->loop_entry_block = append_basic_block(ctx, "LOOP");
flow->next_block = append_basic_block(ctx, "ENDLOOP");
set_basicblock_name(flow->loop_entry_block, "loop", bld_base->pc);
LLVMBuildBr(gallivm->builder, flow->loop_entry_block);
LLVMPositionBuilderAtEnd(gallivm->builder, flow->loop_entry_block);
}
static void brk_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = get_innermost_loop(ctx);
LLVMBuildBr(gallivm->builder, flow->next_block);
}
static void cont_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = get_innermost_loop(ctx);
LLVMBuildBr(gallivm->builder, flow->loop_entry_block);
}
static void else_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *current_branch = get_current_flow(ctx);
LLVMBasicBlockRef endif_block;
assert(!current_branch->loop_entry_block);
endif_block = append_basic_block(ctx, "ENDIF");
emit_default_branch(gallivm->builder, endif_block);
LLVMPositionBuilderAtEnd(gallivm->builder, current_branch->next_block);
set_basicblock_name(current_branch->next_block, "else", bld_base->pc);
current_branch->next_block = endif_block;
}
static void endif_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *current_branch = get_current_flow(ctx);
assert(!current_branch->loop_entry_block);
emit_default_branch(gallivm->builder, current_branch->next_block);
LLVMPositionBuilderAtEnd(gallivm->builder, current_branch->next_block);
set_basicblock_name(current_branch->next_block, "endif", bld_base->pc);
ctx->flow_depth--;
}
static void endloop_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *current_loop = get_current_flow(ctx);
assert(current_loop->loop_entry_block);
emit_default_branch(gallivm->builder, current_loop->loop_entry_block);
LLVMPositionBuilderAtEnd(gallivm->builder, current_loop->next_block);
set_basicblock_name(current_loop->next_block, "endloop", bld_base->pc);
ctx->flow_depth--;
}
static void if_cond_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data,
LLVMValueRef cond)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = push_flow(ctx);
LLVMBasicBlockRef if_block;
if_block = append_basic_block(ctx, "IF");
flow->next_block = append_basic_block(ctx, "ELSE");
set_basicblock_name(if_block, "if", bld_base->pc);
LLVMBuildCondBr(gallivm->builder, cond, if_block, flow->next_block);
LLVMPositionBuilderAtEnd(gallivm->builder, if_block);
}
@@ -1382,21 +1382,21 @@ void si_llvm_create_func(struct si_shader_context *ctx,
main_fn_type = LLVMFunctionType(ret_type, ParamTypes, ParamCount, 0);
ctx->main_fn = LLVMAddFunction(ctx->gallivm.module, name, main_fn_type);
main_fn_body = LLVMAppendBasicBlockInContext(ctx->gallivm.context,
ctx->main_fn, "main_body");
LLVMPositionBuilderAtEnd(ctx->gallivm.builder, main_fn_body);
}
void si_llvm_finalize_module(struct si_shader_context *ctx,
bool run_verifier)
{
- struct gallivm_state *gallivm = ctx->bld_base.base.gallivm;
+ struct gallivm_state *gallivm = &ctx->gallivm;
const char *triple = LLVMGetTarget(gallivm->module);
LLVMTargetLibraryInfoRef target_library_info;
/* Create the pass manager */
gallivm->passmgr = LLVMCreatePassManager();
target_library_info = gallivm_create_target_library_info(triple);
LLVMAddTargetLibraryInfo(target_library_info, gallivm->passmgr);
if (run_verifier)
@@ -1417,22 +1417,22 @@ void si_llvm_finalize_module(struct si_shader_context
*ctx,
/* Run the pass */
LLVMRunPassManager(gallivm->passmgr, ctx->gallivm.module);
LLVMDisposeBuilder(gallivm->builder);
LLVMDisposePassManager(gallivm->passmgr);
gallivm_dispose_target_library_info(target_library_info);
}
void si_llvm_dispose(struct si_shader_context *ctx)
{
- LLVMDisposeModule(ctx->bld_base.base.gallivm->module);
- LLVMContextDispose(ctx->bld_base.base.gallivm->context);
+ LLVMDisposeModule(ctx->gallivm.module);
+ LLVMContextDispose(ctx->gallivm.context);
FREE(ctx->temp_arrays);
ctx->temp_arrays = NULL;
FREE(ctx->temp_array_allocas);
ctx->temp_array_allocas = NULL;
FREE(ctx->temps);
ctx->temps = NULL;
ctx->temps_count = 0;
FREE(ctx->imms);
ctx->imms = NULL;
ctx->imms_num = 0;
--
2.7.4
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
