* DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
* OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
* USE OR OTHER DEALINGS IN THE SOFTWARE.
- *
- * Authors:
- * Tom Stellard <thomas.stellard@amd.com>
- * Michel Dänzer <michel.daenzer@amd.com>
- * Christian König <christian.koenig@amd.com>
*/
#include "gallivm/lp_bld_const.h"
static bool is_merged_shader(struct si_shader *shader)
{
- if (shader->selector->screen->b.chip_class <= VI)
+ if (shader->selector->screen->info.chip_class <= VI)
return false;
return shader->key.as_ls ||
return LLVMConstInt(ctx->i32, stride * 4, 0);
case PIPE_SHADER_TESS_CTRL:
- if (ctx->screen->b.chip_class >= GFX9 &&
+ if (ctx->screen->info.chip_class >= GFX9 &&
ctx->shader->is_monolithic) {
stride = util_last_bit64(ctx->shader->key.part.tcs.ls->outputs_written);
return LLVMConstInt(ctx->i32, stride * 4, 0);
unsigned double_index)
{
LLVMBuilderRef builder = ctx->ac.builder;
- LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->gallivm.context);
+ LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->ac.context);
LLVMValueRef dvec2 = LLVMBuildBitCast(builder, vec4,
LLVMVectorType(f64, 2), "");
LLVMValueRef index = LLVMConstInt(ctx->i32, double_index, 0);
return LLVMBuildFPTrunc(builder, value, ctx->f32, "");
}
+static LLVMValueRef unpack_sint16(struct si_shader_context *ctx,
+ LLVMValueRef i32, unsigned index)
+{
+ assert(index <= 1);
+
+ if (index == 1)
+ return LLVMBuildAShr(ctx->ac.builder, i32,
+ LLVMConstInt(ctx->i32, 16, 0), "");
+
+ return LLVMBuildSExt(ctx->ac.builder,
+ LLVMBuildTrunc(ctx->ac.builder, i32,
+ ctx->ac.i16, ""),
+ ctx->i32, "");
+}
+
void si_llvm_load_input_vs(
struct si_shader_context *ctx,
unsigned input_index,
LLVMValueRef out[4])
{
+ unsigned vs_blit_property =
+ ctx->shader->selector->info.properties[TGSI_PROPERTY_VS_BLIT_SGPRS];
+
+ if (vs_blit_property) {
+ LLVMValueRef vertex_id = ctx->abi.vertex_id;
+ LLVMValueRef sel_x1 = LLVMBuildICmp(ctx->ac.builder,
+ LLVMIntULE, vertex_id,
+ ctx->i32_1, "");
+ /* Use LLVMIntNE, because we have 3 vertices and only
+ * the middle one should use y2.
+ */
+ LLVMValueRef sel_y1 = LLVMBuildICmp(ctx->ac.builder,
+ LLVMIntNE, vertex_id,
+ ctx->i32_1, "");
+
+ if (input_index == 0) {
+ /* Position: */
+ LLVMValueRef x1y1 = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs);
+ LLVMValueRef x2y2 = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 1);
+
+ LLVMValueRef x1 = unpack_sint16(ctx, x1y1, 0);
+ LLVMValueRef y1 = unpack_sint16(ctx, x1y1, 1);
+ LLVMValueRef x2 = unpack_sint16(ctx, x2y2, 0);
+ LLVMValueRef y2 = unpack_sint16(ctx, x2y2, 1);
+
+ LLVMValueRef x = LLVMBuildSelect(ctx->ac.builder, sel_x1,
+ x1, x2, "");
+ LLVMValueRef y = LLVMBuildSelect(ctx->ac.builder, sel_y1,
+ y1, y2, "");
+
+ out[0] = LLVMBuildSIToFP(ctx->ac.builder, x, ctx->f32, "");
+ out[1] = LLVMBuildSIToFP(ctx->ac.builder, y, ctx->f32, "");
+ out[2] = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 2);
+ out[3] = ctx->ac.f32_1;
+ return;
+ }
+
+ /* Color or texture coordinates: */
+ assert(input_index == 1);
+
+ if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_COLOR) {
+ for (int i = 0; i < 4; i++) {
+ out[i] = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 3 + i);
+ }
+ } else {
+ assert(vs_blit_property == SI_VS_BLIT_SGPRS_POS_TEXCOORD);
+ LLVMValueRef x1 = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 3);
+ LLVMValueRef y1 = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 4);
+ LLVMValueRef x2 = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 5);
+ LLVMValueRef y2 = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 6);
+
+ out[0] = LLVMBuildSelect(ctx->ac.builder, sel_x1,
+ x1, x2, "");
+ out[1] = LLVMBuildSelect(ctx->ac.builder, sel_y1,
+ y1, y2, "");
+ out[2] = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 7);
+ out[3] = LLVMGetParam(ctx->main_fn,
+ ctx->param_vs_blit_inputs + 8);
+ }
+ return;
+ }
+
unsigned chan;
unsigned fix_fetch;
unsigned num_fetches;
t_offset = LLVMConstInt(ctx->i32, input_index, 0);
- t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr, t_offset);
+ t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
vertex_index = LLVMGetParam(ctx->main_fn,
ctx->param_vertex_index0 +
return LLVMGetParam(ctx->main_fn,
ctx->param_tes_patch_id);
case PIPE_SHADER_GEOMETRY:
- return LLVMGetParam(ctx->main_fn,
- ctx->param_gs_prim_id);
+ return ctx->abi.gs_prim_id;
default:
assert(0);
return ctx->i32_0;
value2 = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
swizzle * 4 + 4, 1, 0, can_speculate, false);
- return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
+ return si_llvm_emit_fetch_64bit(bld_base, tgsi2llvmtype(bld_base, type),
+ value, value2);
}
/**
LLVMValueRef dw_addr)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef value;
if (swizzle == ~0) {
for (unsigned chan = 0; chan < TGSI_NUM_CHANNELS; chan++)
values[chan] = lds_load(bld_base, type, chan, dw_addr);
- return lp_build_gather_values(gallivm, values,
+ return lp_build_gather_values(&ctx->gallivm, values,
TGSI_NUM_CHANNELS);
}
+ /* Split 64-bit loads. */
+ if (tgsi_type_is_64bit(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);
+ return si_llvm_emit_fetch_64bit(bld_base, tgsi2llvmtype(bld_base, type),
+ lo, hi);
+ }
+
dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
LLVMConstInt(ctx->i32, swizzle, 0));
- value = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
- if (tgsi_type_is_64bit(type)) {
- LLVMValueRef value2;
- dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
- ctx->i32_1);
- value2 = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
- return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
- }
+ value = ac_lds_load(&ctx->ac, dw_addr);
return bitcast(bld_base, type, value);
}
* \param dw_addr address in dwords
* \param value value to store
*/
-static void lds_store(struct lp_build_tgsi_context *bld_base,
+static void lds_store(struct si_shader_context *ctx,
unsigned dw_offset_imm, LLVMValueRef dw_addr,
LLVMValueRef value)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
-
- dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr,
+ dw_addr = lp_build_add(&ctx->bld_base.uint_bld, dw_addr,
LLVMConstInt(ctx->i32, dw_offset_imm, 0));
- value = ac_to_integer(&ctx->ac, value);
- ac_build_indexed_store(&ctx->ac, ctx->lds,
- dw_addr, value);
+ ac_lds_store(&ctx->ac, dw_addr, value);
}
static LLVMValueRef desc_from_addr_base64k(struct si_shader_context *ctx,
LLVMValueRef dst[4])
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_dst_register *reg = &inst->Dst[index];
const struct tgsi_shader_info *sh_info = &ctx->shader->selector->info;
unsigned chan_index;
/* Skip LDS stores if there is no LDS read of this output. */
if (!skip_lds_store)
- lds_store(bld_base, chan_index, dw_addr, value);
+ lds_store(ctx, chan_index, dw_addr, value);
value = ac_to_integer(&ctx->ac, value);
values[chan_index] = value;
}
if (reg->Register.WriteMask == 0xF && !is_tess_factor) {
- LLVMValueRef value = lp_build_gather_values(gallivm,
+ LLVMValueRef value = lp_build_gather_values(&ctx->gallivm,
values, 4);
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr,
base, 0, 1, 0, true, false);
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
struct lp_build_context *uint = &ctx->bld_base.uint_bld;
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef vtx_offset, soffset;
struct tgsi_shader_info *info = &shader->selector->info;
unsigned semantic_name = info->input_semantic_name[reg->Register.Index];
param = si_shader_io_get_unique_index(semantic_name, semantic_index);
/* GFX9 has the ESGS ring in LDS. */
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
unsigned index = reg->Dimension.Index;
switch (index / 2) {
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
values[chan] = fetch_input_gs(bld_base, reg, type, chan);
}
- return lp_build_gather_values(gallivm, values,
+ 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;
- if (vtx_offset_param < 2) {
- vtx_offset_param += ctx->param_gs_vtx0_offset;
- } else {
- assert(vtx_offset_param < 6);
- vtx_offset_param += ctx->param_gs_vtx2_offset - 2;
- }
- vtx_offset = lp_build_mul_imm(uint,
- LLVMGetParam(ctx->main_fn,
- vtx_offset_param),
- 4);
+ LLVMValueRef gs_vtx_offset = ctx->gs_vtx_offset[vtx_offset_param];
+
+ vtx_offset = lp_build_mul_imm(uint, gs_vtx_offset, 4);
soffset = LLVMConstInt(ctx->i32, (param * 4 + swizzle) * 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,
+ return si_llvm_emit_fetch_64bit(bld_base, tgsi2llvmtype(bld_base, type),
value, value2);
}
return bitcast(bld_base, type, value);
static LLVMValueRef load_sample_position(struct si_shader_context *ctx, LLVMValueRef sample_id)
{
struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
- struct gallivm_state *gallivm = &ctx->gallivm;
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_indexed_load_const(&ctx->ac, desc, buf_index);
+ 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);
LLVMConstReal(ctx->f32, 0)
};
- return lp_build_gather_values(gallivm, pos, 4);
+ return lp_build_gather_values(&ctx->gallivm, pos, 4);
}
void si_load_system_value(struct si_shader_context *ctx,
const struct tgsi_full_declaration *decl)
{
struct lp_build_context *bld = &ctx->bld_base.base;
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef value = 0;
assert(index < RADEON_LLVM_MAX_SYSTEM_VALUES);
if (ctx->type == PIPE_SHADER_TESS_CTRL)
value = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
else if (ctx->type == PIPE_SHADER_GEOMETRY)
- value = LLVMGetParam(ctx->main_fn,
- ctx->param_gs_instance_id);
+ value = ctx->abi.gs_invocation_id;
else
assert(!"INVOCATIONID not implemented");
break;
LLVMGetParam(ctx->main_fn,
SI_PARAM_POS_W_FLOAT)),
};
- value = lp_build_gather_values(gallivm, pos, 4);
+ value = lp_build_gather_values(&ctx->gallivm, pos, 4);
break;
}
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(gallivm, pos, 4);
+ value = lp_build_gather_values(&ctx->gallivm, pos, 4);
break;
}
LLVMValueRef coord[4] = {
LLVMGetParam(ctx->main_fn, ctx->param_tes_u),
LLVMGetParam(ctx->main_fn, ctx->param_tes_v),
- bld->zero,
- bld->zero
+ 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, bld->one,
+ coord[2] = lp_build_sub(bld, ctx->ac.f32_1,
lp_build_add(bld, coord[0], coord[1]));
- value = lp_build_gather_values(gallivm, coord, 4);
+ value = lp_build_gather_values(&ctx->gallivm, coord, 4);
break;
}
slot = LLVMConstInt(ctx->i32, SI_HS_CONST_DEFAULT_TESS_LEVELS, 0);
buf = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
- buf = ac_build_indexed_load_const(&ctx->ac, buf, slot);
+ buf = ac_build_load_to_sgpr(&ctx->ac, buf, slot);
offset = decl->Semantic.Name == TGSI_SEMANTIC_DEFAULT_TESSINNER_SI ? 4 : 0;
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(gallivm, val, 4);
+ value = lp_build_gather_values(&ctx->gallivm, val, 4);
break;
}
for (i = 0; i < 3; ++i)
values[i] = LLVMConstInt(ctx->i32, sizes[i], 0);
- value = lp_build_gather_values(gallivm, values, 3);
+ value = lp_build_gather_values(&ctx->gallivm, values, 3);
} else {
value = LLVMGetParam(ctx->main_fn, ctx->param_block_size);
}
ctx->param_block_id[i]);
}
}
- value = lp_build_gather_values(gallivm, values, 3);
+ value = lp_build_gather_values(&ctx->gallivm, values, 3);
break;
}
const struct tgsi_full_declaration *decl)
{
struct si_shader_selector *sel = ctx->shader->selector;
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMTypeRef i8p = LLVMPointerType(ctx->i8, LOCAL_ADDR_SPACE);
LLVMValueRef var;
assert(decl->Declaration.MemType == TGSI_MEMORY_TYPE_SHARED);
assert(decl->Range.First == decl->Range.Last);
- assert(!ctx->shared_memory);
+ assert(!ctx->ac.lds);
- var = LLVMAddGlobalInAddressSpace(gallivm->module,
+ var = LLVMAddGlobalInAddressSpace(ctx->ac.module,
LLVMArrayType(ctx->i8, sel->local_size),
"compute_lds",
LOCAL_ADDR_SPACE);
LLVMSetAlignment(var, 4);
- ctx->shared_memory = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
+ ctx->ac.lds = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
}
static LLVMValueRef load_const_buffer_desc(struct si_shader_context *ctx, int i)
LLVMValueRef list_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_const_and_shader_buffers);
- return ac_build_indexed_load_const(&ctx->ac, list_ptr,
- LLVMConstInt(ctx->i32, si_get_constbuf_slot(i), 0));
+ return ac_build_load_to_sgpr(&ctx->ac, list_ptr,
+ LLVMConstInt(ctx->i32, si_get_constbuf_slot(i), 0));
}
static LLVMValueRef load_ubo(struct ac_shader_abi *abi, LLVMValueRef index)
index = LLVMBuildAdd(ctx->ac.builder, index,
LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS, 0), "");
- return ac_build_indexed_load_const(&ctx->ac, ptr, index);
+ return ac_build_load_to_sgpr(&ctx->ac, ptr, index);
}
static LLVMValueRef
LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS - 1, 0),
index, "");
- return ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
+ return ac_build_load_to_sgpr(&ctx->ac, rsrc_ptr, index);
}
static LLVMValueRef fetch_constant(
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
+ struct si_shader_selector *sel = ctx->shader->selector;
const struct tgsi_ind_register *ireg = ®->Indirect;
unsigned buf, idx;
LLVMValueRef addr, bufp;
- LLVMValueRef result;
if (swizzle == LP_CHAN_ALL) {
unsigned chan;
return lp_build_gather_values(&ctx->gallivm, values, 4);
}
+ /* Split 64-bit loads. */
+ if (tgsi_type_is_64bit(type)) {
+ LLVMValueRef lo, hi;
+
+ 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, tgsi2llvmtype(bld_base, type),
+ lo, hi);
+ }
+
+ idx = reg->Register.Index * 4 + swizzle;
+ if (reg->Register.Indirect) {
+ addr = si_get_indirect_index(ctx, ireg, 16, idx * 4);
+ } else {
+ addr = LLVMConstInt(ctx->i32, idx * 4, 0);
+ }
+
+ /* Fast path when user data SGPRs point to constant buffer 0 directly. */
+ if (sel->info.const_buffers_declared == 1 &&
+ sel->info.shader_buffers_declared == 0) {
+ LLVMValueRef ptr =
+ LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers);
+
+ /* This enables use of s_load_dword and flat_load_dword for const buffer 0
+ * loads, and up to x4 load opcode merging. However, it leads to horrible
+ * code reducing SIMD wave occupancy from 8 to 2 in many cases.
+ *
+ * Using s_buffer_load_dword (x1) seems to be the best option right now.
+ *
+ * LLVM 5.0 on SI doesn't insert a required s_nop between SALU setting
+ * a descriptor and s_buffer_load_dword using it, so we can't expand
+ * the pointer into a full descriptor like below. We have to use
+ * s_load_dword instead. The only case when LLVM 5.0 would select
+ * s_buffer_load_dword (that we have to prevent) is when we use use
+ * a literal offset where we don't need bounds checking.
+ */
+ if (ctx->screen->info.chip_class == SI &&
+ HAVE_LLVM < 0x0600 &&
+ !reg->Register.Indirect) {
+ addr = LLVMBuildLShr(ctx->ac.builder, addr, LLVMConstInt(ctx->i32, 2, 0), "");
+ LLVMValueRef result = ac_build_load_invariant(&ctx->ac, ptr, addr);
+ return bitcast(bld_base, type, result);
+ }
+
+ /* Do the bounds checking with a descriptor, because
+ * doing computation and manual bounds checking of 64-bit
+ * addresses generates horrible VALU code with very high
+ * VGPR usage and very low SIMD occupancy.
+ */
+ ptr = LLVMBuildPtrToInt(ctx->ac.builder, ptr, ctx->i64, "");
+ ptr = LLVMBuildBitCast(ctx->ac.builder, ptr, ctx->v2i32, "");
+
+ LLVMValueRef desc_elems[] = {
+ LLVMBuildExtractElement(ctx->ac.builder, ptr, ctx->i32_0, ""),
+ LLVMBuildExtractElement(ctx->ac.builder, ptr, ctx->i32_1, ""),
+ LLVMConstInt(ctx->i32, (sel->info.const_file_max[0] + 1) * 16, 0),
+ 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) |
+ S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32), 0)
+ };
+ LLVMValueRef desc = ac_build_gather_values(&ctx->ac, desc_elems, 4);
+ LLVMValueRef result = buffer_load_const(ctx, desc, addr);
+ return bitcast(bld_base, type, result);
+ }
+
assert(reg->Register.Dimension);
buf = reg->Dimension.Index;
- idx = reg->Register.Index * 4 + swizzle;
if (reg->Dimension.Indirect) {
LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers);
ctx->num_const_buffers);
index = LLVMBuildAdd(ctx->ac.builder, index,
LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS, 0), "");
- bufp = ac_build_indexed_load_const(&ctx->ac, ptr, index);
+ bufp = ac_build_load_to_sgpr(&ctx->ac, ptr, index);
} else
bufp = load_const_buffer_desc(ctx, buf);
- if (reg->Register.Indirect) {
- addr = si_get_indirect_index(ctx, ireg, 16, idx * 4);
- } else {
- addr = LLVMConstInt(ctx->i32, idx * 4, 0);
- }
-
- result = buffer_load_const(ctx, bufp, addr);
-
- if (!tgsi_type_is_64bit(type))
- result = bitcast(bld_base, type, result);
- else {
- LLVMValueRef addr2, result2;
-
- addr2 = lp_build_add(&bld_base->uint_bld, addr,
- LLVMConstInt(ctx->i32, 4, 0));
- result2 = buffer_load_const(ctx, bufp, addr2);
-
- result = si_llvm_emit_fetch_64bit(bld_base, type,
- result, result2);
- }
- return result;
+ return bitcast(bld_base, type, buffer_load_const(ctx, bufp, addr));
}
/* Upper 16 bits must be zero. */
}
/* Initialize arguments for the shader export intrinsic */
-static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
+static void si_llvm_init_export_args(struct si_shader_context *ctx,
LLVMValueRef *values,
unsigned target,
struct ac_export_args *args)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
- struct lp_build_context *base = &bld_base->base;
+ LLVMValueRef f32undef = LLVMGetUndef(ctx->ac.f32);
LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef val[4];
unsigned spi_shader_col_format = V_028714_SPI_SHADER_32_ABGR;
}
args->compr = false;
- args->out[0] = base->undef;
- args->out[1] = base->undef;
- args->out[2] = base->undef;
- args->out[3] = base->undef;
+ args->out[0] = f32undef;
+ args->out[1] = f32undef;
+ args->out[2] = f32undef;
+ args->out[3] = f32undef;
switch (spi_shader_col_format) {
case V_028714_SPI_SHADER_ZERO:
case V_028714_SPI_SHADER_SNORM16_ABGR:
for (chan = 0; chan < 4; chan++) {
/* Clamp between [-1, 1]. */
- val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_MIN,
+ val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_MIN,
values[chan],
LLVMConstReal(ctx->f32, 1));
- val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_MAX,
+ val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_MAX,
val[chan],
LLVMConstReal(ctx->f32, -1));
/* Convert to a signed integer in [-32767, 32767]. */
val[chan] = LLVMBuildFAdd(builder, val[chan],
LLVMBuildSelect(builder,
LLVMBuildFCmp(builder, LLVMRealOGE,
- val[chan], base->zero, ""),
+ val[chan], ctx->ac.f32_0, ""),
LLVMConstReal(ctx->f32, 0.5),
LLVMConstReal(ctx->f32, -0.5), ""), "");
val[chan] = LLVMBuildFPToSI(builder, val[chan], ctx->i32, "");
/* Clamp. */
for (chan = 0; chan < 4; chan++) {
val[chan] = ac_to_integer(&ctx->ac, values[chan]);
- val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_UMIN,
+ val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_UMIN,
val[chan],
chan == 3 ? max_alpha : max_rgb);
}
/* Clamp. */
for (chan = 0; chan < 4; chan++) {
val[chan] = ac_to_integer(&ctx->ac, values[chan]);
- val[chan] = lp_build_emit_llvm_binary(bld_base,
+ val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base,
TGSI_OPCODE_IMIN,
val[chan], chan == 3 ? max_alpha : max_rgb);
- val[chan] = lp_build_emit_llvm_binary(bld_base,
+ val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base,
TGSI_OPCODE_IMAX,
val[chan], chan == 3 ? min_alpha : min_rgb);
}
struct si_shader_context *ctx = si_shader_context(bld_base);
if (ctx->shader->key.part.ps.epilog.alpha_func != PIPE_FUNC_NEVER) {
+ static LLVMRealPredicate cond_map[PIPE_FUNC_ALWAYS + 1] = {
+ [PIPE_FUNC_LESS] = LLVMRealOLT,
+ [PIPE_FUNC_EQUAL] = LLVMRealOEQ,
+ [PIPE_FUNC_LEQUAL] = LLVMRealOLE,
+ [PIPE_FUNC_GREATER] = LLVMRealOGT,
+ [PIPE_FUNC_NOTEQUAL] = LLVMRealONE,
+ [PIPE_FUNC_GEQUAL] = LLVMRealOGE,
+ };
+ LLVMRealPredicate cond = cond_map[ctx->shader->key.part.ps.epilog.alpha_func];
+ assert(cond);
+
LLVMValueRef alpha_ref = LLVMGetParam(ctx->main_fn,
SI_PARAM_ALPHA_REF);
-
LLVMValueRef alpha_pass =
- lp_build_cmp(&bld_base->base,
- ctx->shader->key.part.ps.epilog.alpha_func,
- alpha, alpha_ref);
- LLVMValueRef arg =
- lp_build_select(&bld_base->base,
- alpha_pass,
- LLVMConstReal(ctx->f32, 1.0f),
- LLVMConstReal(ctx->f32, -1.0f));
-
- ac_build_kill(&ctx->ac, arg);
+ LLVMBuildFCmp(ctx->ac.builder, cond, alpha, alpha_ref, "");
+ ac_build_kill_if_false(&ctx->ac, alpha_pass);
} else {
- ac_build_kill(&ctx->ac, NULL);
+ ac_build_kill_if_false(&ctx->ac, LLVMConstInt(ctx->i1, 0, 0));
}
}
return LLVMBuildFMul(ctx->ac.builder, alpha, coverage, "");
}
-static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
+static void si_llvm_emit_clipvertex(struct si_shader_context *ctx,
struct ac_export_args *pos, LLVMValueRef *out_elts)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
- struct lp_build_context *base = &bld_base->base;
unsigned reg_index;
unsigned chan;
unsigned const_chan;
LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
LLVMValueRef constbuf_index = LLVMConstInt(ctx->i32,
SI_VS_CONST_CLIP_PLANES, 0);
- LLVMValueRef const_resource = ac_build_indexed_load_const(&ctx->ac, ptr, constbuf_index);
+ LLVMValueRef const_resource = ac_build_load_to_sgpr(&ctx->ac, ptr, constbuf_index);
for (reg_index = 0; reg_index < 2; reg_index ++) {
struct ac_export_args *args = &pos[2 + reg_index];
base_elt = buffer_load_const(ctx, const_resource,
addr);
args->out[chan] =
- lp_build_add(base, args->out[chan],
- lp_build_mul(base, base_elt,
+ lp_build_add(&ctx->bld_base.base, args->out[chan],
+ lp_build_mul(&ctx->bld_base.base, base_elt,
out_elts[const_chan]));
}
}
{
struct si_shader_selector *sel = ctx->shader->selector;
struct pipe_stream_output_info *so = &sel->so;
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = ctx->ac.builder;
int i;
struct lp_build_if_state if_ctx;
/* Emit the streamout code conditionally. This actually avoids
* out-of-bounds buffer access. The hw tells us via the SGPR
* (so_vtx_count) which threads are allowed to emit streamout data. */
- lp_build_if(&if_ctx, gallivm, can_emit);
+ lp_build_if(&if_ctx, &ctx->gallivm, can_emit);
{
/* The buffer offset is computed as follows:
* ByteOffset = streamout_offset[buffer_id]*4 +
LLVMValueRef offset = LLVMConstInt(ctx->i32,
SI_VS_STREAMOUT_BUF0 + i, 0);
- so_buffers[i] = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
+ so_buffers[i] = ac_build_load_to_sgpr(&ctx->ac, buf_ptr, offset);
LLVMValueRef so_offset = LLVMGetParam(ctx->main_fn,
ctx->param_streamout_offset[i]);
{
struct ac_export_args args;
- si_llvm_init_export_args(&ctx->bld_base, values,
+ si_llvm_init_export_args(ctx, values,
V_008DFC_SQ_EXP_PARAM + index, &args);
ac_build_export(&ctx->ac, &args);
}
}
/* Generate export instructions for hardware VS shader stage */
-static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
+static void si_llvm_export_vs(struct si_shader_context *ctx,
struct si_shader_output_values *outputs,
unsigned noutput)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
- struct lp_build_context *base = &bld_base->base;
struct ac_export_args pos_args[4] = {};
LLVMValueRef psize_value = NULL, edgeflag_value = NULL, layer_value = NULL, viewport_index_value = NULL;
unsigned pos_idx;
for (i = 0; i < noutput; i++) {
switch (outputs[i].semantic_name) {
case TGSI_SEMANTIC_POSITION:
- si_llvm_init_export_args(bld_base, outputs[i].values,
+ si_llvm_init_export_args(ctx, outputs[i].values,
V_008DFC_SQ_EXP_POS, &pos_args[0]);
break;
case TGSI_SEMANTIC_PSIZE:
case TGSI_SEMANTIC_CLIPDIST:
if (!shader->key.opt.clip_disable) {
unsigned index = 2 + outputs[i].semantic_index;
- si_llvm_init_export_args(bld_base, outputs[i].values,
+ si_llvm_init_export_args(ctx, outputs[i].values,
V_008DFC_SQ_EXP_POS + index,
&pos_args[index]);
}
break;
case TGSI_SEMANTIC_CLIPVERTEX:
if (!shader->key.opt.clip_disable) {
- si_llvm_emit_clipvertex(bld_base, pos_args,
+ si_llvm_emit_clipvertex(ctx, pos_args,
outputs[i].values);
}
break;
pos_args[0].done = 0; /* last export? */
pos_args[0].target = V_008DFC_SQ_EXP_POS;
pos_args[0].compr = 0; /* COMPR flag */
- pos_args[0].out[0] = base->zero; /* X */
- pos_args[0].out[1] = base->zero; /* Y */
- pos_args[0].out[2] = base->zero; /* Z */
- pos_args[0].out[3] = base->one; /* W */
+ pos_args[0].out[0] = ctx->ac.f32_0; /* X */
+ pos_args[0].out[1] = ctx->ac.f32_0; /* Y */
+ pos_args[0].out[2] = ctx->ac.f32_0; /* Z */
+ pos_args[0].out[3] = ctx->ac.f32_1; /* W */
}
/* Write the misc vector (point size, edgeflag, layer, viewport). */
pos_args[1].done = 0; /* last export? */
pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
pos_args[1].compr = 0; /* COMPR flag */
- pos_args[1].out[0] = base->zero; /* X */
- pos_args[1].out[1] = base->zero; /* Y */
- pos_args[1].out[2] = base->zero; /* Z */
- pos_args[1].out[3] = base->zero; /* W */
+ pos_args[1].out[0] = ctx->ac.f32_0; /* X */
+ pos_args[1].out[1] = ctx->ac.f32_0; /* Y */
+ pos_args[1].out[2] = ctx->ac.f32_0; /* Z */
+ pos_args[1].out[3] = ctx->ac.f32_0; /* W */
if (shader->selector->info.writes_psize)
pos_args[1].out[0] = psize_value;
pos_args[1].out[1] = ac_to_float(&ctx->ac, edgeflag_value);
}
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
/* GFX9 has the layer in out.z[10:0] and the viewport
* index in out.z[19:16].
*/
LLVMValueRef invoc0_tf_inner[2])
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_shader *shader = ctx->shader;
unsigned tess_inner_index, tess_outer_index;
LLVMValueRef lds_base, lds_inner, lds_outer, byteoffset, buffer;
* 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,
+ lp_build_if(&if_ctx, &ctx->gallivm,
LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
invocation_id, ctx->i32_0, ""));
}
/* Convert the outputs to vectors for stores. */
- vec0 = lp_build_gather_values(gallivm, out, MIN2(stride, 4));
+ vec0 = lp_build_gather_values(&ctx->gallivm, out, MIN2(stride, 4));
vec1 = NULL;
if (stride > 4)
- vec1 = lp_build_gather_values(gallivm, out+4, stride - 4);
+ vec1 = lp_build_gather_values(&ctx->gallivm, out+4, stride - 4);
/* Get the buffer. */
buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_factor_addr_base64k);
byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
LLVMConstInt(ctx->i32, 4 * stride, 0), "");
- lp_build_if(&inner_if_ctx, gallivm,
+ lp_build_if(&inner_if_ctx, &ctx->gallivm,
LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
rel_patch_id, ctx->i32_0, ""));
/* Store the dynamic HS control word. */
offset = 0;
- if (ctx->screen->b.chip_class <= VI) {
+ if (ctx->screen->info.chip_class <= VI) {
ac_build_buffer_store_dword(&ctx->ac, buffer,
LLVMConstInt(ctx->i32, 0x80000000, 0),
1, ctx->i32_0, tf_base,
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(gallivm, outer,
+ outer_vec = lp_build_gather_values(&ctx->gallivm, 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(gallivm, inner, inner_comps);
+ lp_build_gather_values(&ctx->gallivm, inner, inner_comps);
ac_build_buffer_store_dword(&ctx->ac, buf, inner_vec,
inner_comps, tf_inner_offset,
base, 0, 1, 0, true, false);
invocation_id = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
tf_lds_offset = get_tcs_out_current_patch_data_offset(ctx);
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
LLVMBasicBlockRef blocks[2] = {
LLVMGetInsertBlock(builder),
ctx->merged_wrap_if_state.entry_block
LLVMValueRef ret = ctx->return_value;
unsigned vgpr;
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_offchip_layout,
8 + GFX9_SGPR_TCS_OFFCHIP_LAYOUT);
ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_offchip_addr_base64k,
{
LLVMValueRef ret = ctx->return_value;
- ret = si_insert_input_ptr_as_2xi32(ctx, ret, ctx->param_rw_buffers, 0);
ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_offchip_offset, 2);
ret = si_insert_input_ret(ctx, ret, ctx->param_merged_wave_info, 3);
ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_factor_offset, 4);
ret = si_insert_input_ret(ctx, ret, ctx->param_merged_scratch_offset, 5);
+
+ ret = si_insert_input_ptr_as_2xi32(ctx, ret, ctx->param_rw_buffers,
+ 8 + SI_SGPR_RW_BUFFERS);
ret = si_insert_input_ptr_as_2xi32(ctx, ret,
ctx->param_bindless_samplers_and_images,
8 + SI_SGPR_BINDLESS_SAMPLERS_AND_IMAGES);
{
LLVMValueRef ret = ctx->return_value;
- ret = si_insert_input_ptr_as_2xi32(ctx, ret, ctx->param_rw_buffers, 0);
ret = si_insert_input_ret(ctx, ret, ctx->param_gs2vs_offset, 2);
ret = si_insert_input_ret(ctx, ret, ctx->param_merged_wave_info, 3);
-
ret = si_insert_input_ret(ctx, ret, ctx->param_merged_scratch_offset, 5);
+
+ ret = si_insert_input_ptr_as_2xi32(ctx, ret, ctx->param_rw_buffers,
+ 8 + SI_SGPR_RW_BUFFERS);
ret = si_insert_input_ptr_as_2xi32(ctx, ret,
ctx->param_bindless_samplers_and_images,
8 + SI_SGPR_BINDLESS_SAMPLERS_AND_IMAGES);
ctx->return_value = ret;
}
-static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
+static void si_llvm_emit_ls_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 si_shader *shader = ctx->shader;
struct tgsi_shader_info *info = &shader->selector->info;
unsigned i, chan;
/* Write outputs to LDS. The next shader (TCS aka HS) will read
* its inputs from it. */
for (i = 0; i < info->num_outputs; i++) {
- LLVMValueRef *out_ptr = ctx->outputs[i];
unsigned name = info->output_semantic_name[i];
unsigned index = info->output_semantic_index[i];
LLVMConstInt(ctx->i32, param * 4, 0), "");
for (chan = 0; chan < 4; chan++) {
- lds_store(bld_base, chan, dw_addr,
- LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], ""));
+ if (!(info->output_usagemask[i] & (1 << chan)))
+ continue;
+
+ lds_store(ctx, chan, dw_addr,
+ LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], ""));
}
}
- if (ctx->screen->b.chip_class >= GFX9)
+ if (ctx->screen->info.chip_class >= GFX9)
si_set_ls_return_value_for_tcs(ctx);
}
-static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
+static void si_llvm_emit_es_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 si_shader *es = ctx->shader;
struct tgsi_shader_info *info = &es->selector->info;
LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
unsigned chan;
int i;
- if (ctx->screen->b.chip_class >= GFX9 && info->num_outputs) {
+ if (ctx->screen->info.chip_class >= GFX9 && info->num_outputs) {
unsigned itemsize_dw = es->selector->esgs_itemsize / 4;
LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
LLVMValueRef wave_idx = unpack_param(ctx, ctx->param_merged_wave_info, 24, 4);
}
for (i = 0; i < info->num_outputs; i++) {
- LLVMValueRef *out_ptr = ctx->outputs[i];
int param;
if (info->output_semantic_name[i] == TGSI_SEMANTIC_VIEWPORT_INDEX ||
info->output_semantic_index[i]);
for (chan = 0; chan < 4; chan++) {
- LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], "");
+ LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], "");
out_val = ac_to_integer(&ctx->ac, out_val);
/* GFX9 has the ESGS ring in LDS. */
- if (ctx->screen->b.chip_class >= GFX9) {
- lds_store(bld_base, param * 4 + chan, lds_base, out_val);
+ if (ctx->screen->info.chip_class >= GFX9) {
+ lds_store(ctx, param * 4 + chan, lds_base, out_val);
continue;
}
}
}
- if (ctx->screen->b.chip_class >= GFX9)
+ if (ctx->screen->info.chip_class >= GFX9)
si_set_es_return_value_for_gs(ctx);
}
static LLVMValueRef si_get_gs_wave_id(struct si_shader_context *ctx)
{
- if (ctx->screen->b.chip_class >= GFX9)
+ if (ctx->screen->info.chip_class >= GFX9)
return unpack_param(ctx, ctx->param_merged_wave_info, 16, 8);
else
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));
- if (ctx->screen->b.chip_class >= GFX9)
+ if (ctx->screen->info.chip_class >= GFX9)
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 si_shader_context *ctx = si_shader_context_from_abi(abi);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct tgsi_shader_info *info = &ctx->shader->selector->info;
struct si_shader_output_values *outputs = NULL;
int i,j;
ctx->param_vs_state_bits);
cond = LLVMBuildTrunc(ctx->ac.builder, cond,
ctx->i1, "");
- lp_build_if(&if_ctx, gallivm, cond);
+ lp_build_if(&if_ctx, &ctx->gallivm, cond);
}
for (j = 0; j < 4; j++) {
i++;
}
- si_llvm_export_vs(&ctx->bld_base, outputs, i);
+ si_llvm_export_vs(ctx, outputs, i);
FREE(outputs);
}
/* SI (except OLAND and HAINAN) has a bug that it only looks
* at the X writemask component. */
- if (ctx->screen->b.chip_class == SI &&
- ctx->screen->b.family != CHIP_OLAND &&
- ctx->screen->b.family != CHIP_HAINAN)
+ 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 */
bool is_last, struct si_ps_exports *exp)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct lp_build_context *base = &bld_base->base;
int i;
/* Clamp color */
/* Alpha to one */
if (ctx->shader->key.part.ps.epilog.alpha_to_one)
- color[3] = base->one;
+ color[3] = ctx->ac.f32_1;
/* Alpha test */
if (index == 0 &&
/* Get the export arguments, also find out what the last one is. */
for (c = 0; c <= ctx->shader->key.part.ps.epilog.last_cbuf; c++) {
- si_llvm_init_export_args(bld_base, color,
+ si_llvm_init_export_args(ctx, color,
V_008DFC_SQ_EXP_MRT + c, &args[c]);
if (args[c].enabled_channels)
last = c;
struct ac_export_args args;
/* Export */
- si_llvm_init_export_args(bld_base, color, V_008DFC_SQ_EXP_MRT + index,
+ si_llvm_init_export_args(ctx, color, V_008DFC_SQ_EXP_MRT + index,
&args);
if (is_last) {
args.valid_mask = 1; /* whether the EXEC mask is valid */
LLVMValueRef ret;
if (ctx->postponed_kill)
- ac_build_kill(&ctx->ac, LLVMBuildLoad(builder, ctx->postponed_kill, ""));
+ ac_build_kill_if_false(&ctx->ac, LLVMBuildLoad(builder, ctx->postponed_kill, ""));
/* Read the output values. */
for (i = 0; i < info->num_outputs; i++) {
LLVMValueRef interp_ij)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef result[4], a;
unsigned i;
result[2+i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDY, a);
}
- return lp_build_gather_values(gallivm, result, 4);
+ return lp_build_gather_values(&ctx->gallivm, result, 4);
}
static void interp_fetch_args(
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_instruction *inst = emit_data->inst;
if (inst->Instruction.Opcode == TGSI_OPCODE_INTERP_OFFSET) {
ctx->ac.f32_0,
};
- sample_position = lp_build_gather_values(gallivm, center, 4);
+ sample_position = lp_build_gather_values(&ctx->gallivm, center, 4);
} else {
sample_position = load_sample_position(ctx, sample_id);
}
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
- struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_shader_info *info = &shader->selector->info;
LLVMValueRef interp_param;
const struct tgsi_full_instruction *inst = emit_data->inst;
ij_out[i] = LLVMBuildFAdd(ctx->ac.builder, temp2, temp1, "");
}
- interp_param = lp_build_gather_values(gallivm, ij_out, 2);
+ interp_param = lp_build_gather_values(&ctx->gallivm, ij_out, 2);
}
if (interp_param)
}
/* Emit one vertex from the geometry shader */
-static void si_llvm_emit_vertex(
- const struct lp_build_tgsi_action *action,
- struct lp_build_tgsi_context *bld_base,
- struct lp_build_emit_data *emit_data)
+static void si_llvm_emit_vertex(struct ac_shader_abi *abi,
+ unsigned stream,
+ LLVMValueRef *addrs)
{
- struct si_shader_context *ctx = si_shader_context(bld_base);
- struct lp_build_context *uint = &bld_base->uint_bld;
+ 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 tgsi_shader_info *info = &shader->selector->info;
- struct gallivm_state *gallivm = &ctx->gallivm;
struct lp_build_if_state if_state;
LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
ctx->param_gs2vs_offset);
LLVMValueRef gs_next_vertex;
- LLVMValueRef can_emit, kill;
+ LLVMValueRef can_emit;
unsigned chan, offset;
int i;
- unsigned stream;
-
- stream = si_llvm_get_stream(bld_base, emit_data);
/* Write vertex attribute values to GSVS ring */
gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
bool use_kill = !info->writes_memory;
if (use_kill) {
- kill = lp_build_select(&bld_base->base, can_emit,
- LLVMConstReal(ctx->f32, 1.0f),
- LLVMConstReal(ctx->f32, -1.0f));
-
- ac_build_kill(&ctx->ac, kill);
+ ac_build_kill_if_false(&ctx->ac, can_emit);
} else {
- lp_build_if(&if_state, gallivm, can_emit);
+ lp_build_if(&if_state, &ctx->gallivm, can_emit);
}
offset = 0;
for (i = 0; i < info->num_outputs; i++) {
- LLVMValueRef *out_ptr = ctx->outputs[i];
-
for (chan = 0; chan < 4; chan++) {
if (!(info->output_usagemask[i] & (1 << chan)) ||
((info->output_streams[i] >> (2 * chan)) & 3) != stream)
continue;
- LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], "");
+ LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], "");
LLVMValueRef voffset =
LLVMConstInt(ctx->i32, offset *
shader->selector->gs_max_out_vertices, 0);
lp_build_endif(&if_state);
}
+/* Emit one vertex from the geometry shader */
+static void si_tgsi_emit_vertex(
+ const struct lp_build_tgsi_action *action,
+ struct lp_build_tgsi_context *bld_base,
+ struct lp_build_emit_data *emit_data)
+{
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ unsigned stream = si_llvm_get_stream(bld_base, emit_data);
+
+ si_llvm_emit_vertex(&ctx->abi, stream, ctx->outputs[0]);
+}
+
/* Cut one primitive from the geometry shader */
static void si_llvm_emit_primitive(
const struct lp_build_tgsi_action *action,
* The real barrier instruction isn’t needed, because an entire patch
* always fits into a single wave.
*/
- if (ctx->screen->b.chip_class == SI &&
+ if (ctx->screen->info.chip_class == SI &&
ctx->type == PIPE_SHADER_TESS_CTRL) {
si_emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
return;
"no-signed-zeros-fp-math",
"true");
- if (ctx->screen->b.debug_flags & DBG_UNSAFE_MATH) {
+ if (ctx->screen->debug_flags & DBG(UNSAFE_MATH)) {
/* These were copied from some LLVM test. */
LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
"less-precise-fpmad",
}
}
-static void declare_lds_as_pointer(struct si_shader_context *ctx)
-{
- unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768;
- ctx->lds = LLVMBuildIntToPtr(ctx->ac.builder, ctx->i32_0,
- LLVMPointerType(LLVMArrayType(ctx->i32, lds_size / 4), LOCAL_ADDR_SPACE),
- "lds");
-}
-
static unsigned si_get_max_workgroup_size(const struct si_shader *shader)
{
switch (shader->selector->type) {
case PIPE_SHADER_TESS_CTRL:
/* Return this so that LLVM doesn't remove s_barrier
* instructions on chips where we use s_barrier. */
- return shader->selector->screen->b.chip_class >= CIK ? 128 : 64;
+ return shader->selector->screen->info.chip_class >= CIK ? 128 : 64;
case PIPE_SHADER_GEOMETRY:
- return shader->selector->screen->b.chip_class >= GFX9 ? 128 : 64;
+ return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 64;
case PIPE_SHADER_COMPUTE:
break; /* see below */
struct si_function_info *fninfo,
bool assign_params)
{
+ LLVMTypeRef const_shader_buf_type;
+
+ if (ctx->shader->selector->info.const_buffers_declared == 1 &&
+ ctx->shader->selector->info.shader_buffers_declared == 0)
+ const_shader_buf_type = ctx->f32;
+ else
+ const_shader_buf_type = ctx->v4i32;
+
unsigned const_and_shader_buffers =
add_arg(fninfo, ARG_SGPR,
- si_const_array(ctx->v4i32,
- SI_NUM_SHADER_BUFFERS + SI_NUM_CONST_BUFFERS));
+ si_const_array(const_shader_buf_type, 0));
+
unsigned samplers_and_images =
add_arg(fninfo, ARG_SGPR,
si_const_array(ctx->v8i32,
}
}
-static void declare_default_desc_pointers(struct si_shader_context *ctx,
- struct si_function_info *fninfo)
+static void declare_global_desc_pointers(struct si_shader_context *ctx,
+ struct si_function_info *fninfo)
{
ctx->param_rw_buffers = add_arg(fninfo, ARG_SGPR,
si_const_array(ctx->v4i32, SI_NUM_RW_BUFFERS));
ctx->param_bindless_samplers_and_images = add_arg(fninfo, ARG_SGPR,
si_const_array(ctx->v8i32, 0));
- declare_per_stage_desc_pointers(ctx, fninfo, true);
}
static void declare_vs_specific_input_sgprs(struct si_shader_context *ctx,
unsigned num_returns = 0;
unsigned num_prolog_vgprs = 0;
unsigned type = ctx->type;
+ unsigned vs_blit_property =
+ shader->selector->info.properties[TGSI_PROPERTY_VS_BLIT_SGPRS];
si_init_function_info(&fninfo);
/* Set MERGED shaders. */
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
if (shader->key.as_ls || type == PIPE_SHADER_TESS_CTRL)
type = SI_SHADER_MERGED_VERTEX_TESSCTRL; /* LS or HS */
else if (shader->key.as_es || type == PIPE_SHADER_GEOMETRY)
switch (type) {
case PIPE_SHADER_VERTEX:
- declare_default_desc_pointers(ctx, &fninfo);
+ declare_global_desc_pointers(ctx, &fninfo);
+
+ if (vs_blit_property) {
+ ctx->param_vs_blit_inputs = fninfo.num_params;
+ add_arg(&fninfo, ARG_SGPR, ctx->i32); /* i16 x1, y1 */
+ add_arg(&fninfo, ARG_SGPR, ctx->i32); /* i16 x2, y2 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* depth */
+
+ if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_COLOR) {
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color0 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color1 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color2 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* color3 */
+ } else if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_TEXCOORD) {
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.x1 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.y1 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.x2 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.y2 */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.z */
+ add_arg(&fninfo, ARG_SGPR, ctx->f32); /* texcoord.w */
+ }
+
+ /* VGPRs */
+ declare_vs_input_vgprs(ctx, &fninfo, &num_prolog_vgprs);
+ break;
+ }
+
+ declare_per_stage_desc_pointers(ctx, &fninfo, true);
declare_vs_specific_input_sgprs(ctx, &fninfo);
if (shader->key.as_es) {
- assert(!shader->selector->nir);
ctx->param_es2gs_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
} else if (shader->key.as_ls) {
- assert(!shader->selector->nir);
/* no extra parameters */
} else {
if (shader->is_gs_copy_shader) {
break;
case PIPE_SHADER_TESS_CTRL: /* SI-CI-VI */
- declare_default_desc_pointers(ctx, &fninfo);
+ declare_global_desc_pointers(ctx, &fninfo);
+ declare_per_stage_desc_pointers(ctx, &fninfo, true);
ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_out_lds_offsets = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_out_lds_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
case SI_SHADER_MERGED_VERTEX_TESSCTRL:
/* Merged stages have 8 system SGPRs at the beginning. */
- ctx->param_rw_buffers = /* SPI_SHADER_USER_DATA_ADDR_LO_HS */
- add_arg(&fninfo, ARG_SGPR, si_const_array(ctx->v4i32, SI_NUM_RW_BUFFERS));
+ add_arg(&fninfo, ARG_SGPR, ctx->i32); /* SPI_SHADER_USER_DATA_ADDR_LO_HS */
+ add_arg(&fninfo, ARG_SGPR, ctx->i32); /* SPI_SHADER_USER_DATA_ADDR_HI_HS */
ctx->param_tcs_offchip_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_merged_wave_info = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_factor_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
- add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
- add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
-
- ctx->param_bindless_samplers_and_images =
- add_arg(&fninfo, ARG_SGPR, si_const_array(ctx->v8i32, 0));
-
+ declare_global_desc_pointers(ctx, &fninfo);
declare_per_stage_desc_pointers(ctx, &fninfo,
ctx->type == PIPE_SHADER_VERTEX);
declare_vs_specific_input_sgprs(ctx, &fninfo);
case SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY:
/* Merged stages have 8 system SGPRs at the beginning. */
- ctx->param_rw_buffers = /* SPI_SHADER_USER_DATA_ADDR_LO_GS */
- add_arg(&fninfo, ARG_SGPR, si_const_array(ctx->v4i32, SI_NUM_RW_BUFFERS));
+ add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused (SPI_SHADER_USER_DATA_ADDR_LO_GS) */
+ add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused (SPI_SHADER_USER_DATA_ADDR_HI_GS) */
ctx->param_gs2vs_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_merged_wave_info = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_offchip_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused (SPI_SHADER_PGM_LO/HI_GS << 8) */
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused (SPI_SHADER_PGM_LO/HI_GS >> 24) */
- add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
- add_arg(&fninfo, ARG_SGPR, ctx->i32); /* unused */
-
- ctx->param_bindless_samplers_and_images =
- add_arg(&fninfo, ARG_SGPR, si_const_array(ctx->v8i32, 0));
-
+ declare_global_desc_pointers(ctx, &fninfo);
declare_per_stage_desc_pointers(ctx, &fninfo,
(ctx->type == PIPE_SHADER_VERTEX ||
ctx->type == PIPE_SHADER_TESS_EVAL));
/* VGPRs (first GS, then VS/TES) */
ctx->param_gs_vtx01_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
ctx->param_gs_vtx23_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_prim_id = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_instance_id = add_arg(&fninfo, ARG_VGPR, ctx->i32);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.gs_prim_id);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.gs_invocation_id);
ctx->param_gs_vtx45_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
if (ctx->type == PIPE_SHADER_VERTEX) {
break;
case PIPE_SHADER_TESS_EVAL:
- declare_default_desc_pointers(ctx, &fninfo);
+ declare_global_desc_pointers(ctx, &fninfo);
+ declare_per_stage_desc_pointers(ctx, &fninfo, true);
ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_tcs_offchip_addr_base64k = add_arg(&fninfo, ARG_SGPR, ctx->i32);
break;
case PIPE_SHADER_GEOMETRY:
- declare_default_desc_pointers(ctx, &fninfo);
+ declare_global_desc_pointers(ctx, &fninfo);
+ declare_per_stage_desc_pointers(ctx, &fninfo, true);
ctx->param_gs2vs_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
ctx->param_gs_wave_id = add_arg(&fninfo, ARG_SGPR, ctx->i32);
/* VGPRs */
- ctx->param_gs_vtx0_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_vtx1_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_prim_id = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_vtx2_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_vtx3_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_vtx4_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_vtx5_offset = add_arg(&fninfo, ARG_VGPR, ctx->i32);
- ctx->param_gs_instance_id = add_arg(&fninfo, ARG_VGPR, ctx->i32);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->gs_vtx_offset[0]);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->gs_vtx_offset[1]);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.gs_prim_id);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->gs_vtx_offset[2]);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->gs_vtx_offset[3]);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->gs_vtx_offset[4]);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->gs_vtx_offset[5]);
+ add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.gs_invocation_id);
break;
case PIPE_SHADER_FRAGMENT:
- declare_default_desc_pointers(ctx, &fninfo);
+ declare_global_desc_pointers(ctx, &fninfo);
+ declare_per_stage_desc_pointers(ctx, &fninfo, true);
add_arg_checked(&fninfo, ARG_SGPR, ctx->f32, SI_PARAM_ALPHA_REF);
add_arg_checked(&fninfo, ARG_SGPR, ctx->i32, SI_PARAM_PRIM_MASK);
break;
case PIPE_SHADER_COMPUTE:
- declare_default_desc_pointers(ctx, &fninfo);
+ declare_global_desc_pointers(ctx, &fninfo);
+ declare_per_stage_desc_pointers(ctx, &fninfo, true);
if (shader->selector->info.uses_grid_size)
ctx->param_grid_size = add_arg(&fninfo, ARG_SGPR, v3i32);
if (shader->selector->info.uses_block_size)
if (shader->key.as_ls ||
ctx->type == PIPE_SHADER_TESS_CTRL ||
/* GFX9 has the ESGS ring buffer in LDS. */
- (ctx->screen->b.chip_class >= GFX9 &&
- (shader->key.as_es ||
- ctx->type == PIPE_SHADER_GEOMETRY)))
- declare_lds_as_pointer(ctx);
+ type == SI_SHADER_MERGED_VERTEX_OR_TESSEVAL_GEOMETRY)
+ ac_declare_lds_as_pointer(&ctx->ac);
}
/**
LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_rw_buffers);
- if (ctx->screen->b.chip_class <= VI &&
+ if (ctx->screen->info.chip_class <= VI &&
(ctx->shader->key.as_es || ctx->type == PIPE_SHADER_GEOMETRY)) {
unsigned ring =
ctx->type == PIPE_SHADER_GEOMETRY ? SI_GS_RING_ESGS
LLVMValueRef offset = LLVMConstInt(ctx->i32, ring, 0);
ctx->esgs_ring =
- ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
+ ac_build_load_to_sgpr(&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);
+ ac_build_load_to_sgpr(&ctx->ac, buf_ptr, offset);
} else if (ctx->type == PIPE_SHADER_GEOMETRY) {
const struct si_shader_selector *sel = ctx->shader->selector;
LLVMValueRef offset = LLVMConstInt(ctx->i32, SI_RING_GSVS, 0);
LLVMValueRef base_ring;
- base_ring = ac_build_indexed_load_const(&ctx->ac, buf_ptr, offset);
+ base_ring = ac_build_load_to_sgpr(&ctx->ac, buf_ptr, offset);
/* The conceptual layout of the GSVS ring is
* v0c0 .. vLv0 v0c1 .. vLc1 ..
/* Load the buffer descriptor. */
slot = LLVMConstInt(ctx->i32, SI_PS_CONST_POLY_STIPPLE, 0);
- desc = ac_build_indexed_load_const(&ctx->ac, param_rw_buffers, slot);
+ desc = ac_build_load_to_sgpr(&ctx->ac, param_rw_buffers, slot);
/* The stipple pattern is 32x32, each row has 32 bits. */
offset = LLVMBuildMul(builder, address[1],
row = ac_to_integer(&ctx->ac, row);
bit = LLVMBuildLShr(builder, row, address[0], "");
bit = LLVMBuildTrunc(builder, bit, ctx->i1, "");
-
- /* The intrinsic kills the thread if arg < 0. */
- bit = LLVMBuildSelect(builder, bit, LLVMConstReal(ctx->f32, 0),
- LLVMConstReal(ctx->f32, -1), "");
- ac_build_kill(&ctx->ac, bit);
+ ac_build_kill_if_false(&ctx->ac, bit);
}
void si_shader_binary_read_config(struct ac_shader_binary *binary,
r600_resource_reference(&shader->bo, NULL);
shader->bo = (struct r600_resource*)
- pipe_buffer_create(&sscreen->b.b, 0,
+ pipe_buffer_create(&sscreen->b, 0,
PIPE_USAGE_IMMUTABLE,
align(bo_size, SI_CPDMA_ALIGNMENT));
if (!shader->bo)
return -ENOMEM;
/* Upload. */
- ptr = sscreen->b.ws->buffer_map(shader->bo->buf, NULL,
+ ptr = sscreen->ws->buffer_map(shader->bo->buf, NULL,
PIPE_TRANSFER_READ_WRITE |
PIPE_TRANSFER_UNSYNCHRONIZED);
else if (mainb->rodata_size > 0)
memcpy(ptr, mainb->rodata, mainb->rodata_size);
- sscreen->b.ws->buffer_unmap(shader->bo->buf);
+ sscreen->ws->buffer_unmap(shader->bo->buf);
return 0;
}
const struct si_shader_config *conf = &shader->config;
unsigned num_inputs = shader->selector ? shader->selector->info.num_inputs : 0;
unsigned code_size = si_get_shader_binary_size(shader);
- unsigned lds_increment = sscreen->b.chip_class >= CIK ? 512 : 256;
+ unsigned lds_increment = sscreen->info.chip_class >= CIK ? 512 : 256;
unsigned lds_per_wave = 0;
unsigned max_simd_waves;
- switch (sscreen->b.family) {
+ switch (sscreen->info.family) {
/* These always have 8 waves: */
case CHIP_POLARIS10:
case CHIP_POLARIS11:
/* Compute the per-SIMD wave counts. */
if (conf->num_sgprs) {
- if (sscreen->b.chip_class >= VI)
+ if (sscreen->info.chip_class >= VI)
max_simd_waves = MIN2(max_simd_waves, 800 / conf->num_sgprs);
else
max_simd_waves = MIN2(max_simd_waves, 512 / conf->num_sgprs);
max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave);
if (!check_debug_option ||
- si_can_dump_shader(&sscreen->b, processor)) {
+ si_can_dump_shader(sscreen, processor)) {
if (processor == PIPE_SHADER_FRAGMENT) {
fprintf(file, "*** SHADER CONFIG ***\n"
"SPI_PS_INPUT_ADDR = 0x%04x\n"
FILE *file, bool check_debug_option)
{
if (!check_debug_option ||
- si_can_dump_shader(&sscreen->b, processor))
+ si_can_dump_shader(sscreen, processor))
si_dump_shader_key(processor, shader, file);
if (!check_debug_option && shader->binary.llvm_ir_string) {
}
if (!check_debug_option ||
- (si_can_dump_shader(&sscreen->b, processor) &&
- !(sscreen->b.debug_flags & DBG_NO_ASM))) {
+ (si_can_dump_shader(sscreen, processor) &&
+ !(sscreen->debug_flags & DBG(NO_ASM)))) {
fprintf(file, "\n%s:\n", si_get_shader_name(shader, processor));
if (shader->prolog)
const char *name)
{
int r = 0;
- unsigned count = p_atomic_inc_return(&sscreen->b.num_compilations);
+ unsigned count = p_atomic_inc_return(&sscreen->num_compilations);
- if (si_can_dump_shader(&sscreen->b, processor)) {
+ if (si_can_dump_shader(sscreen, processor)) {
fprintf(stderr, "radeonsi: Compiling shader %d\n", count);
- if (!(sscreen->b.debug_flags & (DBG_NO_IR | DBG_PREOPT_IR))) {
+ if (!(sscreen->debug_flags & (DBG(NO_IR) | DBG(PREOPT_IR)))) {
fprintf(stderr, "%s LLVM IR:\n\n", name);
ac_dump_module(mod);
fprintf(stderr, "\n");
{
struct si_shader_context ctx;
struct si_shader *shader;
- struct gallivm_state *gallivm = &ctx.gallivm;
LLVMBuilderRef builder;
struct lp_build_tgsi_context *bld_base = &ctx.bld_base;
struct lp_build_context *uint = &bld_base->uint_bld;
return NULL;
}
+ /* We can leave the fence as permanently signaled because the GS copy
+ * shader only becomes visible globally after it has been compiled. */
+ util_queue_fence_init(&shader->ready);
shader->selector = gs_selector;
shader->is_gs_copy_shader = true;
LLVMBasicBlockRef end_bb;
LLVMValueRef switch_inst;
- end_bb = LLVMAppendBasicBlockInContext(gallivm->context, ctx.main_fn, "end");
+ end_bb = LLVMAppendBasicBlockInContext(ctx.ac.context, ctx.main_fn, "end");
switch_inst = LLVMBuildSwitch(builder, stream_id, end_bb, 4);
for (int stream = 0; stream < 4; stream++) {
if (stream > 0 && !gs_selector->so.num_outputs)
continue;
- bb = LLVMInsertBasicBlockInContext(gallivm->context, end_bb, "out");
+ bb = LLVMInsertBasicBlockInContext(ctx.ac.context, end_bb, "out");
LLVMAddCase(switch_inst, LLVMConstInt(ctx.i32, stream, 0), bb);
LLVMPositionBuilderAtEnd(builder, bb);
}
if (stream == 0)
- si_llvm_export_vs(bld_base, outputs, gsinfo->num_outputs);
+ si_llvm_export_vs(&ctx, outputs, gsinfo->num_outputs);
LLVMBuildBr(builder, end_bb);
}
debug, PIPE_SHADER_GEOMETRY,
"GS Copy Shader");
if (!r) {
- if (si_can_dump_shader(&sscreen->b, PIPE_SHADER_GEOMETRY))
+ if (si_can_dump_shader(sscreen, PIPE_SHADER_GEOMETRY))
fprintf(stderr, "GS Copy Shader:\n");
si_shader_dump(sscreen, ctx.shader, debug,
PIPE_SHADER_GEOMETRY, stderr, true);
break;
case PIPE_SHADER_TESS_CTRL:
- if (shader->selector->screen->b.chip_class >= GFX9) {
+ if (shader->selector->screen->info.chip_class >= GFX9) {
si_dump_shader_key_vs(key, &key->part.tcs.ls_prolog,
"part.tcs.ls_prolog", f);
}
if (shader->is_gs_copy_shader)
break;
- if (shader->selector->screen->b.chip_class >= GFX9 &&
+ if (shader->selector->screen->info.chip_class >= GFX9 &&
key->part.gs.es->type == PIPE_SHADER_VERTEX) {
si_dump_shader_key_vs(key, &key->part.gs.vs_prolog,
"part.gs.vs_prolog", f);
bld_base->op_actions[TGSI_OPCODE_READ_INVOC].fetch_args = read_invoc_fetch_args;
bld_base->op_actions[TGSI_OPCODE_READ_INVOC].emit = read_lane_emit;
- bld_base->op_actions[TGSI_OPCODE_EMIT].emit = si_llvm_emit_vertex;
+ 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_BARRIER].emit = si_llvm_emit_barrier;
}
case PIPE_SHADER_VERTEX:
ctx->load_input = declare_input_vs;
if (shader->key.as_ls)
- bld_base->emit_epilogue = si_llvm_emit_ls_epilogue;
+ ctx->abi.emit_outputs = si_llvm_emit_ls_epilogue;
else if (shader->key.as_es)
- bld_base->emit_epilogue = si_llvm_emit_es_epilogue;
- else {
+ ctx->abi.emit_outputs = si_llvm_emit_es_epilogue;
+ else
ctx->abi.emit_outputs = si_llvm_emit_vs_epilogue;
- bld_base->emit_epilogue = si_tgsi_emit_epilogue;
- }
+ bld_base->emit_epilogue = si_tgsi_emit_epilogue;
break;
case PIPE_SHADER_TESS_CTRL:
bld_base->emit_fetch_funcs[TGSI_FILE_INPUT] = fetch_input_tcs;
case PIPE_SHADER_TESS_EVAL:
bld_base->emit_fetch_funcs[TGSI_FILE_INPUT] = fetch_input_tes;
if (shader->key.as_es)
- bld_base->emit_epilogue = si_llvm_emit_es_epilogue;
- else {
+ ctx->abi.emit_outputs = si_llvm_emit_es_epilogue;
+ else
ctx->abi.emit_outputs = si_llvm_emit_vs_epilogue;
- bld_base->emit_epilogue = si_tgsi_emit_epilogue;
- }
+ bld_base->emit_epilogue = si_tgsi_emit_epilogue;
break;
case PIPE_SHADER_GEOMETRY:
bld_base->emit_fetch_funcs[TGSI_FILE_INPUT] = fetch_input_gs;
- bld_base->emit_epilogue = si_llvm_emit_gs_epilogue;
+ ctx->abi.emit_vertex = si_llvm_emit_vertex;
+ 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;
* For monolithic merged shaders, the first shader is wrapped in an
* if-block together with its prolog in si_build_wrapper_function.
*/
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
if (!is_monolithic &&
sel->info.num_instructions > 1 && /* not empty shader */
(shader->key.as_es || shader->key.as_ls) &&
}
}
- if (ctx->type == PIPE_SHADER_FRAGMENT && sel->info.uses_kill &&
- ctx->screen->b.debug_flags & DBG_FS_CORRECT_DERIVS_AFTER_KILL) {
- /* This is initialized to 0.0 = not kill. */
- ctx->postponed_kill = lp_build_alloca(&ctx->gallivm, ctx->f32, "");
+ if (sel->force_correct_derivs_after_kill) {
+ ctx->postponed_kill = lp_build_alloca_undef(&ctx->gallivm, ctx->i1, "");
+ /* true = don't kill. */
+ LLVMBuildStore(ctx->ac.builder, LLVMConstInt(ctx->i1, 1, 0),
+ ctx->postponed_kill);
}
if (sel->tokens) {
key->vs_prolog.num_input_sgprs = num_input_sgprs;
key->vs_prolog.last_input = MAX2(1, info->num_inputs) - 1;
key->vs_prolog.as_ls = shader_out->key.as_ls;
+ key->vs_prolog.as_es = shader_out->key.as_es;
if (shader_out->selector->type == PIPE_SHADER_TESS_CTRL) {
key->vs_prolog.as_ls = 1;
key->vs_prolog.num_merged_next_stage_vgprs = 2;
} else if (shader_out->selector->type == PIPE_SHADER_GEOMETRY) {
+ key->vs_prolog.as_es = 1;
key->vs_prolog.num_merged_next_stage_vgprs = 5;
}
si_init_function_info(&fninfo);
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
num_sgprs = 8 + GFX9_GS_NUM_USER_SGPR;
num_vgprs = 5; /* ES inputs are not needed by GS */
} else {
* with registers here. The main shader part will set the correct EXEC
* mask.
*/
- if (ctx->screen->b.chip_class >= GFX9 && !key->gs_prolog.is_monolithic)
+ if (ctx->screen->info.chip_class >= GFX9 && !key->gs_prolog.is_monolithic)
si_init_exec_full_mask(ctx);
/* Copy inputs to outputs. This should be no-op, as the registers match,
LLVMValueRef vtx_in[6], vtx_out[6];
LLVMValueRef prim_id, rotate;
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
for (unsigned i = 0; i < 3; i++) {
vtx_in[i*2] = unpack_param(ctx, gfx9_vtx_params[i], 0, 16);
vtx_in[i*2+1] = unpack_param(ctx, gfx9_vtx_params[i], 16, 16);
vtx_out[i] = LLVMBuildSelect(builder, rotate, rotated, base, "");
}
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
for (unsigned i = 0; i < 3; i++) {
LLVMValueRef hi, out;
unsigned main_part,
unsigned next_shader_first_part)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMBuilderRef builder = ctx->ac.builder;
/* PS epilog has one arg per color component; gfx9 merged shader
* prologs need to forward 32 user SGPRs.
if (param_size == 1)
arg = out[out_idx];
else
- arg = lp_build_gather_values(gallivm, &out[out_idx], param_size);
+ arg = lp_build_gather_values(&ctx->gallivm, &out[out_idx], param_size);
if (LLVMTypeOf(arg) != param_type) {
if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
/* Dump TGSI code before doing TGSI->LLVM conversion in case the
* conversion fails. */
- if (si_can_dump_shader(&sscreen->b, sel->info.processor) &&
- !(sscreen->b.debug_flags & DBG_NO_TGSI)) {
+ if (si_can_dump_shader(sscreen, sel->info.processor) &&
+ !(sscreen->debug_flags & DBG(NO_TGSI))) {
if (sel->tokens)
tgsi_dump(sel->tokens, 0);
else
si_build_wrapper_function(&ctx, parts + !need_prolog,
1 + need_prolog, need_prolog, 0);
} else if (is_monolithic && ctx.type == PIPE_SHADER_TESS_CTRL) {
- if (sscreen->b.chip_class >= GFX9) {
+ if (sscreen->info.chip_class >= GFX9) {
struct si_shader_selector *ls = shader->key.part.tcs.ls;
LLVMValueRef parts[4];
bool vs_needs_prolog =
si_build_wrapper_function(&ctx, parts, 2, 0, 0);
}
} else if (is_monolithic && ctx.type == PIPE_SHADER_GEOMETRY) {
- if (ctx.screen->b.chip_class >= GFX9) {
+ if (ctx.screen->info.chip_class >= GFX9) {
struct si_shader_selector *es = shader->key.part.gs.es;
LLVMValueRef es_prolog = NULL;
LLVMValueRef es_main = NULL;
union si_shader_part_key vs_prolog_key;
si_get_vs_prolog_key(&es->info,
shader->info.num_input_sgprs,
- &shader->key.part.tcs.ls_prolog,
+ &shader->key.part.gs.vs_prolog,
shader, &vs_prolog_key);
vs_prolog_key.vs_prolog.is_monolithic = true;
si_build_vs_prolog_function(&ctx, &vs_prolog_key);
si_optimize_vs_outputs(&ctx);
if ((debug && debug->debug_message) ||
- si_can_dump_shader(&sscreen->b, ctx.type))
+ si_can_dump_shader(sscreen, ctx.type))
si_count_scratch_private_memory(&ctx);
/* Compile to bytecode. */
if (sel->type == PIPE_SHADER_COMPUTE) {
unsigned wave_size = 64;
unsigned max_vgprs = 256;
- unsigned max_sgprs = sscreen->b.chip_class >= VI ? 800 : 512;
+ unsigned max_sgprs = sscreen->info.chip_class >= VI ? 800 : 512;
unsigned max_sgprs_per_wave = 128;
unsigned max_block_threads = si_get_max_workgroup_size(shader);
unsigned min_waves_per_cu = DIV_ROUND_UP(max_block_threads, wave_size);
struct si_shader shader = {};
struct si_shader_context ctx;
- struct gallivm_state *gallivm = &ctx.gallivm;
si_init_shader_ctx(&ctx, sscreen, tm);
ctx.shader = &shader;
switch (type) {
case PIPE_SHADER_VERTEX:
+ shader.key.as_ls = key->vs_prolog.as_ls;
+ shader.key.as_es = key->vs_prolog.as_es;
break;
case PIPE_SHADER_TESS_CTRL:
assert(!prolog);
si_llvm_optimize_module(&ctx);
if (si_compile_llvm(sscreen, &result->binary, &result->config, tm,
- gallivm->module, debug, ctx.type, name)) {
+ ctx.ac.module, debug, ctx.type, name)) {
FREE(result);
result = NULL;
goto out;
static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef ptr[2], list;
+ bool is_merged_shader =
+ ctx->screen->info.chip_class >= GFX9 &&
+ (ctx->type == PIPE_SHADER_TESS_CTRL ||
+ ctx->type == PIPE_SHADER_GEOMETRY ||
+ ctx->shader->key.as_ls || ctx->shader->key.as_es);
/* Get the pointer to rw buffers. */
- ptr[0] = LLVMGetParam(ctx->main_fn, SI_SGPR_RW_BUFFERS);
- ptr[1] = LLVMGetParam(ctx->main_fn, SI_SGPR_RW_BUFFERS_HI);
- list = lp_build_gather_values(gallivm, ptr, 2);
+ 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_HI);
+ list = lp_build_gather_values(&ctx->gallivm, ptr, 2);
list = LLVMBuildBitCast(ctx->ac.builder, list, ctx->i64, "");
list = LLVMBuildIntToPtr(ctx->ac.builder, list,
si_const_array(ctx->v4i32, SI_NUM_RW_BUFFERS), "");
si_init_exec_from_input(ctx, 3, 0);
if (key->vs_prolog.as_ls &&
- (ctx->screen->b.family == CHIP_VEGA10 ||
- ctx->screen->b.family == CHIP_RAVEN)) {
+ ctx->screen->has_ls_vgpr_init_bug) {
/* If there are no HS threads, SPI loads the LS VGPRs
* starting at VGPR 0. Shift them back to where they
* belong.
LLVMValueRef buf_index =
LLVMConstInt(ctx->i32, SI_VS_CONST_INSTANCE_DIVISORS, 0);
instance_divisor_constbuf =
- ac_build_indexed_load_const(&ctx->ac, list, buf_index);
+ ac_build_load_to_sgpr(&ctx->ac, list, buf_index);
}
for (i = 0; i <= key->vs_prolog.last_input; i++) {
si_init_function_info(&fninfo);
- if (ctx->screen->b.chip_class >= GFX9) {
+ if (ctx->screen->info.chip_class >= GFX9) {
add_arg(&fninfo, ARG_SGPR, ctx->i64);
ctx->param_tcs_offchip_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* wave info */
/* Create the function. */
si_create_function(ctx, "tcs_epilog", NULL, 0, &fninfo,
- ctx->screen->b.chip_class >= CIK ? 128 : 64);
- declare_lds_as_pointer(ctx);
+ ctx->screen->info.chip_class >= CIK ? 128 : 64);
+ ac_declare_lds_as_pointer(&ctx->ac);
func = ctx->main_fn;
LLVMValueRef invoc0_tess_factors[6];
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
- if (sscreen->b.chip_class >= GFX9) {
+ if (sscreen->info.chip_class >= GFX9) {
struct si_shader *ls_main_part =
shader->key.part.tcs.ls->main_shader_part_ls;
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
- if (sscreen->b.chip_class >= GFX9) {
+ if (sscreen->info.chip_class >= GFX9) {
struct si_shader *es_main_part =
shader->key.part.gs.es->main_shader_part_es;
static void si_build_ps_prolog_function(struct si_shader_context *ctx,
union si_shader_part_key *key)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_function_info fninfo;
LLVMValueRef ret, func;
int num_returns, i, num_color_channels;
interp_vgpr, "");
interp[1] = LLVMBuildExtractValue(ctx->ac.builder, ret,
interp_vgpr + 1, "");
- interp_ij = lp_build_gather_values(gallivm, interp, 2);
+ interp_ij = lp_build_gather_values(&ctx->gallivm, interp, 2);
}
/* Use the absolute location of the input. */
* Make sure we have at least 4k of LDS in use to avoid the bug.
* It applies to workgroup sizes of more than one wavefront.
*/
- if (sscreen->b.family == CHIP_BONAIRE ||
- sscreen->b.family == CHIP_KABINI ||
- sscreen->b.family == CHIP_MULLINS)
+ if (sscreen->info.family == CHIP_BONAIRE ||
+ sscreen->info.family == CHIP_KABINI ||
+ sscreen->info.family == CHIP_MULLINS)
*lds_size = MAX2(*lds_size, 8);
}
r600_resource_reference(&shader->bo, NULL);
if (!shader->is_binary_shared)
- si_radeon_shader_binary_clean(&shader->binary);
+ ac_shader_binary_clean(&shader->binary);
free(shader->shader_log);
}
projects / mesa.git / blobdiff