[Mesa-dev] [PATCH 1/3] radeonsi: use i32_0/1 instead of *int_bld.zero/one in most places
Dave Airlie
airlied at gmail.com
Mon Apr 3 18:50:17 UTC 2017
On 3 April 2017 at 19:52, Marek Olšák <maraeo at gmail.com> wrote:
> From: Marek Olšák <marek.olsak at amd.com>
radv uses i32zero and i32one, it might be nice to be consistent, but I
don't mind which way.
Dave.
>
> ---
> 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 at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
More information about the mesa-dev
mailing list