From: Marek Olšák <marek.ol...@amd.com>
---
src/gallium/drivers/radeonsi/si_shader.c | 88 ++++++++++------------
.../drivers/radeonsi/si_shader_tgsi_setup.c | 14 ++--
2 files changed, 47 insertions(+), 55 deletions(-)
diff --git a/src/gallium/drivers/radeonsi/si_shader.c
b/src/gallium/drivers/radeonsi/si_shader.c
index 21efd9a..a5d370b 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -538,38 +538,38 @@ static void declare_input_vs(
break;
}
}
static LLVMValueRef get_primitive_id(struct lp_build_tgsi_context *bld_base,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
if (swizzle > 0)
- return bld_base->uint_bld.zero;
+ return ctx->i32_0;
switch (ctx->type) {
case PIPE_SHADER_VERTEX:
return LLVMGetParam(ctx->main_fn,
ctx->param_vs_prim_id);
case PIPE_SHADER_TESS_CTRL:
return LLVMGetParam(ctx->main_fn,
SI_PARAM_PATCH_ID);
case PIPE_SHADER_TESS_EVAL:
return LLVMGetParam(ctx->main_fn,
ctx->param_tes_patch_id);
case PIPE_SHADER_GEOMETRY:
return LLVMGetParam(ctx->main_fn,
SI_PARAM_PRIMITIVE_ID);
default:
assert(0);
- return bld_base->uint_bld.zero;
+ return ctx->i32_0;
}
}
/**
* 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)
@@ -1096,28 +1096,28 @@ static LLVMValueRef fetch_input_gs(
vtx_offset_param += SI_PARAM_VTX2_OFFSET - 2;
}
vtx_offset = lp_build_mul_imm(uint,
LLVMGetParam(ctx->main_fn,
vtx_offset_param),
4);
param = si_shader_io_get_unique_index(semantic_name, semantic_index);
soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle) * 256, 0);
- value = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1, uint->zero,
+ value = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1, ctx->i32_0,
vtx_offset, soffset, 0, 1, 0, true);
if (tgsi_type_is_64bit(type)) {
LLVMValueRef value2;
soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle + 1) *
256, 0);
value2 = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1,
- uint->zero, vtx_offset, soffset,
+ ctx->i32_0, vtx_offset, soffset,
0, 1, 0, true);
return si_llvm_emit_fetch_64bit(bld_base, type,
value, value2);
}
return LLVMBuildBitCast(gallivm->builder,
value,
tgsi2llvmtype(bld_base, type), "");
}
static int lookup_interp_param_index(unsigned interpolate, unsigned location)
@@ -1169,21 +1169,20 @@ static void interp_fs_input(struct si_shader_context
*ctx,
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 lp_build_context *uint = &bld_base->uint_bld;
struct gallivm_state *gallivm = base->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).
@@ -1198,41 +1197,41 @@ static void interp_fs_input(struct si_shader_context
*ctx,
*/
bool interp = interp_param != NULL;
attr_number = LLVMConstInt(ctx->i32, input_index, 0);
if (interp) {
interp_param = LLVMBuildBitCast(gallivm->builder, interp_param,
LLVMVectorType(ctx->f32, 2),
"");
i = LLVMBuildExtractElement(gallivm->builder, interp_param,
- uint->zero, "");
+ ctx->i32_0, "");
j = LLVMBuildExtractElement(gallivm->builder, interp_param,
- uint->one, "");
+ ctx->i32_1, "");
}
if (semantic_name == TGSI_SEMANTIC_COLOR &&
ctx->shader->key.part.ps.prolog.color_two_side) {
LLVMValueRef is_face_positive;
LLVMValueRef back_attr_number;
/* If BCOLOR0 is used, BCOLOR1 is at offset "num_inputs + 1",
* otherwise it's at offset "num_inputs".
*/
unsigned back_attr_offset = num_interp_inputs;
if (semantic_index == 1 && colors_read_mask & 0xf)
back_attr_offset += 1;
back_attr_number = LLVMConstInt(ctx->i32, back_attr_offset, 0);
is_face_positive = LLVMBuildICmp(gallivm->builder, LLVMIntNE,
- face, uint->zero, "");
+ face, ctx->i32_0, "");
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan,
0);
LLVMValueRef front, back;
if (interp) {
front = ac_build_fs_interp(&ctx->ac, llvm_chan,
attr_number, prim_mask,
i, j);
back = ac_build_fs_interp(&ctx->ac, llvm_chan,
@@ -1248,24 +1247,24 @@ static void interp_fs_input(struct si_shader_context
*ctx,
}
result[chan] = LLVMBuildSelect(gallivm->builder,
is_face_positive,
front,
back,
"");
}
} else if (semantic_name == TGSI_SEMANTIC_FOG) {
if (interp) {
- result[0] = ac_build_fs_interp(&ctx->ac, uint->zero,
+ result[0] = ac_build_fs_interp(&ctx->ac, ctx->i32_0,
attr_number, prim_mask,
i, j);
} else {
- result[0] = ac_build_fs_interp_mov(&ctx->ac, uint->zero,
+ result[0] = ac_build_fs_interp_mov(&ctx->ac, ctx->i32_0,
LLVMConstInt(ctx->i32, 2, 0), /* P0 */
attr_number,
prim_mask);
}
result[1] =
result[2] = LLVMConstReal(ctx->f32, 0.0f);
result[3] = LLVMConstReal(ctx->f32, 1.0f);
} else {
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan,
0);
@@ -2306,21 +2305,21 @@ handle_semantic:
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,
ctx->i32, "");
edgeflag_value = lp_build_min(&bld_base->int_bld,
edgeflag_value,
- bld_base->int_bld.one);
+ ctx->i32_1);
/* The LLVM intrinsic expects a float. */
pos_args[1].out[1] =
LLVMBuildBitCast(base->gallivm->builder,
edgeflag_value,
ctx->f32, "");
}
if (shader->selector->info.writes_layer)
pos_args[1].out[2] = layer_value;
@@ -2412,21 +2411,21 @@ static void si_write_tess_factors(struct
lp_build_tgsi_context *bld_base,
si_llvm_emit_barrier(NULL, bld_base, NULL);
/* Do this only for invocation 0, because the tess levels are per-patch,
* not per-vertex.
*
* This can't jump, because invocation 0 executes this. It should
* at least mask out the loads and stores for other invocations.
*/
lp_build_if(&if_ctx, gallivm,
LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
- invocation_id, bld_base->uint_bld.zero, ""));
+ invocation_id, ctx->i32_0, ""));
/* Determine the layout of one tess factor element in the buffer. */
switch (shader->key.part.tcs.epilog.prim_mode) {
case PIPE_PRIM_LINES:
stride = 2; /* 2 dwords, 1 vec2 store */
outer_comps = 2;
inner_comps = 0;
break;
case PIPE_PRIM_TRIANGLES:
stride = 4; /* 4 dwords, 1 vec4 store */
@@ -2493,21 +2492,21 @@ static void si_write_tess_factors(struct
lp_build_tgsi_context *bld_base,
LLVMConstInt(ctx->i32, SI_HS_RING_TESS_FACTOR, 0));
/* Get the offset. */
tf_base = LLVMGetParam(ctx->main_fn,
SI_PARAM_TESS_FACTOR_OFFSET);
byteoffset = LLVMBuildMul(gallivm->builder, rel_patch_id,
LLVMConstInt(ctx->i32, 4 * stride, 0), "");
lp_build_if(&inner_if_ctx, gallivm,
LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
- rel_patch_id, bld_base->uint_bld.zero, ""));
+ rel_patch_id, ctx->i32_0, ""));
/* Store the dynamic HS control word. */
ac_build_buffer_store_dword(&ctx->ac, buffer,
LLVMConstInt(ctx->i32, 0x80000000, 0),
1, ctx->i32_0, tf_base,
0, 1, 0, true, false);
lp_build_endif(&inner_if_ctx);
/* Store the tessellation factors. */
@@ -2575,23 +2574,23 @@ static void si_llvm_emit_tcs_epilogue(struct
lp_build_tgsi_context *bld_base)
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,
- bld_base->uint_bld.zero, "");
+ ctx->i32_0, "");
rw1 = LLVMBuildExtractElement(builder, rw_buffers,
- bld_base->uint_bld.one, "");
+ ctx->i32_1, "");
ret = LLVMBuildInsertValue(builder, ret, rw0, 0, "");
ret = LLVMBuildInsertValue(builder, ret, rw1, 1, "");
/* Tess offchip and factor buffer soffset are after user SGPRs. */
offchip_layout = LLVMGetParam(ctx->main_fn,
SI_PARAM_TCS_OFFCHIP_LAYOUT);
offchip_soffset = LLVMGetParam(ctx->main_fn, ctx->param_oc_lds);
tf_soffset = LLVMGetParam(ctx->main_fn,
SI_PARAM_TESS_FACTOR_OFFSET);
ret = LLVMBuildInsertValue(builder, ret, offchip_layout,
@@ -3314,25 +3313,25 @@ static LLVMValueRef image_fetch_coords(
for (chan = 0; chan < num_coords; ++chan) {
tmp = lp_build_emit_fetch(bld_base, inst, src, chan);
tmp = LLVMBuildBitCast(builder, tmp,
bld_base->uint_bld.elem_type, "");
coords[chan] = tmp;
}
/* 1D textures are allocated and used as 2D on GFX9. */
if (ctx->screen->b.chip_class >= GFX9) {
if (target == TGSI_TEXTURE_1D) {
- coords[1] = bld_base->uint_bld.zero;
+ coords[1] = ctx->i32_0;
num_coords++;
} else if (target == TGSI_TEXTURE_1D_ARRAY) {
coords[2] = coords[1];
- coords[1] = bld_base->uint_bld.zero;
+ coords[1] = ctx->i32_0;
}
}
if (num_coords == 1)
return coords[0];
if (num_coords == 3) {
/* LLVM has difficulties lowering 3-element vectors. */
coords[3] = bld_base->uint_bld.undef;
num_coords = 4;
@@ -3427,31 +3426,31 @@ static void load_fetch_args(
if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef offset;
LLVMValueRef tmp;
rsrc = shader_buffer_fetch_rsrc(ctx, &inst->Src[0]);
tmp = lp_build_emit_fetch(bld_base, inst, 1, 0);
offset = LLVMBuildBitCast(builder, tmp,
bld_base->uint_bld.elem_type, "");
- buffer_append_args(ctx, emit_data, rsrc,
bld_base->uint_bld.zero,
+ buffer_append_args(ctx, emit_data, rsrc, ctx->i32_0,
offset, false, false);
} else if (inst->Src[0].Register.File == TGSI_FILE_IMAGE) {
LLVMValueRef coords;
image_fetch_rsrc(bld_base, &inst->Src[0], false, target, &rsrc);
coords = image_fetch_coords(bld_base, inst, 1);
if (target == TGSI_TEXTURE_BUFFER) {
buffer_append_args(ctx, emit_data, rsrc, coords,
- bld_base->uint_bld.zero, false,
false);
+ ctx->i32_0, false, false);
} else {
emit_data->args[0] = coords;
emit_data->args[1] = rsrc;
emit_data->args[2] = LLVMConstInt(ctx->i32, 15, 0); /*
dmask */
emit_data->arg_count = 3;
image_append_args(ctx, emit_data, target, false, false);
}
}
}
@@ -3695,61 +3694,60 @@ static void store_fetch_args(
if (inst->Dst[0].Register.File == TGSI_FILE_BUFFER) {
LLVMValueRef offset;
LLVMValueRef tmp;
rsrc = shader_buffer_fetch_rsrc(ctx, &memory);
tmp = lp_build_emit_fetch(bld_base, inst, 0, 0);
offset = LLVMBuildBitCast(builder, tmp,
bld_base->uint_bld.elem_type, "");
- buffer_append_args(ctx, emit_data, rsrc,
bld_base->uint_bld.zero,
+ buffer_append_args(ctx, emit_data, rsrc, ctx->i32_0,
offset, false, false);
} else if (inst->Dst[0].Register.File == TGSI_FILE_IMAGE) {
unsigned target = inst->Memory.Texture;
LLVMValueRef coords;
/* 8bit/16bit TC L1 write corruption bug on SI.
* All store opcodes not aligned to a dword are affected.
*
* The only way to get unaligned stores in radeonsi is through
* shader images.
*/
bool force_glc = ctx->screen->b.chip_class == SI;
coords = image_fetch_coords(bld_base, inst, 0);
if (target == TGSI_TEXTURE_BUFFER) {
image_fetch_rsrc(bld_base, &memory, true, target,
&rsrc);
buffer_append_args(ctx, emit_data, rsrc, coords,
- bld_base->uint_bld.zero, false,
force_glc);
+ ctx->i32_0, false, force_glc);
} else {
emit_data->args[1] = coords;
image_fetch_rsrc(bld_base, &memory, true, target,
&emit_data->args[2]);
emit_data->args[3] = LLVMConstInt(ctx->i32, 15, 0); /*
dmask */
emit_data->arg_count = 4;
image_append_args(ctx, emit_data, target, false,
force_glc);
}
}
}
static void store_emit_buffer(
struct si_shader_context *ctx,
struct lp_build_emit_data *emit_data,
bool writeonly_memory)
{
const struct tgsi_full_instruction *inst = emit_data->inst;
struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = gallivm->builder;
- struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
LLVMValueRef base_data = emit_data->args[0];
LLVMValueRef base_offset = emit_data->args[3];
unsigned writemask = inst->Dst[0].Register.WriteMask;
while (writemask) {
int start, count;
const char *intrinsic_name;
LLVMValueRef data;
LLVMValueRef offset;
LLVMValueRef tmp;
@@ -3767,27 +3765,27 @@ static void store_emit_buffer(
data = base_data;
intrinsic_name = "llvm.amdgcn.buffer.store.v4f32";
} else if (count == 2) {
LLVMTypeRef v2f32 = LLVMVectorType(ctx->f32, 2);
tmp = LLVMBuildExtractElement(
builder, base_data,
LLVMConstInt(ctx->i32, start, 0), "");
data = LLVMBuildInsertElement(
builder, LLVMGetUndef(v2f32), tmp,
- uint_bld->zero, "");
+ ctx->i32_0, "");
tmp = LLVMBuildExtractElement(
builder, base_data,
LLVMConstInt(ctx->i32, start + 1, 0), "");
data = LLVMBuildInsertElement(
- builder, data, tmp, uint_bld->one, "");
+ builder, data, tmp, ctx->i32_1, "");
intrinsic_name = "llvm.amdgcn.buffer.store.v2f32";
} else {
assert(count == 1);
data = LLVMBuildExtractElement(
builder, base_data,
LLVMConstInt(ctx->i32, start, 0), "");
intrinsic_name = "llvm.amdgcn.buffer.store.f32";
}
@@ -3917,32 +3915,32 @@ static void atomic_fetch_args(
emit_data->args[emit_data->arg_count++] = data1;
if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
LLVMValueRef offset;
rsrc = shader_buffer_fetch_rsrc(ctx, &inst->Src[0]);
tmp = lp_build_emit_fetch(bld_base, inst, 1, 0);
offset = LLVMBuildBitCast(builder, tmp,
bld_base->uint_bld.elem_type, "");
- buffer_append_args(ctx, emit_data, rsrc,
bld_base->uint_bld.zero,
+ buffer_append_args(ctx, emit_data, rsrc, ctx->i32_0,
offset, true, false);
} else if (inst->Src[0].Register.File == TGSI_FILE_IMAGE) {
unsigned target = inst->Memory.Texture;
LLVMValueRef coords;
image_fetch_rsrc(bld_base, &inst->Src[0], true, target, &rsrc);
coords = image_fetch_coords(bld_base, inst, 1);
if (target == TGSI_TEXTURE_BUFFER) {
buffer_append_args(ctx, emit_data, rsrc, coords,
- bld_base->uint_bld.zero, true,
false);
+ ctx->i32_0, true, false);
} else {
emit_data->args[emit_data->arg_count++] = coords;
emit_data->args[emit_data->arg_count++] = rsrc;
image_append_args(ctx, emit_data, target, true, false);
}
}
}
static void atomic_emit_memory(struct si_shader_context *ctx,
@@ -4139,21 +4137,21 @@ static void resq_fetch_args(
unsigned image_target;
if (inst->Memory.Texture == TGSI_TEXTURE_3D)
image_target = TGSI_TEXTURE_2D_ARRAY;
else
image_target = inst->Memory.Texture;
image_fetch_rsrc(bld_base, reg, false, inst->Memory.Texture,
&res_ptr);
set_tex_fetch_args(ctx, emit_data, image_target,
- res_ptr, NULL, &bld_base->uint_bld.zero, 1,
+ res_ptr, NULL, &ctx->i32_0, 1,
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);
@@ -4372,21 +4370,21 @@ static void tex_fetch_args(
bool has_offset = inst->Texture.NumOffsets > 0;
LLVMValueRef res_ptr, samp_ptr, fmask_ptr = NULL;
unsigned dmask = 0xf;
tex_fetch_ptrs(bld_base, emit_data, &res_ptr, &samp_ptr, &fmask_ptr);
if (target == TGSI_TEXTURE_BUFFER) {
emit_data->dst_type = ctx->v4f32;
emit_data->args[0] = LLVMBuildBitCast(gallivm->builder, res_ptr,
ctx->v16i8, "");
- emit_data->args[1] = bld_base->uint_bld.zero;
+ emit_data->args[1] = ctx->i32_0;
emit_data->args[2] = lp_build_emit_fetch(bld_base,
emit_data->inst, 0, TGSI_CHAN_X);
emit_data->arg_count = 3;
return;
}
/* Fetch and project texture coordinates */
coords[3] = lp_build_emit_fetch(bld_base, emit_data->inst, 0,
TGSI_CHAN_W);
for (chan = 0; chan < 3; chan++ ) {
coords[chan] = lp_build_emit_fetch(bld_base,
emit_data->inst, 0,
@@ -4538,21 +4536,21 @@ static void tex_fetch_args(
address[count++] = coords[1];
if (num_coords > 2)
address[count++] = coords[2];
/* 1D textures are allocated and used as 2D on GFX9. */
if (ctx->screen->b.chip_class >= GFX9) {
LLVMValueRef filler;
/* Use 0.5, so that we don't sample the border color. */
if (opcode == TGSI_OPCODE_TXF)
- filler = bld_base->uint_bld.zero;
+ filler = ctx->i32_0;
else
filler = LLVMConstReal(ctx->f32, 0.5);
if (target == TGSI_TEXTURE_1D ||
target == TGSI_TEXTURE_SHADOW1D) {
address[count++] = filler;
} else if (target == TGSI_TEXTURE_1D_ARRAY ||
target == TGSI_TEXTURE_SHADOW1D_ARRAY) {
address[count] = address[count - 1];
address[count - 1] = filler;
@@ -4585,21 +4583,20 @@ static void tex_fetch_args(
* For example, 0x11111100 means there are only 2 samples stored and
* the second sample covers 3/4 of the pixel. When reading samples 0
* and 1, return physical sample 0 (determined by the first two 0s
* in FMASK), otherwise return physical sample 1.
*
* The sample index should be adjusted as follows:
* sample_index = (fmask >> (sample_index * 4)) & 0xF;
*/
if (target == TGSI_TEXTURE_2D_MSAA ||
target == TGSI_TEXTURE_2D_ARRAY_MSAA) {
- struct lp_build_context *uint_bld = &bld_base->uint_bld;
struct lp_build_emit_data txf_emit_data = *emit_data;
LLVMValueRef txf_address[4];
/* We only need .xy for non-arrays, and .xyz for arrays. */
unsigned txf_count = target == TGSI_TEXTURE_2D_MSAA ? 2 : 3;
struct tgsi_full_instruction inst = {};
memcpy(txf_address, address, sizeof(txf_address));
/* Read FMASK using TXF_LZ. */
inst.Instruction.Opcode = TGSI_OPCODE_TXF_LZ;
@@ -4612,21 +4609,21 @@ static void tex_fetch_args(
build_tex_intrinsic(&tex_action, bld_base, &txf_emit_data);
/* Initialize some constants. */
LLVMValueRef four = LLVMConstInt(ctx->i32, 4, 0);
LLVMValueRef F = LLVMConstInt(ctx->i32, 0xF, 0);
/* Apply the formula. */
LLVMValueRef fmask =
LLVMBuildExtractElement(gallivm->builder,
txf_emit_data.output[0],
- uint_bld->zero, "");
+ ctx->i32_0, "");
unsigned sample_chan = txf_count; /* the sample index is last */
LLVMValueRef sample_index4 =
LLVMBuildMul(gallivm->builder, address[sample_chan],
four, "");
LLVMValueRef shifted_fmask =
LLVMBuildLShr(gallivm->builder, fmask, sample_index4,
"");
LLVMValueRef final_sample =
@@ -4634,25 +4631,25 @@ static void tex_fetch_args(
/* Don't rewrite the sample index if WORD1.DATA_FORMAT of the
FMASK
* resource descriptor is 0 (invalid),
*/
LLVMValueRef fmask_desc =
LLVMBuildBitCast(gallivm->builder, fmask_ptr,
ctx->v8i32, "");
LLVMValueRef fmask_word1 =
LLVMBuildExtractElement(gallivm->builder, fmask_desc,
- uint_bld->one, "");
+ ctx->i32_1, "");
LLVMValueRef word1_is_nonzero =
LLVMBuildICmp(gallivm->builder, LLVMIntNE,
- fmask_word1, uint_bld->zero, "");
+ fmask_word1, ctx->i32_0, "");
/* Replace the MSAA sample index. */
address[sample_chan] =
LLVMBuildSelect(gallivm->builder, word1_is_nonzero,
final_sample, address[sample_chan], "");
}
if (opcode == TGSI_OPCODE_TXF ||
opcode == TGSI_OPCODE_TXF_LZ) {
/* add tex offsets */
@@ -4748,22 +4745,21 @@ static void si_lower_gather4_integer(struct
si_shader_context *ctx,
half_texel[0] = half_texel[1] = LLVMConstReal(ctx->f32, -0.5);
} else {
struct tgsi_full_instruction txq_inst = {};
struct lp_build_emit_data txq_emit_data = {};
/* Query the texture size. */
txq_inst.Texture.Texture = target;
txq_emit_data.inst = &txq_inst;
txq_emit_data.dst_type = ctx->v4i32;
set_tex_fetch_args(ctx, &txq_emit_data, target,
- args->resource, NULL,
- &ctx->bld_base.uint_bld.zero,
+ args->resource, NULL, &ctx->i32_0,
1, 0xf);
txq_emit(NULL, &ctx->bld_base, &txq_emit_data);
/* Compute -0.5 / size. */
for (c = 0; c < 2; c++) {
half_texel[c] =
LLVMBuildExtractElement(builder,
txq_emit_data.output[0],
LLVMConstInt(ctx->i32,
c, 0), "");
half_texel[c] = LLVMBuildUIToFP(builder, half_texel[c],
ctx->f32, "");
half_texel[c] =
@@ -5005,21 +5001,20 @@ static void interp_fetch_args(
}
}
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 lp_build_context *uint = &bld_base->uint_bld;
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;
@@ -5084,23 +5079,23 @@ static void build_interp_intrinsic(const struct
lp_build_tgsi_action *action,
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) {
interp_param = LLVMBuildBitCast(gallivm->builder,
interp_param, LLVMVectorType(ctx->f32, 2), "");
LLVMValueRef i = LLVMBuildExtractElement(
- gallivm->builder, interp_param, uint->zero, "");
+ gallivm->builder, interp_param, ctx->i32_0, "");
LLVMValueRef j = LLVMBuildExtractElement(
- gallivm->builder, interp_param, uint->one, "");
+ gallivm->builder, interp_param, ctx->i32_1, "");
emit_data->output[chan] = ac_build_fs_interp(&ctx->ac,
llvm_chan, attr_number, params,
i, j);
} else {
emit_data->output[chan] =
ac_build_fs_interp_mov(&ctx->ac,
LLVMConstInt(ctx->i32, 2, 0), /* P0 */
llvm_chan, attr_number, params);
}
}
}
@@ -5433,25 +5428,23 @@ static unsigned llvm_get_type_size(LLVMTypeRef type)
llvm_get_type_size(LLVMGetElementType(type));
default:
assert(0);
return 0;
}
}
static void declare_tess_lds(struct si_shader_context *ctx)
{
struct gallivm_state *gallivm = &ctx->gallivm;
- struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
- struct lp_build_context *uint = &bld_base->uint_bld;
unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768;
- ctx->lds = LLVMBuildIntToPtr(gallivm->builder, uint->zero,
+ ctx->lds = LLVMBuildIntToPtr(gallivm->builder, ctx->i32_0,
LLVMPointerType(LLVMArrayType(ctx->i32, lds_size / 4),
LOCAL_ADDR_SPACE),
"tess_lds");
}
static unsigned si_get_max_workgroup_size(struct si_shader *shader)
{
const unsigned *properties = shader->selector->info.properties;
unsigned max_work_group_size =
properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] *
properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT] *
@@ -5754,21 +5747,20 @@ static void preload_ring_buffers(struct
si_shader_context *ctx)
ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
}
if (ctx->shader->is_gs_copy_shader) {
LLVMValueRef offset = LLVMConstInt(ctx->i32, SI_RING_GSVS, 0);
ctx->gsvs_ring[0] =
ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
} else if (ctx->type == PIPE_SHADER_GEOMETRY) {
const struct si_shader_selector *sel = ctx->shader->selector;
- struct lp_build_context *uint = &ctx->bld_base.uint_bld;
LLVMValueRef offset = LLVMConstInt(ctx->i32, SI_RING_GSVS, 0);
LLVMValueRef base_ring;
base_ring = ac_build_indexed_load_const(&ctx->ac, buf_ptr,
offset);
/* The conceptual layout of the GSVS ring is
* v0c0 .. vLv0 v0c1 .. vLc1 ..
* but the real memory layout is swizzled across
* threads:
* t0v0c0 .. t15v0c0 t0v1c0 .. t15v1c0 ... t15vLcL
@@ -5789,34 +5781,34 @@ static void preload_ring_buffers(struct
si_shader_context *ctx)
continue;
stride = 4 * num_components * sel->gs_max_out_vertices;
/* Limit on the stride field for <= CIK. */
assert(stride < (1 << 14));
num_records = 64;
ring = LLVMBuildBitCast(builder, base_ring, v2i64, "");
- tmp = LLVMBuildExtractElement(builder, ring,
uint->zero, "");
+ tmp = LLVMBuildExtractElement(builder, ring,
ctx->i32_0, "");
tmp = LLVMBuildAdd(builder, tmp,
LLVMConstInt(ctx->i64,
stream_offset, 0), "");
stream_offset += stride * 64;
- ring = LLVMBuildInsertElement(builder, ring, tmp,
uint->zero, "");
+ ring = LLVMBuildInsertElement(builder, ring, tmp,
ctx->i32_0, "");
ring = LLVMBuildBitCast(builder, ring, ctx->v4i32, "");
- tmp = LLVMBuildExtractElement(builder, ring, uint->one,
"");
+ tmp = LLVMBuildExtractElement(builder, ring,
ctx->i32_1, "");
tmp = LLVMBuildOr(builder, tmp,
LLVMConstInt(ctx->i32,
S_008F04_STRIDE(stride) |
S_008F04_SWIZZLE_ENABLE(1), 0),
"");
- ring = LLVMBuildInsertElement(builder, ring, tmp,
uint->one, "");
+ ring = LLVMBuildInsertElement(builder, ring, tmp,
ctx->i32_1, "");
ring = LLVMBuildInsertElement(builder, ring,
LLVMConstInt(ctx->i32, num_records, 0),
LLVMConstInt(ctx->i32, 2, 0), "");
ring = LLVMBuildInsertElement(builder, ring,
LLVMConstInt(ctx->i32,
S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) |
S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
@@ -6371,21 +6363,21 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
LLVMValueRef voffset =
lp_build_mul_imm(uint, LLVMGetParam(ctx.main_fn,
ctx.param_vertex_id), 4);
/* Fetch the vertex stream ID.*/
LLVMValueRef stream_id;
if (gs_selector->so.num_outputs)
stream_id = unpack_param(&ctx, ctx.param_streamout_config, 24,
2);
else
- stream_id = uint->zero;
+ stream_id = ctx.i32_0;
/* Fill in output information. */
for (i = 0; i < gsinfo->num_outputs; ++i) {
outputs[i].semantic_name = gsinfo->output_semantic_name[i];
outputs[i].semantic_index = gsinfo->output_semantic_index[i];
for (int chan = 0; chan < 4; chan++) {
outputs[i].vertex_stream[chan] =
(gsinfo->output_streams[i] >> (2 * chan)) & 3;
}
@@ -6421,21 +6413,21 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
continue;
}
LLVMValueRef soffset = LLVMConstInt(ctx.i32,
offset *
gs_selector->gs_max_out_vertices * 16 * 4, 0);
offset++;
outputs[i].values[chan] =
ac_build_buffer_load(&ctx.ac,
ctx.gsvs_ring[0],
1,
- uint->zero,
voffset,
+ ctx.i32_0, voffset,
soffset, 0, 1, 1,
true);
}
}
/* Streamout and exports. */
if (gs_selector->so.num_outputs) {
si_llvm_emit_streamout(&ctx, outputs,
gsinfo->num_outputs,
stream);
}
diff --git a/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
b/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
index f576a5e..3442a4b 100644
--- a/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
+++ b/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
@@ -491,21 +491,21 @@ get_pointer_into_array(struct si_shader_context *ctx,
index = LLVMBuildMul(
builder, index,
LLVMConstInt(ctx->i32, util_bitcount(array->writemask), 0),
"");
index = LLVMBuildAdd(
builder, index,
LLVMConstInt(ctx->i32,
util_bitcount(array->writemask & ((1 << swizzle) -
1)), 0),
"");
- idxs[0] = ctx->bld_base.uint_bld.zero;
+ idxs[0] = ctx->i32_0;
idxs[1] = index;
return LLVMBuildGEP(builder, alloca, idxs, 2, "");
}
LLVMValueRef
si_llvm_emit_fetch_64bit(struct lp_build_tgsi_context *bld_base,
enum tgsi_opcode_type type,
LLVMValueRef ptr,
LLVMValueRef ptr2)
{
@@ -562,21 +562,21 @@ load_value_from_array(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;
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,
&bld_base->uint_bld.one, 1, "");
+ ptr_hi = LLVMBuildGEP(builder, ptr, &ctx->i32_1, 1, "");
val_hi = LLVMBuildLoad(builder, ptr_hi, "");
val = si_llvm_emit_fetch_64bit(bld_base, type, val,
val_hi);
}
return val;
} else {
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);
@@ -683,24 +683,24 @@ LLVMValueRef si_llvm_emit_fetch(struct
lp_build_tgsi_context *bld_base,
return bitcast(bld_base, type, load);
}
switch(reg->Register.File) {
case TGSI_FILE_IMMEDIATE: {
LLVMTypeRef ctype = tgsi2llvmtype(bld_base, type);
if (tgsi_type_is_64bit(type)) {
result =
LLVMGetUndef(LLVMVectorType(LLVMIntTypeInContext(bld_base->base.gallivm->context,
32), bld_base->base.type.length * 2));
result = LLVMConstInsertElement(result,
ctx->imms[reg->Register.Index * TGSI_NUM_CHANNELS + swizzle],
- bld_base->int_bld.zero);
+ ctx->i32_0);
result = LLVMConstInsertElement(result,
ctx->imms[reg->Register.Index * TGSI_NUM_CHANNELS + swizzle + 1],
- bld_base->int_bld.one);
+ ctx->i32_1);
return LLVMConstBitCast(result, ctype);
} else {
return LLVMConstBitCast(ctx->imms[reg->Register.Index *
TGSI_NUM_CHANNELS + swizzle], ctype);
}
}
case TGSI_FILE_INPUT: {
unsigned index = reg->Register.Index;
LLVMValueRef input[4];
@@ -852,21 +852,21 @@ static void emit_declaration(struct lp_build_tgsi_context
*bld_base,
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,
bld_base->base.vec_type,
name);
}
} else {
LLVMValueRef idxs[2] = {
- bld_base->uint_bld.zero,
+ 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.
@@ -1028,23 +1028,23 @@ void si_llvm_emit_store(struct lp_build_tgsi_context
*bld_base,
default:
return;
}
if (!tgsi_type_is_64bit(dtype))
LLVMBuildStore(builder, value, temp_ptr);
else {
LLVMValueRef ptr = LLVMBuildBitCast(builder,
value,
LLVMVectorType(LLVMIntTypeInContext(bld_base->base.gallivm->context, 32), 2),
"");
LLVMValueRef val2;
value = LLVMBuildExtractElement(builder, ptr,
-
bld_base->uint_bld.zero, "");
+ ctx->i32_0, "");
val2 = LLVMBuildExtractElement(builder, ptr,
-
bld_base->uint_bld.one, "");
+ ctx->i32_1, "");
LLVMBuildStore(builder, bitcast(bld_base,
TGSI_TYPE_FLOAT, value), temp_ptr);
LLVMBuildStore(builder, bitcast(bld_base,
TGSI_TYPE_FLOAT, val2), temp_ptr2);
}
}
}
}
static void set_basicblock_name(LLVMBasicBlockRef bb, const char *base, int pc)
{
--
2.7.4
_______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
