* USE OR OTHER DEALINGS IN THE SOFTWARE.
*/
-#include "gallivm/lp_bld_const.h"
-#include "gallivm/lp_bld_gather.h"
-#include "gallivm/lp_bld_intr.h"
-#include "gallivm/lp_bld_logic.h"
-#include "gallivm/lp_bld_arit.h"
-#include "gallivm/lp_bld_flow.h"
-#include "gallivm/lp_bld_misc.h"
#include "util/u_memory.h"
#include "util/u_string.h"
#include "tgsi/tgsi_build.h"
#include "tgsi/tgsi_util.h"
#include "tgsi/tgsi_dump.h"
-#include "ac_binary.h"
-#include "ac_llvm_util.h"
#include "ac_exp_param.h"
#include "ac_shader_util.h"
+#include "ac_llvm_util.h"
#include "si_shader_internal.h"
#include "si_pipe.h"
#include "sid.h"
static void si_init_shader_ctx(struct si_shader_context *ctx,
struct si_screen *sscreen,
- struct si_compiler *compiler);
+ struct ac_llvm_compiler *compiler);
static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
* less than 64, so that a 64-bit bitmask of used inputs or outputs can be
* calculated.
*/
-unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index)
+unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index,
+ unsigned is_varying)
{
switch (semantic_name) {
case TGSI_SEMANTIC_POSITION:
return SI_MAX_IO_GENERIC + 6;
case TGSI_SEMANTIC_PRIMID:
return SI_MAX_IO_GENERIC + 7;
- case TGSI_SEMANTIC_COLOR: /* these alias */
- case TGSI_SEMANTIC_BCOLOR:
+ case TGSI_SEMANTIC_COLOR:
assert(index < 2);
return SI_MAX_IO_GENERIC + 8 + index;
+ case TGSI_SEMANTIC_BCOLOR:
+ assert(index < 2);
+ /* If it's a varying, COLOR and BCOLOR alias. */
+ if (is_varying)
+ return SI_MAX_IO_GENERIC + 8 + index;
+ else
+ return SI_MAX_IO_GENERIC + 10 + index;
case TGSI_SEMANTIC_TEXCOORD:
assert(index < 8);
- assert(SI_MAX_IO_GENERIC + 10 + index < 64);
- return SI_MAX_IO_GENERIC + 10 + index;
+ STATIC_ASSERT(SI_MAX_IO_GENERIC + 12 + 8 <= 63);
+ return SI_MAX_IO_GENERIC + 12 + index;
+ case TGSI_SEMANTIC_CLIPVERTEX:
+ return 63;
default:
+ fprintf(stderr, "invalid semantic name = %u\n", semantic_name);
assert(!"invalid semantic name");
return 0;
}
static LLVMValueRef
get_tcs_out_patch0_offset(struct si_shader_context *ctx)
{
- return lp_build_mul_imm(&ctx->bld_base.uint_bld,
- si_unpack_param(ctx,
- ctx->param_tcs_out_lds_offsets,
- 0, 16),
- 4);
+ return LLVMBuildMul(ctx->ac.builder,
+ si_unpack_param(ctx,
+ ctx->param_tcs_out_lds_offsets,
+ 0, 16),
+ LLVMConstInt(ctx->i32, 4, 0), "");
}
static LLVMValueRef
get_tcs_out_patch0_patch_data_offset(struct si_shader_context *ctx)
{
- return lp_build_mul_imm(&ctx->bld_base.uint_bld,
- si_unpack_param(ctx,
- ctx->param_tcs_out_lds_offsets,
- 16, 16),
- 4);
+ return LLVMBuildMul(ctx->ac.builder,
+ si_unpack_param(ctx,
+ ctx->param_tcs_out_lds_offsets,
+ 16, 16),
+ LLVMConstInt(ctx->i32, 4, 0), "");
}
static LLVMValueRef
si_shader_io_get_unique_index_patch(name[input_index],
index[input_index]) :
si_shader_io_get_unique_index(name[input_index],
- index[input_index]);
+ index[input_index], false);
/* Add the base address of the element. */
return LLVMBuildAdd(ctx->ac.builder, base_addr,
param_index_base = is_patch ?
si_shader_io_get_unique_index_patch(name[param_base], index[param_base]) :
- si_shader_io_get_unique_index(name[param_base], index[param_base]);
+ si_shader_io_get_unique_index(name[param_base], index[param_base], false);
if (param_index) {
param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; chan++)
values[chan] = lds_load(bld_base, type, chan, dw_addr);
- return lp_build_gather_values(&ctx->gallivm, values,
+ return ac_build_gather_values(&ctx->ac, values,
TGSI_NUM_CHANNELS);
}
return si_llvm_emit_fetch_64bit(bld_base, type, lo, hi);
}
- dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
- LLVMConstInt(ctx->i32, swizzle, 0));
+ dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+ LLVMConstInt(ctx->i32, swizzle, 0), "");
value = ac_lds_load(&ctx->ac, dw_addr);
unsigned dw_offset_imm, LLVMValueRef dw_addr,
LLVMValueRef value)
{
- dw_addr = lp_build_add(&ctx->bld_base.uint_bld, dw_addr,
- LLVMConstInt(ctx->i32, dw_offset_imm, 0));
+ dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+ LLVMConstInt(ctx->i32, dw_offset_imm, 0), "");
ac_lds_store(&ctx->ac, dw_addr, value);
}
}
if (reg->Register.WriteMask == 0xF && !is_tess_factor) {
- LLVMValueRef value = lp_build_gather_values(&ctx->gallivm,
+ LLVMValueRef value = ac_build_gather_values(&ctx->ac,
values, 4);
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr,
base, 0, 1, 0, true, false);
}
if (writemask == 0xF && !is_tess_factor) {
- LLVMValueRef value = lp_build_gather_values(&ctx->gallivm,
+ LLVMValueRef value = ac_build_gather_values(&ctx->ac,
values, 4);
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, addr,
base, 0, 1, 0, true, false);
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
struct si_shader *shader = ctx->shader;
- struct lp_build_context *uint = &ctx->bld_base.uint_bld;
LLVMValueRef vtx_offset, soffset;
struct tgsi_shader_info *info = &shader->selector->info;
unsigned semantic_name = info->input_semantic_name[input_index];
unsigned param;
LLVMValueRef value;
- param = si_shader_io_get_unique_index(semantic_name, semantic_index);
+ param = si_shader_io_get_unique_index(semantic_name, semantic_index, false);
/* GFX9 has the ESGS ring in LDS. */
if (ctx->screen->info.chip_class >= GFX9) {
values[chan] = si_llvm_load_input_gs(abi, input_index, vtx_offset_param,
type, chan);
}
- return lp_build_gather_values(&ctx->gallivm, values,
+ return ac_build_gather_values(&ctx->ac, values,
TGSI_NUM_CHANNELS);
}
/* Get the vertex offset parameter on GFX6. */
LLVMValueRef gs_vtx_offset = ctx->gs_vtx_offset[vtx_offset_param];
- vtx_offset = lp_build_mul_imm(uint, gs_vtx_offset, 4);
+ vtx_offset = LLVMBuildMul(ctx->ac.builder, gs_vtx_offset,
+ LLVMConstInt(ctx->i32, 4, 0), "");
soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle) * 256, 0);
unsigned input_index,
LLVMValueRef out[4])
{
- struct lp_build_context *base = &ctx->bld_base.base;
struct si_shader *shader = ctx->shader;
struct tgsi_shader_info *info = &shader->selector->info;
LLVMValueRef main_fn = ctx->main_fn;
unsigned mask = colors_read >> (semantic_index * 4);
unsigned offset = SI_PARAM_POS_FIXED_PT + 1 +
(semantic_index ? util_bitcount(colors_read & 0xf) : 0);
+ LLVMValueRef undef = LLVMGetUndef(ctx->f32);
- out[0] = mask & 0x1 ? LLVMGetParam(main_fn, offset++) : base->undef;
- out[1] = mask & 0x2 ? LLVMGetParam(main_fn, offset++) : base->undef;
- out[2] = mask & 0x4 ? LLVMGetParam(main_fn, offset++) : base->undef;
- out[3] = mask & 0x8 ? LLVMGetParam(main_fn, offset++) : base->undef;
+ out[0] = mask & 0x1 ? LLVMGetParam(main_fn, offset++) : undef;
+ out[1] = mask & 0x2 ? LLVMGetParam(main_fn, offset++) : undef;
+ out[2] = mask & 0x4 ? LLVMGetParam(main_fn, offset++) : undef;
+ out[3] = mask & 0x8 ? LLVMGetParam(main_fn, offset++) : undef;
return;
}
for (i = 0; i < 3; ++i)
values[i] = LLVMConstInt(ctx->i32, sizes[i], 0);
- result = lp_build_gather_values(&ctx->gallivm, values, 3);
+ result = ac_build_gather_values(&ctx->ac, values, 3);
} else {
result = LLVMGetParam(ctx->main_fn, ctx->param_block_size);
}
static LLVMValueRef load_sample_position(struct ac_shader_abi *abi, LLVMValueRef sample_id)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
- struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
LLVMValueRef desc = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
LLVMValueRef buf_index = LLVMConstInt(ctx->i32, SI_PS_CONST_SAMPLE_POSITIONS, 0);
LLVMValueRef resource = ac_build_load_to_sgpr(&ctx->ac, desc, buf_index);
/* offset = sample_id * 8 (8 = 2 floats containing samplepos.xy) */
- LLVMValueRef offset0 = lp_build_mul_imm(uint_bld, sample_id, 8);
+ LLVMValueRef offset0 = LLVMBuildMul(ctx->ac.builder, sample_id, LLVMConstInt(ctx->i32, 8, 0), "");
LLVMValueRef offset1 = LLVMBuildAdd(ctx->ac.builder, offset0, LLVMConstInt(ctx->i32, 4, 0), "");
LLVMValueRef pos[4] = {
LLVMConstReal(ctx->f32, 0)
};
- return lp_build_gather_values(&ctx->gallivm, pos, 4);
+ return ac_build_gather_values(&ctx->ac, pos, 4);
}
static LLVMValueRef load_sample_mask_in(struct ac_shader_abi *abi)
static LLVMValueRef si_load_tess_coord(struct ac_shader_abi *abi)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
- struct lp_build_context *bld = &ctx->bld_base.base;
-
LLVMValueRef coord[4] = {
LLVMGetParam(ctx->main_fn, ctx->param_tes_u),
LLVMGetParam(ctx->main_fn, ctx->param_tes_v),
/* For triangles, the vector should be (u, v, 1-u-v). */
if (ctx->shader->selector->info.properties[TGSI_PROPERTY_TES_PRIM_MODE] ==
- PIPE_PRIM_TRIANGLES)
- coord[2] = lp_build_sub(bld, ctx->ac.f32_1,
- lp_build_add(bld, coord[0], coord[1]));
-
- return lp_build_gather_values(&ctx->gallivm, coord, 4);
+ PIPE_PRIM_TRIANGLES) {
+ coord[2] = LLVMBuildFSub(ctx->ac.builder, ctx->ac.f32_1,
+ LLVMBuildFAdd(ctx->ac.builder,
+ coord[0], coord[1], ""), "");
+ }
+ return ac_build_gather_values(&ctx->ac, coord, 4);
}
static LLVMValueRef load_tess_level(struct si_shader_context *ctx,
LLVMGetParam(ctx->main_fn, SI_PARAM_POS_X_FLOAT),
LLVMGetParam(ctx->main_fn, SI_PARAM_POS_Y_FLOAT),
LLVMGetParam(ctx->main_fn, SI_PARAM_POS_Z_FLOAT),
- lp_build_emit_llvm_unary(&ctx->bld_base, TGSI_OPCODE_RCP,
- LLVMGetParam(ctx->main_fn,
- SI_PARAM_POS_W_FLOAT)),
+ ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
+ LLVMGetParam(ctx->main_fn, SI_PARAM_POS_W_FLOAT)),
};
- value = lp_build_gather_values(&ctx->gallivm, pos, 4);
+ value = ac_build_gather_values(&ctx->ac, pos, 4);
break;
}
LLVMConstReal(ctx->f32, 0),
LLVMConstReal(ctx->f32, 0)
};
- pos[0] = lp_build_emit_llvm_unary(&ctx->bld_base,
- TGSI_OPCODE_FRC, pos[0]);
- pos[1] = lp_build_emit_llvm_unary(&ctx->bld_base,
- TGSI_OPCODE_FRC, pos[1]);
- value = lp_build_gather_values(&ctx->gallivm, pos, 4);
+ pos[0] = ac_build_fract(&ctx->ac, pos[0], 32);
+ pos[1] = ac_build_fract(&ctx->ac, pos[1], 32);
+ value = ac_build_gather_values(&ctx->ac, pos, 4);
break;
}
for (i = 0; i < 4; i++)
val[i] = buffer_load_const(ctx, buf,
LLVMConstInt(ctx->i32, (offset + i) * 4, 0));
- value = lp_build_gather_values(&ctx->gallivm, val, 4);
+ value = ac_build_gather_values(&ctx->ac, val, 4);
break;
}
values[i] = ctx->abi.workgroup_ids[i];
}
}
- value = lp_build_gather_values(&ctx->gallivm, values, 3);
+ value = ac_build_gather_values(&ctx->ac, values, 3);
break;
}
break;
case TGSI_SEMANTIC_HELPER_INVOCATION:
- value = lp_build_intrinsic(ctx->ac.builder,
+ value = ac_build_intrinsic(&ctx->ac,
"llvm.amdgcn.ps.live",
ctx->i1, NULL, 0,
- LP_FUNC_ATTR_READNONE);
+ AC_FUNC_ATTR_READNONE);
value = LLVMBuildNot(ctx->ac.builder, value, "");
value = LLVMBuildSExt(ctx->ac.builder, value, ctx->i32, "");
break;
void si_declare_compute_memory(struct si_shader_context *ctx)
{
struct si_shader_selector *sel = ctx->shader->selector;
+ unsigned lds_size = sel->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE];
LLVMTypeRef i8p = LLVMPointerType(ctx->i8, AC_LOCAL_ADDR_SPACE);
LLVMValueRef var;
assert(!ctx->ac.lds);
var = LLVMAddGlobalInAddressSpace(ctx->ac.module,
- LLVMArrayType(ctx->i8, sel->local_size),
+ LLVMArrayType(ctx->i8, lds_size),
"compute_lds",
AC_LOCAL_ADDR_SPACE);
LLVMSetAlignment(var, 4);
for (chan = 0; chan < TGSI_NUM_CHANNELS; ++chan)
values[chan] = fetch_constant(bld_base, reg, type, chan);
- return lp_build_gather_values(&ctx->gallivm, values, 4);
+ return ac_build_gather_values(&ctx->ac, values, 4);
}
/* Split 64-bit loads. */
samplemask_param);
coverage = ac_to_integer(&ctx->ac, coverage);
- coverage = lp_build_intrinsic(ctx->ac.builder, "llvm.ctpop.i32",
+ coverage = ac_build_intrinsic(&ctx->ac, "llvm.ctpop.i32",
ctx->i32,
- &coverage, 1, LP_FUNC_ATTR_READNONE);
+ &coverage, 1, AC_FUNC_ATTR_READNONE);
coverage = LLVMBuildUIToFP(ctx->ac.builder, coverage,
ctx->f32, "");
base_elt = buffer_load_const(ctx, const_resource,
addr);
args->out[chan] =
- lp_build_add(&ctx->bld_base.base, args->out[chan],
- lp_build_mul(&ctx->bld_base.base, base_elt,
- out_elts[const_chan]));
+ LLVMBuildFAdd(ctx->ac.builder, args->out[chan],
+ LLVMBuildFMul(ctx->ac.builder, base_elt,
+ out_elts[const_chan], ""), "");
}
}
if ((semantic_name != TGSI_SEMANTIC_GENERIC ||
semantic_index < SI_MAX_IO_GENERIC) &&
shader->key.opt.kill_outputs &
- (1ull << si_shader_io_get_unique_index(semantic_name, semantic_index)))
+ (1ull << si_shader_io_get_unique_index(semantic_name,
+ semantic_index, true)))
continue;
si_export_param(ctx, param_count, outputs[i].values);
}
/* Convert the outputs to vectors for stores. */
- vec0 = lp_build_gather_values(&ctx->gallivm, out, MIN2(stride, 4));
+ vec0 = ac_build_gather_values(&ctx->ac, out, MIN2(stride, 4));
vec1 = NULL;
if (stride > 4)
- vec1 = lp_build_gather_values(&ctx->gallivm, out+4, stride - 4);
+ vec1 = ac_build_gather_values(&ctx->ac, out+4, stride - 4);
/* Get the buffer. */
buffer = get_tess_ring_descriptor(ctx, TCS_FACTOR_RING);
tf_outer_offset = get_tcs_tes_buffer_address(ctx, rel_patch_id, NULL,
LLVMConstInt(ctx->i32, param_outer, 0));
- outer_vec = lp_build_gather_values(&ctx->gallivm, outer,
+ outer_vec = ac_build_gather_values(&ctx->ac, outer,
util_next_power_of_two(outer_comps));
ac_build_buffer_store_dword(&ctx->ac, buf, outer_vec,
LLVMConstInt(ctx->i32, param_inner, 0));
inner_vec = inner_comps == 1 ? inner[0] :
- lp_build_gather_values(&ctx->gallivm, inner, inner_comps);
+ ac_build_gather_values(&ctx->ac, inner, inner_comps);
ac_build_buffer_store_dword(&ctx->ac, buf, inner_vec,
inner_comps, tf_inner_offset,
base, 0, 1, 0, true, false);
8 + SI_SGPR_VS_STATE_BITS);
#if !HAVE_32BIT_POINTERS
- ret = si_insert_input_ptr(ctx, ret, ctx->param_vs_state_bits + 1,
+ ret = si_insert_input_ptr(ctx, ret, ctx->param_vs_state_bits + 4,
8 + GFX9_SGPR_2ND_SAMPLERS_AND_IMAGES);
#endif
8 + SI_SGPR_BINDLESS_SAMPLERS_AND_IMAGES);
#if !HAVE_32BIT_POINTERS
- ret = si_insert_input_ptr(ctx, ret, ctx->param_vs_state_bits + 1,
+ ret = si_insert_input_ptr(ctx, ret, ctx->param_vs_state_bits + 4,
8 + GFX9_SGPR_2ND_SAMPLERS_AND_IMAGES);
#endif
name == TGSI_SEMANTIC_VIEWPORT_INDEX)
continue;
- int param = si_shader_io_get_unique_index(name, index);
+ int param = si_shader_io_get_unique_index(name, index, false);
LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
LLVMConstInt(ctx->i32, param * 4, 0), "");
continue;
param = si_shader_io_get_unique_index(info->output_semantic_name[i],
- info->output_semantic_index[i]);
+ info->output_semantic_index[i], false);
for (chan = 0; chan < 4; chan++) {
+ if (!(info->output_usagemask[i] & (1 << chan)))
+ continue;
+
LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], "");
out_val = ac_to_integer(&ctx->ac, out_val);
* an IF statement is added that clamps all colors if the constant
* is true.
*/
- if (ctx->type == PIPE_SHADER_VERTEX) {
- struct lp_build_if_state if_ctx;
- LLVMValueRef cond = NULL;
- LLVMValueRef addr, val;
-
- for (i = 0; i < info->num_outputs; i++) {
- if (info->output_semantic_name[i] != TGSI_SEMANTIC_COLOR &&
- info->output_semantic_name[i] != TGSI_SEMANTIC_BCOLOR)
- continue;
+ struct lp_build_if_state if_ctx;
+ LLVMValueRef cond = NULL;
+ LLVMValueRef addr, val;
- /* We've found a color. */
- if (!cond) {
- /* The state is in the first bit of the user SGPR. */
- cond = LLVMGetParam(ctx->main_fn,
- ctx->param_vs_state_bits);
- cond = LLVMBuildTrunc(ctx->ac.builder, cond,
- ctx->i1, "");
- lp_build_if(&if_ctx, &ctx->gallivm, cond);
- }
+ for (i = 0; i < info->num_outputs; i++) {
+ if (info->output_semantic_name[i] != TGSI_SEMANTIC_COLOR &&
+ info->output_semantic_name[i] != TGSI_SEMANTIC_BCOLOR)
+ continue;
- for (j = 0; j < 4; j++) {
- addr = addrs[4 * i + j];
- val = LLVMBuildLoad(ctx->ac.builder, addr, "");
- val = ac_build_clamp(&ctx->ac, val);
- LLVMBuildStore(ctx->ac.builder, val, addr);
- }
+ /* We've found a color. */
+ if (!cond) {
+ /* The state is in the first bit of the user SGPR. */
+ cond = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_state_bits);
+ cond = LLVMBuildTrunc(ctx->ac.builder, cond,
+ ctx->i1, "");
+ lp_build_if(&if_ctx, &ctx->gallivm, cond);
}
- if (cond)
- lp_build_endif(&if_ctx);
+ for (j = 0; j < 4; j++) {
+ addr = addrs[4 * i + j];
+ val = LLVMBuildLoad(ctx->ac.builder, addr, "");
+ val = ac_build_clamp(&ctx->ac, val);
+ LLVMBuildStore(ctx->ac.builder, val, addr);
+ }
}
+ if (cond)
+ lp_build_endif(&if_ctx);
+
for (i = 0; i < info->num_outputs; i++) {
outputs[i].semantic_name = info->output_semantic_name[i];
outputs[i].semantic_index = info->output_semantic_index[i];
for (i = 0; i < 2; i++) {
a = LLVMBuildExtractElement(ctx->ac.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);
+ result[i] = ac_build_ddxy(&ctx->ac, AC_TID_MASK_TOP_LEFT, 1,
+ ac_to_integer(&ctx->ac, a)); /* DDX */
+ result[2+i] = ac_build_ddxy(&ctx->ac, AC_TID_MASK_TOP_LEFT, 2,
+ ac_to_integer(&ctx->ac, a)); /* DDY */
}
- return lp_build_gather_values(&ctx->gallivm, result, 4);
+ return ac_build_gather_values(&ctx->ac, result, 4);
}
static void interp_fetch_args(
ctx->ac.f32_0,
};
- sample_position = lp_build_gather_values(&ctx->gallivm, center, 4);
+ sample_position = ac_build_gather_values(&ctx->ac, center, 4);
} else {
sample_position = load_sample_position(&ctx->abi, sample_id);
}
ij_out[i] = LLVMBuildFAdd(ctx->ac.builder, temp2, temp1, "");
}
- interp_param = lp_build_gather_values(&ctx->gallivm, ij_out, 2);
+ interp_param = ac_build_gather_values(&ctx->ac, ij_out, 2);
}
if (interp_param)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
struct tgsi_shader_info *info = &ctx->shader->selector->info;
- struct lp_build_context *uint = &ctx->bld_base.uint_bld;
struct si_shader *shader = ctx->shader;
struct lp_build_if_state if_state;
LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
shader->selector->gs_max_out_vertices, 0);
offset++;
- voffset = lp_build_add(uint, voffset, gs_next_vertex);
- voffset = lp_build_mul_imm(uint, voffset, 4);
+ voffset = LLVMBuildAdd(ctx->ac.builder, voffset, gs_next_vertex, "");
+ voffset = LLVMBuildMul(ctx->ac.builder, voffset,
+ LLVMConstInt(ctx->i32, 4, 0), "");
out_val = ac_to_integer(&ctx->ac, out_val);
}
}
- gs_next_vertex = lp_build_add(uint, gs_next_vertex,
- ctx->i32_1);
-
+ gs_next_vertex = LLVMBuildAdd(ctx->ac.builder, gs_next_vertex, ctx->i32_1, "");
LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
/* Signal vertex emission */
return;
}
- lp_build_intrinsic(ctx->ac.builder,
+ ac_build_intrinsic(&ctx->ac,
"llvm.amdgcn.s.barrier",
- ctx->voidt, NULL, 0, LP_FUNC_ATTR_CONVERGENT);
+ ctx->voidt, NULL, 0, AC_FUNC_ATTR_CONVERGENT);
}
static const struct lp_build_tgsi_action interp_action = {
* allows the optimization passes to move loads and reduces
* SGPR spilling significantly.
*/
- lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_INREG);
+ ac_add_function_attr(ctx->ac.context, ctx->main_fn, i + 1,
+ AC_FUNC_ATTR_INREG);
if (LLVMGetTypeKind(LLVMTypeOf(P)) == LLVMPointerTypeKind) {
- lp_add_function_attr(ctx->main_fn, i + 1, LP_FUNC_ATTR_NOALIAS);
+ ac_add_function_attr(ctx->ac.context, ctx->main_fn, i + 1,
+ AC_FUNC_ATTR_NOALIAS);
ac_add_attr_dereferenceable(P, UINT64_MAX);
}
}
static void declare_vs_specific_input_sgprs(struct si_shader_context *ctx,
struct si_function_info *fninfo)
{
+ ctx->param_vs_state_bits = add_arg(fninfo, ARG_SGPR, ctx->i32);
add_arg_assign(fninfo, ARG_SGPR, ctx->i32, &ctx->abi.base_vertex);
add_arg_assign(fninfo, ARG_SGPR, ctx->i32, &ctx->abi.start_instance);
add_arg_assign(fninfo, ARG_SGPR, ctx->i32, &ctx->abi.draw_id);
- ctx->param_vs_state_bits = add_arg(fninfo, ARG_SGPR, ctx->i32);
}
static void declare_vs_input_vgprs(struct si_shader_context *ctx,
/* no extra parameters */
} else {
if (shader->is_gs_copy_shader) {
- fninfo.num_params = ctx->param_rw_buffers + 1;
+ fninfo.num_params = ctx->param_vs_state_bits + 1;
fninfo.num_sgpr_params = fninfo.num_params;
}
if (ctx->type == PIPE_SHADER_VERTEX) {
declare_vs_specific_input_sgprs(ctx, &fninfo);
} else {
+ ctx->param_vs_state_bits = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tes_offchip_addr = add_arg(&fninfo, ARG_SGPR, ctx->i32);
- if (!HAVE_32BIT_POINTERS) {
- /* Declare as many input SGPRs as the VS has. */
+ /* Declare as many input SGPRs as the VS has. */
+ if (!HAVE_32BIT_POINTERS)
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
- ctx->param_vs_state_bits = add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
- }
}
if (!HAVE_32BIT_POINTERS) {
case PIPE_SHADER_TESS_EVAL:
declare_global_desc_pointers(ctx, &fninfo);
declare_per_stage_desc_pointers(ctx, &fninfo, true);
+ ctx->param_vs_state_bits = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tes_offchip_addr = add_arg(&fninfo, ARG_SGPR, ctx->i32);
si_get_max_workgroup_size(shader));
/* Reserve register locations for VGPR inputs the PS prolog may need. */
- if (ctx->type == PIPE_SHADER_FRAGMENT &&
- ctx->separate_prolog) {
+ if (ctx->type == PIPE_SHADER_FRAGMENT && !ctx->shader->is_monolithic) {
ac_llvm_add_target_dep_function_attr(ctx->main_fn,
"InitialPSInputAddr",
S_0286D0_PERSP_SAMPLE_ENA(1) |
static int si_compile_llvm(struct si_screen *sscreen,
struct ac_shader_binary *binary,
struct si_shader_config *conf,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
LLVMModuleRef mod,
struct pipe_debug_callback *debug,
unsigned processor,
/* Generate code for the hardware VS shader stage to go with a geometry shader */
struct si_shader *
si_generate_gs_copy_shader(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader_selector *gs_selector,
struct pipe_debug_callback *debug)
{
struct si_shader_context ctx;
struct si_shader *shader;
LLVMBuilderRef builder;
- struct lp_build_tgsi_context *bld_base = &ctx.bld_base;
- struct lp_build_context *uint = &bld_base->uint_bld;
- struct si_shader_output_values *outputs;
+ struct si_shader_output_values outputs[SI_MAX_VS_OUTPUTS];
struct tgsi_shader_info *gsinfo = &gs_selector->info;
int i, r;
- outputs = MALLOC(gsinfo->num_outputs * sizeof(outputs[0]));
-
- if (!outputs)
- return NULL;
shader = CALLOC_STRUCT(si_shader);
- if (!shader) {
- FREE(outputs);
+ if (!shader)
return NULL;
- }
/* We can leave the fence as permanently signaled because the GS copy
* shader only becomes visible globally after it has been compiled. */
preload_ring_buffers(&ctx);
LLVMValueRef voffset =
- lp_build_mul_imm(uint, ctx.abi.vertex_id, 4);
+ LLVMBuildMul(ctx.ac.builder, ctx.abi.vertex_id,
+ LLVMConstInt(ctx.i32, 4, 0), "");
/* Fetch the vertex stream ID.*/
LLVMValueRef stream_id;
for (unsigned chan = 0; chan < 4; chan++) {
if (!(gsinfo->output_usagemask[i] & (1 << chan)) ||
outputs[i].vertex_stream[chan] != stream) {
- outputs[i].values[chan] = ctx.bld_base.base.undef;
+ outputs[i].values[chan] = LLVMGetUndef(ctx.f32);
continue;
}
stream);
}
- if (stream == 0)
+ if (stream == 0) {
+ /* Vertex color clamping.
+ *
+ * This uses a state constant loaded in a user data SGPR and
+ * an IF statement is added that clamps all colors if the constant
+ * is true.
+ */
+ struct lp_build_if_state if_ctx;
+ LLVMValueRef v[2], cond = NULL;
+ LLVMBasicBlockRef blocks[2];
+
+ for (unsigned i = 0; i < gsinfo->num_outputs; i++) {
+ if (gsinfo->output_semantic_name[i] != TGSI_SEMANTIC_COLOR &&
+ gsinfo->output_semantic_name[i] != TGSI_SEMANTIC_BCOLOR)
+ continue;
+
+ /* We've found a color. */
+ if (!cond) {
+ /* The state is in the first bit of the user SGPR. */
+ cond = LLVMGetParam(ctx.main_fn,
+ ctx.param_vs_state_bits);
+ cond = LLVMBuildTrunc(ctx.ac.builder, cond,
+ ctx.i1, "");
+ lp_build_if(&if_ctx, &ctx.gallivm, cond);
+ /* Remember blocks for Phi. */
+ blocks[0] = if_ctx.true_block;
+ blocks[1] = if_ctx.entry_block;
+ }
+
+ for (unsigned j = 0; j < 4; j++) {
+ /* Insert clamp into the true block. */
+ v[0] = ac_build_clamp(&ctx.ac, outputs[i].values[j]);
+ v[1] = outputs[i].values[j];
+
+ /* Insert Phi into the endif block. */
+ LLVMPositionBuilderAtEnd(ctx.ac.builder, if_ctx.merge_block);
+ outputs[i].values[j] = ac_build_phi(&ctx.ac, ctx.f32, 2, v, blocks);
+ LLVMPositionBuilderAtEnd(ctx.ac.builder, if_ctx.true_block);
+ }
+ }
+ if (cond)
+ lp_build_endif(&if_ctx);
+
si_llvm_export_vs(&ctx, outputs, gsinfo->num_outputs);
+ }
LLVMBuildBr(builder, end_bb);
}
r = si_compile_llvm(sscreen, &ctx.shader->binary,
&ctx.shader->config, ctx.compiler,
- ctx.gallivm.module,
+ ctx.ac.module,
debug, PIPE_SHADER_GEOMETRY,
"GS Copy Shader");
if (!r) {
si_llvm_dispose(&ctx);
- FREE(outputs);
-
if (r != 0) {
FREE(shader);
shader = NULL;
static void si_init_shader_ctx(struct si_shader_context *ctx,
struct si_screen *sscreen,
- struct si_compiler *compiler)
+ struct ac_llvm_compiler *compiler)
{
struct lp_build_tgsi_context *bld_base;
LLVMGetParam(ctx->main_fn, param),
LLVMConstInt(ctx->i32, bitoffset, 0),
};
- lp_build_intrinsic(ctx->ac.builder,
+ ac_build_intrinsic(&ctx->ac,
"llvm.amdgcn.init.exec.from.input",
- ctx->voidt, args, 2, LP_FUNC_ATTR_CONVERGENT);
+ ctx->voidt, args, 2, AC_FUNC_ATTR_CONVERGENT);
}
static bool si_vs_needs_prolog(const struct si_shader_selector *sel,
return sel->vs_needs_prolog || key->ls_vgpr_fix;
}
-static bool si_compile_tgsi_main(struct si_shader_context *ctx,
- bool is_monolithic)
+static bool si_compile_tgsi_main(struct si_shader_context *ctx)
{
struct si_shader *shader = ctx->shader;
struct si_shader_selector *sel = shader->selector;
* if-block together with its prolog in si_build_wrapper_function.
*/
if (ctx->screen->info.chip_class >= GFX9) {
- if (!is_monolithic &&
+ if (!shader->is_monolithic &&
sel->info.num_instructions > 1 && /* not empty shader */
(shader->key.as_es || shader->key.as_ls) &&
(ctx->type == PIPE_SHADER_TESS_EVAL ||
ctx->param_merged_wave_info, 0);
} else if (ctx->type == PIPE_SHADER_TESS_CTRL ||
ctx->type == PIPE_SHADER_GEOMETRY) {
- if (!is_monolithic)
+ if (!shader->is_monolithic)
ac_init_exec_full_mask(&ctx->ac);
- /* The barrier must execute for all shaders in a
- * threadgroup.
- */
- si_llvm_emit_barrier(NULL, bld_base, NULL);
-
LLVMValueRef num_threads = si_unpack_param(ctx, ctx->param_merged_wave_info, 8, 8);
LLVMValueRef ena =
LLVMBuildICmp(ctx->ac.builder, LLVMIntULT,
ac_get_thread_id(&ctx->ac), num_threads, "");
lp_build_if(&ctx->merged_wrap_if_state, &ctx->gallivm, ena);
+
+ /* The barrier must execute for all shaders in a
+ * threadgroup.
+ *
+ * Execute the barrier inside the conditional block,
+ * so that empty waves can jump directly to s_endpgm,
+ * which will also signal the barrier.
+ *
+ * If the shader is TCS and the TCS epilog is present
+ * and contains a barrier, it will wait there and then
+ * reach s_endpgm.
+ */
+ si_llvm_emit_barrier(NULL, bld_base, NULL);
}
}
sel->tcs_info.tessfactors_are_def_in_all_invocs) {
for (unsigned i = 0; i < 6; i++) {
ctx->invoc0_tess_factors[i] =
- lp_build_alloca_undef(&ctx->gallivm, ctx->i32, "");
+ ac_build_alloca_undef(&ctx->ac, ctx->i32, "");
}
}
int i;
for (i = 0; i < 4; i++) {
ctx->gs_next_vertex[i] =
- lp_build_alloca(&ctx->gallivm,
- ctx->i32, "");
+ ac_build_alloca(&ctx->ac, ctx->i32, "");
}
}
if (sel->force_correct_derivs_after_kill) {
- ctx->postponed_kill = lp_build_alloca_undef(&ctx->gallivm, ctx->i1, "");
+ ctx->postponed_kill = ac_build_alloca_undef(&ctx->ac, ctx->i1, "");
/* true = don't kill. */
LLVMBuildStore(ctx->ac.builder, LLVMConstInt(ctx->i1, 1, 0),
ctx->postponed_kill);
si_init_function_info(&fninfo);
for (unsigned i = 0; i < num_parts; ++i) {
- lp_add_function_attr(parts[i], -1, LP_FUNC_ATTR_ALWAYSINLINE);
+ ac_add_function_attr(ctx->ac.context, parts[i], -1,
+ AC_FUNC_ATTR_ALWAYSINLINE);
LLVMSetLinkage(parts[i], LLVMPrivateLinkage);
}
param_size = ac_get_type_size(param_type) / 4;
is_sgpr = ac_is_sgpr_param(param);
- if (is_sgpr)
- lp_add_function_attr(parts[part], param_idx + 1, LP_FUNC_ATTR_INREG);
- else if (out_idx < num_out_sgpr) {
+ if (is_sgpr) {
+ ac_add_function_attr(ctx->ac.context, parts[part],
+ param_idx + 1, AC_FUNC_ATTR_INREG);
+ } else if (out_idx < num_out_sgpr) {
/* Skip returned SGPRs the current part doesn't
* declare on the input. */
out_idx = num_out_sgpr;
if (param_size == 1)
arg = out[out_idx];
else
- arg = lp_build_gather_values(&ctx->gallivm, &out[out_idx], param_size);
+ arg = ac_build_gather_values(&ctx->ac, &out[out_idx], param_size);
if (LLVMTypeOf(arg) != param_type) {
if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
}
int si_compile_tgsi_shader(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader *shader,
- bool is_monolithic,
struct pipe_debug_callback *debug)
{
struct si_shader_selector *sel = shader->selector;
si_init_shader_ctx(&ctx, sscreen, compiler);
si_llvm_context_set_tgsi(&ctx, shader);
- ctx.separate_prolog = !is_monolithic;
memset(shader->info.vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
sizeof(shader->info.vs_output_param_offset));
shader->info.uses_instanceid = sel->info.uses_instanceid;
- if (!si_compile_tgsi_main(&ctx, is_monolithic)) {
+ if (!si_compile_tgsi_main(&ctx)) {
si_llvm_dispose(&ctx);
return -1;
}
- if (is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
+ if (shader->is_monolithic && ctx.type == PIPE_SHADER_VERTEX) {
LLVMValueRef parts[2];
bool need_prolog = sel->vs_needs_prolog;
si_build_wrapper_function(&ctx, parts + !need_prolog,
1 + need_prolog, need_prolog, 0);
- } else if (is_monolithic && ctx.type == PIPE_SHADER_TESS_CTRL) {
+ } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_TESS_CTRL) {
if (sscreen->info.chip_class >= GFX9) {
struct si_shader_selector *ls = shader->key.part.tcs.ls;
LLVMValueRef parts[4];
shader_ls.key.as_ls = 1;
shader_ls.key.mono = shader->key.mono;
shader_ls.key.opt = shader->key.opt;
+ shader_ls.is_monolithic = true;
si_llvm_context_set_tgsi(&ctx, &shader_ls);
- if (!si_compile_tgsi_main(&ctx, true)) {
+ if (!si_compile_tgsi_main(&ctx)) {
si_llvm_dispose(&ctx);
return -1;
}
si_build_wrapper_function(&ctx, parts, 2, 0, 0);
}
- } else if (is_monolithic && ctx.type == PIPE_SHADER_GEOMETRY) {
+ } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_GEOMETRY) {
if (ctx.screen->info.chip_class >= GFX9) {
struct si_shader_selector *es = shader->key.part.gs.es;
LLVMValueRef es_prolog = NULL;
shader_es.key.as_es = 1;
shader_es.key.mono = shader->key.mono;
shader_es.key.opt = shader->key.opt;
+ shader_es.is_monolithic = true;
si_llvm_context_set_tgsi(&ctx, &shader_es);
- if (!si_compile_tgsi_main(&ctx, true)) {
+ if (!si_compile_tgsi_main(&ctx)) {
si_llvm_dispose(&ctx);
return -1;
}
si_build_wrapper_function(&ctx, parts, 2, 1, 0);
}
- } else if (is_monolithic && ctx.type == PIPE_SHADER_FRAGMENT) {
+ } else if (shader->is_monolithic && ctx.type == PIPE_SHADER_FRAGMENT) {
LLVMValueRef parts[3];
union si_shader_part_key prolog_key;
union si_shader_part_key epilog_key;
/* Compile to bytecode. */
r = si_compile_llvm(sscreen, &shader->binary, &shader->config, compiler,
- ctx.gallivm.module, debug, ctx.type, "TGSI shader");
+ ctx.ac.module, debug, ctx.type, "TGSI shader");
si_llvm_dispose(&ctx);
if (r) {
fprintf(stderr, "LLVM failed to compile shader\n");
enum pipe_shader_type type,
bool prolog,
union si_shader_part_key *key,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct pipe_debug_callback *debug,
void (*build)(struct si_shader_context *,
union si_shader_part_key *),
/* Get the pointer to rw buffers. */
ptr[0] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
ptr[1] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS + 1);
- list = lp_build_gather_values(&ctx->gallivm, ptr, 2);
+ list = ac_build_gather_values(&ctx->ac, ptr, 2);
list = LLVMBuildBitCast(ctx->ac.builder, list, ctx->i64, "");
list = LLVMBuildIntToPtr(ctx->ac.builder, list,
ac_array_in_const_addr_space(ctx->v4i32), "");
}
static bool si_get_vs_prolog(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader *shader,
struct pipe_debug_callback *debug,
struct si_shader *main_part,
* Select and compile (or reuse) vertex shader parts (prolog & epilog).
*/
static bool si_shader_select_vs_parts(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
* Select and compile (or reuse) TCS parts (epilog).
*/
static bool si_shader_select_tcs_parts(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
* Select and compile (or reuse) GS parts (prolog).
*/
static bool si_shader_select_gs_parts(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
interp_vgpr, "");
interp[1] = LLVMBuildExtractValue(ctx->ac.builder, ret,
interp_vgpr + 1, "");
- interp_ij = lp_build_gather_values(&ctx->gallivm, interp, 2);
+ interp_ij = ac_build_gather_values(&ctx->ac, interp, 2);
}
/* Use the absolute location of the input. */
* Select and compile (or reuse) pixel shader parts (prolog & epilog).
*/
static bool si_shader_select_ps_parts(struct si_screen *sscreen,
- struct si_compiler *compiler,
+ struct ac_llvm_compiler *compiler,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
}
}
-int si_shader_create(struct si_screen *sscreen, struct si_compiler *compiler,
+int si_shader_create(struct si_screen *sscreen, struct ac_llvm_compiler *compiler,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
/* Monolithic shader (compiled as a whole, has many variants,
* may take a long time to compile).
*/
- r = si_compile_tgsi_shader(sscreen, compiler, shader, true, debug);
+ r = si_compile_tgsi_shader(sscreen, compiler, shader, debug);
if (r)
return r;
} else {
projects / mesa.git / blobdiff