#include "ac_binary.h"
#include "ac_llvm_util.h"
#include "ac_exp_param.h"
+#include "ac_shader_util.h"
#include "si_shader_internal.h"
#include "si_pipe.h"
#include "sid.h"
LOCAL_ADDR_SPACE = 3,
};
+static bool llvm_type_is_64bit(struct si_shader_context *ctx,
+ LLVMTypeRef type)
+{
+ if (type == ctx->ac.i64 || type == ctx->ac.f64)
+ return true;
+
+ return false;
+}
+
static bool is_merged_shader(struct si_shader *shader)
{
if (shader->selector->screen->info.chip_class <= VI)
/**
* Get the value of a shader input parameter and extract a bitfield.
*/
-static LLVMValueRef unpack_param(struct si_shader_context *ctx,
- unsigned param, unsigned rshift,
- unsigned bitwidth)
+static LLVMValueRef unpack_llvm_param(struct si_shader_context *ctx,
+ LLVMValueRef value, unsigned rshift,
+ unsigned bitwidth)
{
- LLVMValueRef value = LLVMGetParam(ctx->main_fn,
- param);
-
if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind)
value = ac_to_integer(&ctx->ac, value);
return value;
}
+static LLVMValueRef unpack_param(struct si_shader_context *ctx,
+ unsigned param, unsigned rshift,
+ unsigned bitwidth)
+{
+ LLVMValueRef value = LLVMGetParam(ctx->main_fn, param);
+
+ return unpack_llvm_param(ctx, value, rshift, bitwidth);
+}
+
static LLVMValueRef get_rel_patch_id(struct si_shader_context *ctx)
{
switch (ctx->type) {
case PIPE_SHADER_TESS_CTRL:
- return unpack_param(ctx, ctx->param_tcs_rel_ids, 0, 8);
+ return unpack_llvm_param(ctx, ctx->abi.tcs_rel_ids, 0, 8);
case PIPE_SHADER_TESS_EVAL:
return LLVMGetParam(ctx->main_fn,
return LLVMGetParam(ctx->main_fn,
ctx->param_vs_prim_id);
case PIPE_SHADER_TESS_CTRL:
- return LLVMGetParam(ctx->main_fn,
- ctx->param_tcs_patch_id);
+ return ctx->abi.tcs_patch_id;
case PIPE_SHADER_TESS_EVAL:
- return LLVMGetParam(ctx->main_fn,
- ctx->param_tes_patch_id);
+ return ctx->abi.tes_patch_id;
case PIPE_SHADER_GEOMETRY:
return ctx->abi.gs_prim_id;
default:
return si_llvm_bound_index(ctx, result, num);
}
+static LLVMValueRef get_dw_address_from_generic_indices(struct si_shader_context *ctx,
+ LLVMValueRef vertex_dw_stride,
+ LLVMValueRef base_addr,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned input_index,
+ ubyte *name,
+ ubyte *index,
+ bool is_patch)
+{
+ if (vertex_dw_stride) {
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
+ LLVMBuildMul(ctx->ac.builder, vertex_index,
+ vertex_dw_stride, ""), "");
+ }
+
+ if (param_index) {
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
+ LLVMBuildMul(ctx->ac.builder, param_index,
+ LLVMConstInt(ctx->i32, 4, 0), ""), "");
+ }
+
+ int param = is_patch ?
+ 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]);
+
+ /* Add the base address of the element. */
+ return LLVMBuildAdd(ctx->ac.builder, base_addr,
+ LLVMConstInt(ctx->i32, param * 4, 0), "");
+}
/**
* Calculate a dword address given an input or output register and a stride.
{
struct tgsi_shader_info *info = &ctx->shader->selector->info;
ubyte *name, *index, *array_first;
- int first, param;
+ int input_index;
struct tgsi_full_dst_register reg;
+ LLVMValueRef vertex_index = NULL;
+ LLVMValueRef ind_index = NULL;
/* Set the register description. The address computation is the same
* for sources and destinations. */
/* If the register is 2-dimensional (e.g. an array of vertices
* in a primitive), calculate the base address of the vertex. */
if (reg.Register.Dimension) {
- LLVMValueRef index;
-
if (reg.Dimension.Indirect)
- index = si_get_indirect_index(ctx, ®.DimIndirect,
+ vertex_index = si_get_indirect_index(ctx, ®.DimIndirect,
1, reg.Dimension.Index);
else
- index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0);
-
- base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
- LLVMBuildMul(ctx->ac.builder, index,
- vertex_dw_stride, ""), "");
+ vertex_index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0);
}
/* Get information about the register. */
if (reg.Register.Indirect) {
/* Add the relative address of the element. */
- LLVMValueRef ind_index;
-
if (reg.Indirect.ArrayID)
- first = array_first[reg.Indirect.ArrayID];
+ input_index = array_first[reg.Indirect.ArrayID];
else
- first = reg.Register.Index;
+ input_index = reg.Register.Index;
ind_index = si_get_indirect_index(ctx, ®.Indirect,
- 1, reg.Register.Index - first);
-
- base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
- LLVMBuildMul(ctx->ac.builder, ind_index,
- LLVMConstInt(ctx->i32, 4, 0), ""), "");
-
- param = reg.Register.Dimension ?
- si_shader_io_get_unique_index(name[first], index[first]) :
- si_shader_io_get_unique_index_patch(name[first], index[first]);
+ 1, reg.Register.Index - input_index);
} else {
- param = reg.Register.Dimension ?
- si_shader_io_get_unique_index(name[reg.Register.Index],
- index[reg.Register.Index]) :
- si_shader_io_get_unique_index_patch(name[reg.Register.Index],
- index[reg.Register.Index]);
+ input_index = reg.Register.Index;
}
- /* Add the base address of the element. */
- return LLVMBuildAdd(ctx->ac.builder, base_addr,
- LLVMConstInt(ctx->i32, param * 4, 0), "");
+ return get_dw_address_from_generic_indices(ctx, vertex_dw_stride,
+ base_addr, vertex_index,
+ ind_index, input_index,
+ name, index,
+ !reg.Register.Dimension);
}
/* The offchip buffer layout for TCS->TES is
return base_addr;
}
+/* This is a generic helper that can be shared by the NIR and TGSI backends */
+static LLVMValueRef get_tcs_tes_buffer_address_from_generic_indices(
+ struct si_shader_context *ctx,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned param_base,
+ ubyte *name,
+ ubyte *index,
+ bool is_patch)
+{
+ unsigned param_index_base;
+
+ 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]);
+
+ if (param_index) {
+ param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
+ LLVMConstInt(ctx->i32, param_index_base, 0),
+ "");
+ } else {
+ param_index = LLVMConstInt(ctx->i32, param_index_base, 0);
+ }
+
+ return get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx),
+ vertex_index, param_index);
+}
+
static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
struct si_shader_context *ctx,
const struct tgsi_full_dst_register *dst,
struct tgsi_full_src_register reg;
LLVMValueRef vertex_index = NULL;
LLVMValueRef param_index = NULL;
- unsigned param_index_base, param_base;
+ unsigned param_base;
reg = src ? *src : tgsi_full_src_register_from_dst(dst);
} else {
param_base = reg.Register.Index;
- param_index = ctx->i32_0;
}
- param_index_base = reg.Register.Dimension ?
- si_shader_io_get_unique_index(name[param_base], index[param_base]) :
- si_shader_io_get_unique_index_patch(name[param_base], index[param_base]);
-
- param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
- LLVMConstInt(ctx->i32, param_index_base, 0),
- "");
-
- return get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx),
- vertex_index, param_index);
+ return get_tcs_tes_buffer_address_from_generic_indices(ctx, vertex_index,
+ param_index, param_base,
+ name, index, !reg.Register.Dimension);
}
static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
- enum tgsi_opcode_type type, unsigned swizzle,
+ LLVMTypeRef type, unsigned swizzle,
LLVMValueRef buffer, LLVMValueRef offset,
LLVMValueRef base, bool can_speculate)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef value, value2;
- LLVMTypeRef llvm_type = tgsi2llvmtype(bld_base, type);
- LLVMTypeRef vec_type = LLVMVectorType(llvm_type, 4);
+ LLVMTypeRef vec_type = LLVMVectorType(type, 4);
if (swizzle == ~0) {
value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
return LLVMBuildBitCast(ctx->ac.builder, value, vec_type, "");
}
- if (!tgsi_type_is_64bit(type)) {
+ if (!llvm_type_is_64bit(ctx, type)) {
value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
0, 1, 0, can_speculate, false);
* \param dw_addr address in dwords
*/
static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base,
- enum tgsi_opcode_type type, unsigned swizzle,
+ LLVMTypeRef type, unsigned swizzle,
LLVMValueRef dw_addr)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
}
/* Split 64-bit loads. */
- if (tgsi_type_is_64bit(type)) {
+ if (llvm_type_is_64bit(ctx, type)) {
LLVMValueRef lo, hi;
- lo = lds_load(bld_base, TGSI_TYPE_UNSIGNED, swizzle, dw_addr);
- hi = lds_load(bld_base, TGSI_TYPE_UNSIGNED, swizzle + 1, dw_addr);
+ lo = lds_load(bld_base, ctx->i32, swizzle, dw_addr);
+ hi = lds_load(bld_base, ctx->i32, swizzle + 1, dw_addr);
return si_llvm_emit_fetch_64bit(bld_base, type, lo, hi);
}
value = ac_lds_load(&ctx->ac, dw_addr);
- return bitcast(bld_base, type, value);
+ return LLVMBuildBitCast(ctx->ac.builder, value, type, "");
}
/**
dw_addr = get_tcs_in_current_patch_offset(ctx);
dw_addr = get_dw_address(ctx, NULL, reg, stride, dw_addr);
- return lds_load(bld_base, type, swizzle, dw_addr);
+ return lds_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle, dw_addr);
+}
+
+static LLVMValueRef si_nir_load_tcs_varyings(struct ac_shader_abi *abi,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned const_index,
+ unsigned location,
+ unsigned driver_location,
+ unsigned component,
+ unsigned num_components,
+ bool is_patch,
+ bool is_compact,
+ bool load_input)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ struct tgsi_shader_info *info = &ctx->shader->selector->info;
+ struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
+ LLVMValueRef dw_addr, stride;
+
+ driver_location = driver_location / 4;
+
+ if (load_input) {
+ stride = get_tcs_in_vertex_dw_stride(ctx);
+ dw_addr = get_tcs_in_current_patch_offset(ctx);
+ } else {
+ if (is_patch) {
+ stride = NULL;
+ dw_addr = get_tcs_out_current_patch_data_offset(ctx);
+ } else {
+ stride = get_tcs_out_vertex_dw_stride(ctx);
+ dw_addr = get_tcs_out_current_patch_offset(ctx);
+ }
+ }
+
+ if (param_index) {
+ /* Add the constant index to the indirect index */
+ param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
+ LLVMConstInt(ctx->i32, const_index, 0), "");
+ } else {
+ param_index = LLVMConstInt(ctx->i32, const_index, 0);
+ }
+
+ dw_addr = get_dw_address_from_generic_indices(ctx, stride, dw_addr,
+ vertex_index, param_index,
+ driver_location,
+ info->input_semantic_name,
+ info->input_semantic_index,
+ is_patch);
+
+ LLVMValueRef value[4];
+ for (unsigned i = 0; i < num_components + component; i++) {
+ value[i] = lds_load(bld_base, ctx->i32, i, dw_addr);
+ }
+
+ return ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
}
static LLVMValueRef fetch_output_tcs(
dw_addr = get_dw_address(ctx, NULL, reg, NULL, dw_addr);
}
- return lds_load(bld_base, type, swizzle, dw_addr);
+ return lds_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle, dw_addr);
}
static LLVMValueRef fetch_input_tes(
base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
addr = get_tcs_tes_buffer_address_from_reg(ctx, NULL, reg);
- return buffer_load(bld_base, type, swizzle, buffer, base, addr, true);
+ return buffer_load(bld_base, tgsi2llvmtype(bld_base, type), swizzle,
+ buffer, base, addr, true);
+}
+
+LLVMValueRef si_nir_load_input_tes(struct ac_shader_abi *abi,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned const_index,
+ unsigned location,
+ unsigned driver_location,
+ unsigned component,
+ unsigned num_components,
+ bool is_patch,
+ bool is_compact,
+ bool load_input)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ struct tgsi_shader_info *info = &ctx->shader->selector->info;
+ LLVMValueRef buffer, base, addr;
+
+ driver_location = driver_location / 4;
+
+ buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k);
+
+ base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
+
+ if (param_index) {
+ /* Add the constant index to the indirect index */
+ param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
+ LLVMConstInt(ctx->i32, const_index, 0), "");
+ } else {
+ param_index = LLVMConstInt(ctx->i32, const_index, 0);
+ }
+
+ addr = get_tcs_tes_buffer_address_from_generic_indices(ctx, vertex_index,
+ param_index, driver_location,
+ info->input_semantic_name,
+ info->input_semantic_index,
+ is_patch);
+
+ /* TODO: This will generate rather ordinary llvm code, although it
+ * should be easy for the optimiser to fix up. In future we might want
+ * to refactor buffer_load(), but for now this maximises code sharing
+ * between the NIR and TGSI backends.
+ */
+ LLVMValueRef value[4];
+ for (unsigned i = component; i < num_components + component; i++) {
+ value[i] = buffer_load(&ctx->bld_base, ctx->i32, i, buffer, base, addr, true);
+ }
+
+ return ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
}
static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
}
}
-static LLVMValueRef fetch_input_gs(
- struct lp_build_tgsi_context *bld_base,
- const struct tgsi_full_src_register *reg,
- enum tgsi_opcode_type type,
- unsigned swizzle)
+static void si_nir_store_output_tcs(struct ac_shader_abi *abi,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned const_index,
+ unsigned location,
+ unsigned driver_location,
+ LLVMValueRef src,
+ unsigned component,
+ bool is_patch,
+ bool is_compact,
+ unsigned writemask)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ struct tgsi_shader_info *info = &ctx->shader->selector->info;
+ LLVMValueRef dw_addr, stride;
+ LLVMValueRef buffer, base, addr;
+ LLVMValueRef values[4];
+ bool skip_lds_store;
+ bool is_tess_factor = false, is_tess_inner = false;
+
+ driver_location = driver_location / 4;
+
+ if (param_index) {
+ /* Add the constant index to the indirect index */
+ param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
+ LLVMConstInt(ctx->i32, const_index, 0), "");
+ } else {
+ if (const_index != 0)
+ param_index = LLVMConstInt(ctx->i32, const_index, 0);
+ }
+
+ if (!is_patch) {
+ stride = get_tcs_out_vertex_dw_stride(ctx);
+ dw_addr = get_tcs_out_current_patch_offset(ctx);
+ dw_addr = get_dw_address_from_generic_indices(ctx, stride, dw_addr,
+ vertex_index, param_index,
+ driver_location,
+ info->output_semantic_name,
+ info->output_semantic_index,
+ is_patch);
+
+ skip_lds_store = !info->reads_pervertex_outputs;
+ } else {
+ dw_addr = get_tcs_out_current_patch_data_offset(ctx);
+ dw_addr = get_dw_address_from_generic_indices(ctx, NULL, dw_addr,
+ vertex_index, param_index,
+ driver_location,
+ info->output_semantic_name,
+ info->output_semantic_index,
+ is_patch);
+
+ skip_lds_store = !info->reads_perpatch_outputs;
+
+ if (!param_index) {
+ int name = info->output_semantic_name[driver_location];
+
+ /* Always write tess factors into LDS for the TCS epilog. */
+ if (name == TGSI_SEMANTIC_TESSINNER ||
+ name == TGSI_SEMANTIC_TESSOUTER) {
+ /* The epilog doesn't read LDS if invocation 0 defines tess factors. */
+ skip_lds_store = !info->reads_tessfactor_outputs &&
+ ctx->shader->selector->tcs_info.tessfactors_are_def_in_all_invocs;
+ is_tess_factor = true;
+ is_tess_inner = name == TGSI_SEMANTIC_TESSINNER;
+ }
+ }
+ }
+
+ buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k);
+
+ base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
+
+ addr = get_tcs_tes_buffer_address_from_generic_indices(ctx, vertex_index,
+ param_index, driver_location,
+ info->output_semantic_name,
+ info->output_semantic_index,
+ is_patch);
+
+ for (unsigned chan = 0; chan < 4; chan++) {
+ if (!(writemask & (1 << chan)))
+ continue;
+ LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
+
+ /* Skip LDS stores if there is no LDS read of this output. */
+ if (!skip_lds_store)
+ ac_lds_store(&ctx->ac, dw_addr, value);
+
+ value = ac_to_integer(&ctx->ac, value);
+ values[chan] = value;
+
+ if (writemask != 0xF && !is_tess_factor) {
+ ac_build_buffer_store_dword(&ctx->ac, buffer, value, 1,
+ addr, base,
+ 4 * chan, 1, 0, true, false);
+ }
+
+ /* Write tess factors into VGPRs for the epilog. */
+ if (is_tess_factor &&
+ ctx->shader->selector->tcs_info.tessfactors_are_def_in_all_invocs) {
+ if (!is_tess_inner) {
+ LLVMBuildStore(ctx->ac.builder, value, /* outer */
+ ctx->invoc0_tess_factors[chan]);
+ } else if (chan < 2) {
+ LLVMBuildStore(ctx->ac.builder, value, /* inner */
+ ctx->invoc0_tess_factors[4 + chan]);
+ }
+ }
+ }
+
+ if (writemask == 0xF && !is_tess_factor) {
+ LLVMValueRef value = lp_build_gather_values(&ctx->gallivm,
+ values, 4);
+ ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, addr,
+ base, 0, 1, 0, true, false);
+ }
+}
+
+LLVMValueRef si_llvm_load_input_gs(struct ac_shader_abi *abi,
+ unsigned input_index,
+ unsigned vtx_offset_param,
+ LLVMTypeRef type,
+ unsigned swizzle)
+{
+ 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[reg->Register.Index];
- unsigned semantic_index = info->input_semantic_index[reg->Register.Index];
+ unsigned semantic_name = info->input_semantic_name[input_index];
+ unsigned semantic_index = info->input_semantic_index[input_index];
unsigned param;
LLVMValueRef value;
- if (swizzle != ~0 && semantic_name == TGSI_SEMANTIC_PRIMID)
- return get_primitive_id(ctx, swizzle);
-
- if (!reg->Register.Dimension)
- return NULL;
-
param = si_shader_io_get_unique_index(semantic_name, semantic_index);
/* GFX9 has the ESGS ring in LDS. */
if (ctx->screen->info.chip_class >= GFX9) {
- unsigned index = reg->Dimension.Index;
+ unsigned index = vtx_offset_param;
switch (index / 2) {
case 0:
LLVMValueRef values[TGSI_NUM_CHANNELS];
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
- values[chan] = fetch_input_gs(bld_base, reg, type, chan);
+ values[chan] = si_llvm_load_input_gs(abi, input_index, vtx_offset_param,
+ type, chan);
}
return lp_build_gather_values(&ctx->gallivm, values,
TGSI_NUM_CHANNELS);
}
/* Get the vertex offset parameter on GFX6. */
- unsigned vtx_offset_param = reg->Dimension.Index;
LLVMValueRef gs_vtx_offset = ctx->gs_vtx_offset[vtx_offset_param];
vtx_offset = lp_build_mul_imm(uint, gs_vtx_offset, 4);
value = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1, ctx->i32_0,
vtx_offset, soffset, 0, 1, 0, true, false);
- if (tgsi_type_is_64bit(type)) {
+ if (llvm_type_is_64bit(ctx, type)) {
LLVMValueRef value2;
soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle + 1) * 256, 0);
value2 = ac_build_buffer_load(&ctx->ac, ctx->esgs_ring, 1,
ctx->i32_0, vtx_offset, soffset,
0, 1, 0, true, false);
- return si_llvm_emit_fetch_64bit(bld_base, type,
- value, value2);
+ return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
}
- return bitcast(bld_base, type, value);
+ return LLVMBuildBitCast(ctx->ac.builder, value, type, "");
+}
+
+static LLVMValueRef fetch_input_gs(
+ struct lp_build_tgsi_context *bld_base,
+ const struct tgsi_full_src_register *reg,
+ enum tgsi_opcode_type type,
+ unsigned swizzle)
+{
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ struct tgsi_shader_info *info = &ctx->shader->selector->info;
+
+ unsigned semantic_name = info->input_semantic_name[reg->Register.Index];
+ if (swizzle != ~0 && semantic_name == TGSI_SEMANTIC_PRIMID)
+ return get_primitive_id(ctx, swizzle);
+
+ if (!reg->Register.Dimension)
+ return NULL;
+
+ return si_llvm_load_input_gs(&ctx->abi, reg->Register.Index,
+ reg->Dimension.Index,
+ tgsi2llvmtype(bld_base, type),
+ swizzle);
}
static int lookup_interp_param_index(unsigned interpolate, unsigned location)
return lp_build_gather_values(&ctx->gallivm, pos, 4);
}
+static LLVMValueRef si_load_tess_coord(struct ac_shader_abi *abi,
+ LLVMTypeRef type,
+ unsigned num_components)
+{
+ 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),
+ ctx->ac.f32_0,
+ ctx->ac.f32_0
+ };
+
+ /* 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);
+}
+
+static LLVMValueRef load_tess_level(struct si_shader_context *ctx,
+ unsigned semantic_name)
+{
+ LLVMValueRef buffer, base, addr;
+
+ int param = si_shader_io_get_unique_index_patch(semantic_name, 0);
+
+ buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k);
+
+ base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
+ addr = get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx), NULL,
+ LLVMConstInt(ctx->i32, param, 0));
+
+ return buffer_load(&ctx->bld_base, ctx->f32,
+ ~0, buffer, base, addr, true);
+
+}
+
+static LLVMValueRef si_load_tess_level(struct ac_shader_abi *abi,
+ unsigned varying_id)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ unsigned semantic_name;
+
+ switch (varying_id) {
+ case VARYING_SLOT_TESS_LEVEL_INNER:
+ semantic_name = TGSI_SEMANTIC_TESSINNER;
+ break;
+ case VARYING_SLOT_TESS_LEVEL_OUTER:
+ semantic_name = TGSI_SEMANTIC_TESSOUTER;
+ break;
+ default:
+ unreachable("unknown tess level");
+ }
+
+ return load_tess_level(ctx, semantic_name);
+
+}
+
+static LLVMValueRef si_load_patch_vertices_in(struct ac_shader_abi *abi)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ if (ctx->type == PIPE_SHADER_TESS_CTRL)
+ return unpack_param(ctx, ctx->param_tcs_out_lds_layout, 26, 6);
+ else if (ctx->type == PIPE_SHADER_TESS_EVAL)
+ return get_num_tcs_out_vertices(ctx);
+ else
+ unreachable("invalid shader stage for TGSI_SEMANTIC_VERTICESIN");
+}
+
void si_load_system_value(struct si_shader_context *ctx,
unsigned index,
const struct tgsi_full_declaration *decl)
{
- struct lp_build_context *bld = &ctx->bld_base.base;
LLVMValueRef value = 0;
assert(index < RADEON_LLVM_MAX_SYSTEM_VALUES);
case TGSI_SEMANTIC_INVOCATIONID:
if (ctx->type == PIPE_SHADER_TESS_CTRL)
- value = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
+ value = unpack_llvm_param(ctx, ctx->abi.tcs_rel_ids, 8, 5);
else if (ctx->type == PIPE_SHADER_GEOMETRY)
value = ctx->abi.gs_invocation_id;
else
break;
case TGSI_SEMANTIC_TESSCOORD:
- {
- LLVMValueRef coord[4] = {
- LLVMGetParam(ctx->main_fn, ctx->param_tes_u),
- LLVMGetParam(ctx->main_fn, ctx->param_tes_v),
- ctx->ac.f32_0,
- ctx->ac.f32_0
- };
-
- /* 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]));
-
- value = lp_build_gather_values(&ctx->gallivm, coord, 4);
+ value = si_load_tess_coord(&ctx->abi, NULL, 4);
break;
- }
case TGSI_SEMANTIC_VERTICESIN:
- if (ctx->type == PIPE_SHADER_TESS_CTRL)
- value = unpack_param(ctx, ctx->param_tcs_out_lds_layout, 26, 6);
- else if (ctx->type == PIPE_SHADER_TESS_EVAL)
- value = get_num_tcs_out_vertices(ctx);
- else
- assert(!"invalid shader stage for TGSI_SEMANTIC_VERTICESIN");
+ value = si_load_patch_vertices_in(&ctx->abi);
break;
case TGSI_SEMANTIC_TESSINNER:
case TGSI_SEMANTIC_TESSOUTER:
- {
- LLVMValueRef buffer, base, addr;
- int param = si_shader_io_get_unique_index_patch(decl->Semantic.Name, 0);
-
- buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k);
-
- base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
- addr = get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx), NULL,
- LLVMConstInt(ctx->i32, param, 0));
-
- value = buffer_load(&ctx->bld_base, TGSI_TYPE_FLOAT,
- ~0, buffer, base, addr, true);
-
+ value = load_tess_level(ctx, decl->Semantic.Name);
break;
- }
case TGSI_SEMANTIC_DEFAULT_TESSOUTER_SI:
case TGSI_SEMANTIC_DEFAULT_TESSINNER_SI:
lo = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, swizzle);
hi = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, swizzle + 1);
- return si_llvm_emit_fetch_64bit(bld_base, type, lo, hi);
+ return si_llvm_emit_fetch_64bit(bld_base, tgsi2llvmtype(bld_base, type),
+ lo, hi);
}
idx = reg->Register.Index * 4 + swizzle;
LLVMValueRef lds_vertex_stride, lds_vertex_offset, lds_base;
uint64_t inputs;
- invocation_id = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
+ invocation_id = unpack_llvm_param(ctx, ctx->abi.tcs_rel_ids, 8, 5);
buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k);
buffer_offset = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
invocation_id,
LLVMConstInt(ctx->i32, i, 0));
- LLVMValueRef value = lds_load(bld_base, TGSI_TYPE_SIGNED, ~0,
+ LLVMValueRef value = lds_load(bld_base, ctx->ac.i32, ~0,
lds_ptr);
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buffer_addr,
for (i = 0; i < outer_comps; i++) {
outer[i] = out[i] =
- lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_outer);
+ lds_load(bld_base, ctx->ac.i32, i, lds_outer);
}
for (i = 0; i < inner_comps; i++) {
inner[i] = out[outer_comps+i] =
- lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_inner);
+ lds_load(bld_base, ctx->ac.i32, i, lds_inner);
}
}
}
/* This only writes the tessellation factor levels. */
-static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base)
+static void si_llvm_emit_tcs_epilogue(struct ac_shader_abi *abi,
+ unsigned max_outputs,
+ LLVMValueRef *addrs)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef rel_patch_id, invocation_id, tf_lds_offset;
si_copy_tcs_inputs(bld_base);
rel_patch_id = get_rel_patch_id(ctx);
- invocation_id = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
+ invocation_id = unpack_llvm_param(ctx, ctx->abi.tcs_rel_ids, 8, 5);
tf_lds_offset = get_tcs_out_current_patch_data_offset(ctx);
if (ctx->screen->info.chip_class >= GFX9) {
tf_lds_offset = ac_to_float(&ctx->ac, tf_lds_offset);
/* Leave a hole corresponding to the two input VGPRs. This ensures that
- * the invocation_id output does not alias the param_tcs_rel_ids input,
+ * the invocation_id output does not alias the tcs_rel_ids input,
* which saves a V_MOV on gfx9.
*/
vgpr += 2;
8 + GFX9_SGPR_TCS_SAMPLERS_AND_IMAGES);
unsigned vgpr = 8 + GFX9_TCS_NUM_USER_SGPR;
- ret = si_insert_input_ret_float(ctx, ret,
- ctx->param_tcs_patch_id, vgpr++);
- ret = si_insert_input_ret_float(ctx, ret,
- ctx->param_tcs_rel_ids, vgpr++);
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
+ ac_to_float(&ctx->ac, ctx->abi.tcs_patch_id),
+ vgpr++, "");
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
+ ac_to_float(&ctx->ac, ctx->abi.tcs_rel_ids),
+ vgpr++, "");
ctx->return_value = ret;
}
return LLVMGetParam(ctx->main_fn, ctx->param_gs_wave_id);
}
-static void si_llvm_emit_gs_epilogue(struct lp_build_tgsi_context *bld_base)
+static void emit_gs_epilogue(struct si_shader_context *ctx)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
-
ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_NOP | AC_SENDMSG_GS_DONE,
si_get_gs_wave_id(ctx));
lp_build_endif(&ctx->merged_wrap_if_state);
}
+static void si_llvm_emit_gs_epilogue(struct ac_shader_abi *abi,
+ unsigned max_outputs,
+ LLVMValueRef *addrs)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+ struct tgsi_shader_info UNUSED *info = &ctx->shader->selector->info;
+
+ assert(info->num_outputs <= max_outputs);
+
+ emit_gs_epilogue(ctx);
+}
+
+static void si_tgsi_emit_gs_epilogue(struct lp_build_tgsi_context *bld_base)
+{
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_gs_epilogue(ctx);
+}
+
static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi,
unsigned max_outputs,
LLVMValueRef *addrs)
struct ac_export_args args[10];
};
-unsigned si_get_spi_shader_z_format(bool writes_z, bool writes_stencil,
- bool writes_samplemask)
-{
- if (writes_z) {
- /* Z needs 32 bits. */
- if (writes_samplemask)
- return V_028710_SPI_SHADER_32_ABGR;
- else if (writes_stencil)
- return V_028710_SPI_SHADER_32_GR;
- else
- return V_028710_SPI_SHADER_32_R;
- } else if (writes_stencil || writes_samplemask) {
- /* Both stencil and sample mask need only 16 bits. */
- return V_028710_SPI_SHADER_UINT16_ABGR;
- } else {
- return V_028710_SPI_SHADER_ZERO;
- }
-}
-
static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base,
LLVMValueRef depth, LLVMValueRef stencil,
LLVMValueRef samplemask, struct si_ps_exports *exp)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct lp_build_context *base = &bld_base->base;
struct ac_export_args args;
- unsigned mask = 0;
- unsigned format = si_get_spi_shader_z_format(depth != NULL,
- stencil != NULL,
- samplemask != NULL);
-
- assert(depth || stencil || samplemask);
-
- args.valid_mask = 1; /* whether the EXEC mask is valid */
- args.done = 1; /* DONE bit */
-
- /* Specify the target we are exporting */
- args.target = V_008DFC_SQ_EXP_MRTZ;
-
- args.compr = 0; /* COMP flag */
- args.out[0] = base->undef; /* R, depth */
- args.out[1] = base->undef; /* G, stencil test value[0:7], stencil op value[8:15] */
- args.out[2] = base->undef; /* B, sample mask */
- args.out[3] = base->undef; /* A, alpha to mask */
-
- if (format == V_028710_SPI_SHADER_UINT16_ABGR) {
- assert(!depth);
- args.compr = 1; /* COMPR flag */
-
- if (stencil) {
- /* Stencil should be in X[23:16]. */
- stencil = ac_to_integer(&ctx->ac, stencil);
- stencil = LLVMBuildShl(ctx->ac.builder, stencil,
- LLVMConstInt(ctx->i32, 16, 0), "");
- args.out[0] = ac_to_float(&ctx->ac, stencil);
- mask |= 0x3;
- }
- if (samplemask) {
- /* SampleMask should be in Y[15:0]. */
- args.out[1] = samplemask;
- mask |= 0xc;
- }
- } else {
- if (depth) {
- args.out[0] = depth;
- mask |= 0x1;
- }
- if (stencil) {
- args.out[1] = stencil;
- mask |= 0x2;
- }
- if (samplemask) {
- args.out[2] = samplemask;
- mask |= 0x4;
- }
- }
- /* SI (except OLAND and HAINAN) has a bug that it only looks
- * at the X writemask component. */
- if (ctx->screen->info.chip_class == SI &&
- ctx->screen->info.family != CHIP_OLAND &&
- ctx->screen->info.family != CHIP_HAINAN)
- mask |= 0x1;
-
- /* Specify which components to enable */
- args.enabled_channels = mask;
+ ac_export_mrt_z(&ctx->ac, depth, stencil, samplemask, &args);
memcpy(&exp->args[exp->num++], &args, sizeof(args));
}
ctx->return_value = ret;
}
-void si_emit_waitcnt(struct si_shader_context *ctx, unsigned simm16)
-{
- LLVMValueRef args[1] = {
- LLVMConstInt(ctx->i32, simm16, 0)
- };
- lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.s.waitcnt",
- ctx->voidt, args, 1, 0);
-}
-
static void membar_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
waitcnt &= LGKM_CNT;
if (waitcnt != NOOP_WAITCNT)
- si_emit_waitcnt(ctx, waitcnt);
+ ac_build_waitcnt(&ctx->ac, waitcnt);
}
static void clock_emit(
}
/* Cut one primitive from the geometry shader */
-static void si_llvm_emit_primitive(
+static void si_llvm_emit_primitive(struct ac_shader_abi *abi,
+ unsigned stream)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+
+ /* Signal primitive cut */
+ ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8),
+ si_get_gs_wave_id(ctx));
+}
+
+/* Cut one primitive from the geometry shader */
+static void si_tgsi_emit_primitive(
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);
- unsigned stream;
- /* Signal primitive cut */
- stream = si_llvm_get_stream(bld_base, emit_data);
- ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8),
- si_get_gs_wave_id(ctx));
+ si_llvm_emit_primitive(&ctx->abi, si_llvm_get_stream(bld_base, emit_data));
}
static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
*/
if (ctx->screen->info.chip_class == SI &&
ctx->type == PIPE_SHADER_TESS_CTRL) {
- si_emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
+ ac_build_waitcnt(&ctx->ac, LGKM_CNT & VM_CNT);
return;
}
ctx->param_tes_u = add_arg(fninfo, ARG_VGPR, ctx->f32);
ctx->param_tes_v = add_arg(fninfo, ARG_VGPR, ctx->f32);
ctx->param_tes_rel_patch_id = add_arg(fninfo, ARG_VGPR, ctx->i32);
- ctx->param_tes_patch_id = add_arg(fninfo, ARG_VGPR, ctx->i32);
+ add_arg_assign(fninfo, ARG_VGPR, ctx->i32, &ctx->abi.tes_patch_id);
}
enum {
ctx->param_tcs_factor_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
/* VGPRs */
- ctx->param_tcs_patch_id = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_tcs_rel_ids = add_arg(&fninfo, ARG_VGPR, ctx->i32);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.tcs_patch_id);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.tcs_rel_ids);
/* param_tcs_offchip_offset and param_tcs_factor_offset are
* placed after the user SGPRs.
ctx->type == PIPE_SHADER_TESS_CTRL);
/* VGPRs (first TCS, then VS) */
- ctx->param_tcs_patch_id = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_tcs_rel_ids = add_arg(&fninfo, ARG_VGPR, ctx->i32);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.tcs_patch_id);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.tcs_rel_ids);
if (ctx->type == PIPE_SHADER_VERTEX) {
declare_vs_input_vgprs(ctx, &fninfo,
r600_resource_reference(&shader->bo, NULL);
shader->bo = (struct r600_resource*)
- pipe_buffer_create(&sscreen->b, 0,
- PIPE_USAGE_IMMUTABLE,
- align(bo_size, SI_CPDMA_ALIGNMENT));
+ si_aligned_buffer_create(&sscreen->b,
+ sscreen->cpdma_prefetch_writes_memory ?
+ 0 : R600_RESOURCE_FLAG_READ_ONLY,
+ PIPE_USAGE_IMMUTABLE,
+ align(bo_size, SI_CPDMA_ALIGNMENT),
+ 256);
if (!shader->bo)
return -ENOMEM;
bld_base->op_actions[TGSI_OPCODE_READ_INVOC].emit = read_lane_emit;
bld_base->op_actions[TGSI_OPCODE_EMIT].emit = si_tgsi_emit_vertex;
- bld_base->op_actions[TGSI_OPCODE_ENDPRIM].emit = si_llvm_emit_primitive;
+ bld_base->op_actions[TGSI_OPCODE_ENDPRIM].emit = si_tgsi_emit_primitive;
bld_base->op_actions[TGSI_OPCODE_BARRIER].emit = si_llvm_emit_barrier;
}
}
}
-static void si_init_exec_full_mask(struct si_shader_context *ctx)
-{
- LLVMValueRef full_mask = LLVMConstInt(ctx->i64, ~0ull, 0);
- lp_build_intrinsic(ctx->ac.builder,
- "llvm.amdgcn.init.exec", ctx->voidt,
- &full_mask, 1, LP_FUNC_ATTR_CONVERGENT);
-}
-
static void si_init_exec_from_input(struct si_shader_context *ctx,
unsigned param, unsigned bitoffset)
{
break;
case PIPE_SHADER_TESS_CTRL:
bld_base->emit_fetch_funcs[TGSI_FILE_INPUT] = fetch_input_tcs;
+ ctx->abi.load_tess_varyings = si_nir_load_tcs_varyings;
bld_base->emit_fetch_funcs[TGSI_FILE_OUTPUT] = fetch_output_tcs;
bld_base->emit_store = store_output_tcs;
- bld_base->emit_epilogue = si_llvm_emit_tcs_epilogue;
+ ctx->abi.store_tcs_outputs = si_nir_store_output_tcs;
+ ctx->abi.emit_outputs = si_llvm_emit_tcs_epilogue;
+ ctx->abi.load_patch_vertices_in = si_load_patch_vertices_in;
+ bld_base->emit_epilogue = si_tgsi_emit_epilogue;
break;
case PIPE_SHADER_TESS_EVAL:
bld_base->emit_fetch_funcs[TGSI_FILE_INPUT] = fetch_input_tes;
+ ctx->abi.load_tess_varyings = si_nir_load_input_tes;
+ ctx->abi.load_tess_coord = si_load_tess_coord;
+ ctx->abi.load_tess_level = si_load_tess_level;
+ ctx->abi.load_patch_vertices_in = si_load_patch_vertices_in;
if (shader->key.as_es)
ctx->abi.emit_outputs = si_llvm_emit_es_epilogue;
else
break;
case PIPE_SHADER_GEOMETRY:
bld_base->emit_fetch_funcs[TGSI_FILE_INPUT] = fetch_input_gs;
+ ctx->abi.load_inputs = si_nir_load_input_gs;
ctx->abi.emit_vertex = si_llvm_emit_vertex;
- bld_base->emit_epilogue = si_llvm_emit_gs_epilogue;
+ ctx->abi.emit_outputs = si_llvm_emit_gs_epilogue;
+ bld_base->emit_epilogue = si_tgsi_emit_gs_epilogue;
break;
case PIPE_SHADER_FRAGMENT:
ctx->load_input = declare_input_fs;
} else if (ctx->type == PIPE_SHADER_TESS_CTRL ||
ctx->type == PIPE_SHADER_GEOMETRY) {
if (!is_monolithic)
- si_init_exec_full_mask(ctx);
+ ac_init_exec_full_mask(&ctx->ac);
/* The barrier must execute for all shaders in a
* threadgroup.
* mask.
*/
if (ctx->screen->info.chip_class >= GFX9 && !key->gs_prolog.is_monolithic)
- si_init_exec_full_mask(ctx);
+ ac_init_exec_full_mask(&ctx->ac);
/* Copy inputs to outputs. This should be no-op, as the registers match,
* but it will prevent the compiler from overwriting them unintentionally.
si_get_max_workgroup_size(ctx->shader));
if (is_merged_shader(ctx->shader))
- si_init_exec_full_mask(ctx);
+ ac_init_exec_full_mask(&ctx->ac);
/* Record the arguments of the function as if they were an output of
* a previous part.
projects / mesa.git / blobdiff