X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;ds=sidebyside;f=src%2Fgallium%2Fdrivers%2Fradeonsi%2Fsi_shader.c;h=c3b5f58cd26a9bc788bdc0aea401f898a76f8ebd;hb=7ef1e42c14fb23592e8e003f7a80db9a43cb9bc9;hp=3f4d847d665c05c4b2c67d46288382d387f52bf9;hpb=bbfad34606b8fc91c3671bb9a6e6545d669e6f81;p=mesa.git diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 3f4d847d665..c3b5f58cd26 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -19,11 +19,6 @@ * 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 - * Michel Dänzer - * Christian König */ #include "gallivm/lp_bld_const.h" @@ -89,8 +84,6 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action, static void si_dump_shader_key(unsigned processor, const struct si_shader *shader, FILE *f); -static unsigned llvm_get_type_size(LLVMTypeRef type); - static void si_build_vs_prolog_function(struct si_shader_context *ctx, union si_shader_part_key *key); static void si_build_tcs_epilog_function(struct si_shader_context *ctx, @@ -100,10 +93,10 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, static void si_build_ps_epilog_function(struct si_shader_context *ctx, union si_shader_part_key *key); -/* Ideally pass the sample mask input to the PS epilog as v13, which +/* Ideally pass the sample mask input to the PS epilog as v14, which * is its usual location, so that the shader doesn't have to add v_mov. */ -#define PS_EPILOG_SAMPLEMASK_MIN_LOC 13 +#define PS_EPILOG_SAMPLEMASK_MIN_LOC 14 enum { CONST_ADDR_SPACE = 2, @@ -112,7 +105,7 @@ enum { 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 || @@ -235,20 +228,6 @@ unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index) } } -/** - * Helper function that builds an LLVM IR PHI node and immediately adds - * incoming edges. - */ -static LLVMValueRef -build_phi(struct ac_llvm_context *ctx, LLVMTypeRef type, - unsigned count_incoming, LLVMValueRef *values, - LLVMBasicBlockRef *blocks) -{ - LLVMValueRef phi = LLVMBuildPhi(ctx->builder, type, ""); - LLVMAddIncoming(phi, values, blocks, count_incoming); - return phi; -} - /** * Get the value of a shader input parameter and extract a bitfield. */ @@ -256,21 +235,19 @@ static LLVMValueRef unpack_param(struct si_shader_context *ctx, unsigned param, unsigned rshift, unsigned bitwidth) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef value = LLVMGetParam(ctx->main_fn, param); if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind) - value = bitcast(&ctx->bld_base, - TGSI_TYPE_UNSIGNED, value); + value = ac_to_integer(&ctx->ac, value); if (rshift) - value = LLVMBuildLShr(gallivm->builder, value, + value = LLVMBuildLShr(ctx->ac.builder, value, LLVMConstInt(ctx->i32, rshift, 0), ""); if (rshift + bitwidth < 32) { unsigned mask = (1 << bitwidth) - 1; - value = LLVMBuildAnd(gallivm->builder, value, + value = LLVMBuildAnd(ctx->ac.builder, value, LLVMConstInt(ctx->i32, mask, 0), ""); } @@ -320,10 +297,35 @@ get_tcs_in_patch_stride(struct si_shader_context *ctx) return unpack_param(ctx, ctx->param_vs_state_bits, 8, 13); } -static LLVMValueRef -get_tcs_out_patch_stride(struct si_shader_context *ctx) +static unsigned get_tcs_out_vertex_dw_stride_constant(struct si_shader_context *ctx) { - return unpack_param(ctx, ctx->param_tcs_out_lds_layout, 0, 13); + assert(ctx->type == PIPE_SHADER_TESS_CTRL); + + if (ctx->shader->key.mono.u.ff_tcs_inputs_to_copy) + return util_last_bit64(ctx->shader->key.mono.u.ff_tcs_inputs_to_copy) * 4; + + return util_last_bit64(ctx->shader->selector->outputs_written) * 4; +} + +static LLVMValueRef get_tcs_out_vertex_dw_stride(struct si_shader_context *ctx) +{ + unsigned stride = get_tcs_out_vertex_dw_stride_constant(ctx); + + return LLVMConstInt(ctx->i32, stride, 0); +} + +static LLVMValueRef get_tcs_out_patch_stride(struct si_shader_context *ctx) +{ + if (ctx->shader->key.mono.u.ff_tcs_inputs_to_copy) + return unpack_param(ctx, ctx->param_tcs_out_lds_layout, 0, 13); + + const struct tgsi_shader_info *info = &ctx->shader->selector->info; + unsigned tcs_out_vertices = info->properties[TGSI_PROPERTY_TCS_VERTICES_OUT]; + unsigned vertex_dw_stride = get_tcs_out_vertex_dw_stride_constant(ctx); + unsigned num_patch_outputs = util_last_bit64(ctx->shader->selector->patch_outputs_written); + unsigned patch_dw_stride = tcs_out_vertices * vertex_dw_stride + + num_patch_outputs * 4; + return LLVMConstInt(ctx->i32, patch_dw_stride, 0); } static LLVMValueRef @@ -349,23 +351,21 @@ get_tcs_out_patch0_patch_data_offset(struct si_shader_context *ctx) static LLVMValueRef get_tcs_in_current_patch_offset(struct si_shader_context *ctx) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx); LLVMValueRef rel_patch_id = get_rel_patch_id(ctx); - return LLVMBuildMul(gallivm->builder, patch_stride, rel_patch_id, ""); + return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, ""); } static LLVMValueRef get_tcs_out_current_patch_offset(struct si_shader_context *ctx) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef patch0_offset = get_tcs_out_patch0_offset(ctx); LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx); LLVMValueRef rel_patch_id = get_rel_patch_id(ctx); - return LLVMBuildAdd(gallivm->builder, patch0_offset, - LLVMBuildMul(gallivm->builder, patch_stride, + return LLVMBuildAdd(ctx->ac.builder, patch0_offset, + LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, ""), ""); } @@ -373,31 +373,64 @@ get_tcs_out_current_patch_offset(struct si_shader_context *ctx) static LLVMValueRef get_tcs_out_current_patch_data_offset(struct si_shader_context *ctx) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef patch0_patch_data_offset = get_tcs_out_patch0_patch_data_offset(ctx); LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx); LLVMValueRef rel_patch_id = get_rel_patch_id(ctx); - return LLVMBuildAdd(gallivm->builder, patch0_patch_data_offset, - LLVMBuildMul(gallivm->builder, patch_stride, + return LLVMBuildAdd(ctx->ac.builder, patch0_patch_data_offset, + LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, ""), ""); } +static LLVMValueRef get_num_tcs_out_vertices(struct si_shader_context *ctx) +{ + unsigned tcs_out_vertices = + ctx->shader->selector ? + ctx->shader->selector->info.properties[TGSI_PROPERTY_TCS_VERTICES_OUT] : 0; + + /* If !tcs_out_vertices, it's either the fixed-func TCS or the TCS epilog. */ + if (ctx->type == PIPE_SHADER_TESS_CTRL && tcs_out_vertices) + return LLVMConstInt(ctx->i32, tcs_out_vertices, 0); + + return unpack_param(ctx, ctx->param_tcs_offchip_layout, 6, 6); +} + +static LLVMValueRef get_tcs_in_vertex_dw_stride(struct si_shader_context *ctx) +{ + unsigned stride; + + switch (ctx->type) { + case PIPE_SHADER_VERTEX: + stride = util_last_bit64(ctx->shader->selector->outputs_written); + return LLVMConstInt(ctx->i32, stride * 4, 0); + + case PIPE_SHADER_TESS_CTRL: + 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); + } + return unpack_param(ctx, ctx->param_vs_state_bits, 24, 8); + + default: + assert(0); + return NULL; + } +} + static LLVMValueRef get_instance_index_for_fetch( struct si_shader_context *ctx, unsigned param_start_instance, LLVMValueRef divisor) { - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMValueRef result = ctx->abi.instance_id; /* The division must be done before START_INSTANCE is added. */ if (divisor != ctx->i32_1) - result = LLVMBuildUDiv(gallivm->builder, result, divisor, ""); + result = LLVMBuildUDiv(ctx->ac.builder, result, divisor, ""); - return LLVMBuildAdd(gallivm->builder, result, + return LLVMBuildAdd(ctx->ac.builder, result, LLVMGetParam(ctx->main_fn, param_start_instance), ""); } @@ -407,8 +440,8 @@ static LLVMValueRef extract_double_to_float(struct si_shader_context *ctx, LLVMValueRef vec4, unsigned double_index) { - LLVMBuilderRef builder = ctx->gallivm.builder; - LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->gallivm.context); + LLVMBuilderRef builder = ctx->ac.builder; + LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->ac.context); LLVMValueRef dvec2 = LLVMBuildBitCast(builder, vec4, LLVMVectorType(f64, 2), ""); LLVMValueRef index = LLVMConstInt(ctx->i32, double_index, 0); @@ -416,12 +449,96 @@ static LLVMValueRef extract_double_to_float(struct si_shader_context *ctx, 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]) { - struct gallivm_state *gallivm = &ctx->gallivm; + 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; @@ -439,7 +556,7 @@ void si_llvm_load_input_vs( 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 + @@ -483,7 +600,7 @@ void si_llvm_load_input_vs( /* Break up the vec4 into individual components */ for (chan = 0; chan < 4; chan++) { LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, 0); - out[chan] = LLVMBuildExtractElement(gallivm->builder, + out[chan] = LLVMBuildExtractElement(ctx->ac.builder, input[0], llvm_chan, ""); } @@ -499,9 +616,9 @@ void si_llvm_load_input_vs( /* First, recover the sign-extended signed integer value. */ if (fix_fetch == SI_FIX_FETCH_A2_SSCALED) - tmp = LLVMBuildFPToUI(gallivm->builder, tmp, ctx->i32, ""); + tmp = LLVMBuildFPToUI(ctx->ac.builder, tmp, ctx->i32, ""); else - tmp = LLVMBuildBitCast(gallivm->builder, tmp, ctx->i32, ""); + tmp = ac_to_integer(&ctx->ac, tmp); /* For the integer-like cases, do a natural sign extension. * @@ -509,20 +626,20 @@ void si_llvm_load_input_vs( * and happen to contain 0, 1, 2, 3 as the two LSBs of the * exponent. */ - tmp = LLVMBuildShl(gallivm->builder, tmp, + tmp = LLVMBuildShl(ctx->ac.builder, tmp, fix_fetch == SI_FIX_FETCH_A2_SNORM ? LLVMConstInt(ctx->i32, 7, 0) : c30, ""); - tmp = LLVMBuildAShr(gallivm->builder, tmp, c30, ""); + tmp = LLVMBuildAShr(ctx->ac.builder, tmp, c30, ""); /* Convert back to the right type. */ if (fix_fetch == SI_FIX_FETCH_A2_SNORM) { LLVMValueRef clamp; LLVMValueRef neg_one = LLVMConstReal(ctx->f32, -1.0); - tmp = LLVMBuildSIToFP(gallivm->builder, tmp, ctx->f32, ""); - clamp = LLVMBuildFCmp(gallivm->builder, LLVMRealULT, tmp, neg_one, ""); - tmp = LLVMBuildSelect(gallivm->builder, clamp, neg_one, tmp, ""); + tmp = LLVMBuildSIToFP(ctx->ac.builder, tmp, ctx->f32, ""); + clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, tmp, neg_one, ""); + tmp = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, tmp, ""); } else if (fix_fetch == SI_FIX_FETCH_A2_SSCALED) { - tmp = LLVMBuildSIToFP(gallivm->builder, tmp, ctx->f32, ""); + tmp = LLVMBuildSIToFP(ctx->ac.builder, tmp, ctx->f32, ""); } out[3] = tmp; @@ -531,11 +648,10 @@ void si_llvm_load_input_vs( case SI_FIX_FETCH_RGBA_32_UNORM: case SI_FIX_FETCH_RGBX_32_UNORM: for (chan = 0; chan < 4; chan++) { - out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan], - ctx->i32, ""); - out[chan] = LLVMBuildUIToFP(gallivm->builder, + out[chan] = ac_to_integer(&ctx->ac, out[chan]); + out[chan] = LLVMBuildUIToFP(ctx->ac.builder, out[chan], ctx->f32, ""); - out[chan] = LLVMBuildFMul(gallivm->builder, out[chan], + out[chan] = LLVMBuildFMul(ctx->ac.builder, out[chan], LLVMConstReal(ctx->f32, 1.0 / UINT_MAX), ""); } /* RGBX UINT returns 1 in alpha, which would be rounded to 0 by normalizing. */ @@ -553,11 +669,10 @@ void si_llvm_load_input_vs( scale = 1.0 / INT_MAX; for (chan = 0; chan < 4; chan++) { - out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan], - ctx->i32, ""); - out[chan] = LLVMBuildSIToFP(gallivm->builder, + out[chan] = ac_to_integer(&ctx->ac, out[chan]); + out[chan] = LLVMBuildSIToFP(ctx->ac.builder, out[chan], ctx->f32, ""); - out[chan] = LLVMBuildFMul(gallivm->builder, out[chan], + out[chan] = LLVMBuildFMul(ctx->ac.builder, out[chan], LLVMConstReal(ctx->f32, scale), ""); } /* RGBX SINT returns 1 in alpha, which would be rounded to 0 by normalizing. */ @@ -568,17 +683,15 @@ void si_llvm_load_input_vs( } case SI_FIX_FETCH_RGBA_32_USCALED: for (chan = 0; chan < 4; chan++) { - out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan], - ctx->i32, ""); - out[chan] = LLVMBuildUIToFP(gallivm->builder, + out[chan] = ac_to_integer(&ctx->ac, out[chan]); + out[chan] = LLVMBuildUIToFP(ctx->ac.builder, out[chan], ctx->f32, ""); } break; case SI_FIX_FETCH_RGBA_32_SSCALED: for (chan = 0; chan < 4; chan++) { - out[chan] = LLVMBuildBitCast(gallivm->builder, out[chan], - ctx->i32, ""); - out[chan] = LLVMBuildSIToFP(gallivm->builder, + out[chan] = ac_to_integer(&ctx->ac, out[chan]); + out[chan] = LLVMBuildSIToFP(ctx->ac.builder, out[chan], ctx->f32, ""); } break; @@ -606,7 +719,7 @@ void si_llvm_load_input_vs( case SI_FIX_FETCH_RGB_16: case SI_FIX_FETCH_RGB_16_INT: for (chan = 0; chan < 3; chan++) { - out[chan] = LLVMBuildExtractElement(gallivm->builder, + out[chan] = LLVMBuildExtractElement(ctx->ac.builder, input[chan], ctx->i32_0, ""); } @@ -614,8 +727,7 @@ void si_llvm_load_input_vs( fix_fetch == SI_FIX_FETCH_RGB_16) { out[3] = LLVMConstReal(ctx->f32, 1); } else { - out[3] = LLVMBuildBitCast(gallivm->builder, ctx->i32_1, - ctx->f32, ""); + out[3] = ac_to_float(&ctx->ac, ctx->i32_1); } break; } @@ -647,8 +759,7 @@ static LLVMValueRef get_primitive_id(struct si_shader_context *ctx, 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; @@ -661,14 +772,34 @@ static LLVMValueRef get_primitive_id(struct si_shader_context *ctx, */ LLVMValueRef si_get_indirect_index(struct si_shader_context *ctx, const struct tgsi_ind_register *ind, + unsigned addr_mul, int rel_index) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef result; - result = ctx->addrs[ind->Index][ind->Swizzle]; - result = LLVMBuildLoad(gallivm->builder, result, ""); - result = LLVMBuildAdd(gallivm->builder, result, + if (ind->File == TGSI_FILE_ADDRESS) { + result = ctx->addrs[ind->Index][ind->Swizzle]; + result = LLVMBuildLoad(ctx->ac.builder, result, ""); + } else { + struct tgsi_full_src_register src = {}; + + src.Register.File = ind->File; + src.Register.Index = ind->Index; + + /* Set the second index to 0 for constants. */ + if (ind->File == TGSI_FILE_CONSTANT) + src.Register.Dimension = 1; + + result = ctx->bld_base.emit_fetch_funcs[ind->File](&ctx->bld_base, &src, + TGSI_TYPE_SIGNED, + ind->Swizzle); + result = ac_to_integer(&ctx->ac, result); + } + + if (addr_mul != 1) + result = LLVMBuildMul(ctx->ac.builder, result, + LLVMConstInt(ctx->i32, addr_mul, 0), ""); + result = LLVMBuildAdd(ctx->ac.builder, result, LLVMConstInt(ctx->i32, rel_index, 0), ""); return result; } @@ -681,7 +812,7 @@ LLVMValueRef si_get_bounded_indirect_index(struct si_shader_context *ctx, const struct tgsi_ind_register *ind, int rel_index, unsigned num) { - LLVMValueRef result = si_get_indirect_index(ctx, ind, rel_index); + LLVMValueRef result = si_get_indirect_index(ctx, ind, 1, rel_index); return si_llvm_bound_index(ctx, result, num); } @@ -696,7 +827,6 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx, LLVMValueRef vertex_dw_stride, LLVMValueRef base_addr) { - struct gallivm_state *gallivm = &ctx->gallivm; struct tgsi_shader_info *info = &ctx->shader->selector->info; ubyte *name, *index, *array_first; int first, param; @@ -722,12 +852,12 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx, if (reg.Dimension.Indirect) index = si_get_indirect_index(ctx, ®.DimIndirect, - reg.Dimension.Index); + 1, reg.Dimension.Index); else index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0); - base_addr = LLVMBuildAdd(gallivm->builder, base_addr, - LLVMBuildMul(gallivm->builder, index, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, + LLVMBuildMul(ctx->ac.builder, index, vertex_dw_stride, ""), ""); } @@ -755,10 +885,10 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx, first = reg.Register.Index; ind_index = si_get_indirect_index(ctx, ®.Indirect, - reg.Register.Index - first); + 1, reg.Register.Index - first); - base_addr = LLVMBuildAdd(gallivm->builder, base_addr, - LLVMBuildMul(gallivm->builder, ind_index, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, + LLVMBuildMul(ctx->ac.builder, ind_index, LLVMConstInt(ctx->i32, 4, 0), ""), ""); param = reg.Register.Dimension ? @@ -773,7 +903,7 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx, } /* Add the base address of the element. */ - return LLVMBuildAdd(gallivm->builder, base_addr, + return LLVMBuildAdd(ctx->ac.builder, base_addr, LLVMConstInt(ctx->i32, param * 4, 0), ""); } @@ -800,21 +930,20 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx, LLVMValueRef vertex_index, LLVMValueRef param_index) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef base_addr, vertices_per_patch, num_patches, total_vertices; LLVMValueRef param_stride, constant16; - vertices_per_patch = unpack_param(ctx, ctx->param_tcs_offchip_layout, 6, 6); + vertices_per_patch = get_num_tcs_out_vertices(ctx); num_patches = unpack_param(ctx, ctx->param_tcs_offchip_layout, 0, 6); - total_vertices = LLVMBuildMul(gallivm->builder, vertices_per_patch, + total_vertices = LLVMBuildMul(ctx->ac.builder, vertices_per_patch, num_patches, ""); constant16 = LLVMConstInt(ctx->i32, 16, 0); if (vertex_index) { - base_addr = LLVMBuildMul(gallivm->builder, rel_patch_id, + base_addr = LLVMBuildMul(ctx->ac.builder, rel_patch_id, vertices_per_patch, ""); - base_addr = LLVMBuildAdd(gallivm->builder, base_addr, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, vertex_index, ""); param_stride = total_vertices; @@ -823,17 +952,17 @@ static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx, param_stride = num_patches; } - base_addr = LLVMBuildAdd(gallivm->builder, base_addr, - LLVMBuildMul(gallivm->builder, param_index, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, + LLVMBuildMul(ctx->ac.builder, param_index, param_stride, ""), ""); - base_addr = LLVMBuildMul(gallivm->builder, base_addr, constant16, ""); + base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, ""); if (!vertex_index) { LLVMValueRef patch_data_offset = unpack_param(ctx, ctx->param_tcs_offchip_layout, 12, 20); - base_addr = LLVMBuildAdd(gallivm->builder, base_addr, + base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr, patch_data_offset, ""); } return base_addr; @@ -844,7 +973,6 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg( const struct tgsi_full_dst_register *dst, const struct tgsi_full_src_register *src) { - struct gallivm_state *gallivm = &ctx->gallivm; struct tgsi_shader_info *info = &ctx->shader->selector->info; ubyte *name, *index, *array_first; struct tgsi_full_src_register reg; @@ -858,7 +986,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg( if (reg.Dimension.Indirect) vertex_index = si_get_indirect_index(ctx, ®.DimIndirect, - reg.Dimension.Index); + 1, reg.Dimension.Index); else vertex_index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0); } @@ -884,7 +1012,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg( param_base = reg.Register.Index; param_index = si_get_indirect_index(ctx, ®.Indirect, - reg.Register.Index - param_base); + 1, reg.Register.Index - param_base); } else { param_base = reg.Register.Index; @@ -895,7 +1023,7 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg( 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(gallivm->builder, param_index, + param_index = LLVMBuildAdd(ctx->ac.builder, param_index, LLVMConstInt(ctx->i32, param_index_base, 0), ""); @@ -909,7 +1037,6 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base, LLVMValueRef base, bool can_speculate) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef value, value2; LLVMTypeRef llvm_type = tgsi2llvmtype(bld_base, type); LLVMTypeRef vec_type = LLVMVectorType(llvm_type, 4); @@ -918,15 +1045,15 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base, value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset, 0, 1, 0, can_speculate, false); - return LLVMBuildBitCast(gallivm->builder, value, vec_type, ""); + return LLVMBuildBitCast(ctx->ac.builder, value, vec_type, ""); } if (!tgsi_type_is_64bit(type)) { value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset, 0, 1, 0, can_speculate, false); - value = LLVMBuildBitCast(gallivm->builder, value, vec_type, ""); - return LLVMBuildExtractElement(gallivm->builder, value, + value = LLVMBuildBitCast(ctx->ac.builder, value, vec_type, ""); + return LLVMBuildExtractElement(ctx->ac.builder, value, LLVMConstInt(ctx->i32, swizzle, 0), ""); } @@ -936,7 +1063,8 @@ static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base, 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); } /** @@ -951,7 +1079,6 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base, LLVMValueRef dw_addr) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef value; if (swizzle == ~0) { @@ -960,24 +1087,26 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base, 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 LLVMBuildBitCast(gallivm->builder, value, - tgsi2llvmtype(bld_base, type), ""); + return bitcast(bld_base, type, value); } /** @@ -987,25 +1116,20 @@ static LLVMValueRef lds_load(struct lp_build_tgsi_context *bld_base, * \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); - struct gallivm_state *gallivm = &ctx->gallivm; - - 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 = LLVMBuildBitCast(gallivm->builder, value, ctx->i32, ""); - 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, unsigned param) { - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMValueRef addr = LLVMGetParam(ctx->main_fn, param); addr = LLVMBuildZExt(builder, addr, ctx->i64, ""); @@ -1034,7 +1158,7 @@ static LLVMValueRef fetch_input_tcs( struct si_shader_context *ctx = si_shader_context(bld_base); LLVMValueRef dw_addr, stride; - stride = unpack_param(ctx, ctx->param_vs_state_bits, 24, 8); + stride = get_tcs_in_vertex_dw_stride(ctx); dw_addr = get_tcs_in_current_patch_offset(ctx); dw_addr = get_dw_address(ctx, NULL, reg, stride, dw_addr); @@ -1050,7 +1174,7 @@ static LLVMValueRef fetch_output_tcs( LLVMValueRef dw_addr, stride; if (reg->Register.Dimension) { - stride = unpack_param(ctx, ctx->param_tcs_out_lds_layout, 13, 8); + stride = get_tcs_out_vertex_dw_stride(ctx); dw_addr = get_tcs_out_current_patch_offset(ctx); dw_addr = get_dw_address(ctx, NULL, reg, stride, dw_addr); } else { @@ -1080,30 +1204,30 @@ static LLVMValueRef fetch_input_tes( static void store_output_tcs(struct lp_build_tgsi_context *bld_base, const struct tgsi_full_instruction *inst, const struct tgsi_opcode_info *info, + unsigned index, 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[0]; + const struct tgsi_full_dst_register *reg = &inst->Dst[index]; const struct tgsi_shader_info *sh_info = &ctx->shader->selector->info; unsigned chan_index; LLVMValueRef dw_addr, stride; LLVMValueRef buffer, base, buf_addr; LLVMValueRef values[4]; bool skip_lds_store; - bool is_tess_factor = false; + bool is_tess_factor = false, is_tess_inner = false; /* Only handle per-patch and per-vertex outputs here. * Vectors will be lowered to scalars and this function will be called again. */ if (reg->Register.File != TGSI_FILE_OUTPUT || (dst[0] && LLVMGetTypeKind(LLVMTypeOf(dst[0])) == LLVMVectorTypeKind)) { - si_llvm_emit_store(bld_base, inst, info, dst); + si_llvm_emit_store(bld_base, inst, info, index, dst); return; } if (reg->Register.Dimension) { - stride = unpack_param(ctx, ctx->param_tcs_out_lds_layout, 13, 8); + stride = get_tcs_out_vertex_dw_stride(ctx); dw_addr = get_tcs_out_current_patch_offset(ctx); dw_addr = get_dw_address(ctx, reg, NULL, stride, dw_addr); skip_lds_store = !sh_info->reads_pervertex_outputs; @@ -1118,8 +1242,11 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base, /* Always write tess factors into LDS for the TCS epilog. */ if (name == TGSI_SEMANTIC_TESSINNER || name == TGSI_SEMANTIC_TESSOUTER) { - skip_lds_store = false; + /* The epilog doesn't read LDS if invocation 0 defines tess factors. */ + skip_lds_store = !sh_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; } } } @@ -1129,8 +1256,9 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base, base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset); buf_addr = get_tcs_tes_buffer_address_from_reg(ctx, reg, NULL); - - TGSI_FOR_EACH_DST0_ENABLED_CHANNEL(inst, chan_index) { + uint32_t writemask = reg->Register.WriteMask; + while (writemask) { + chan_index = u_bit_scan(&writemask); LLVMValueRef value = dst[chan_index]; if (inst->Instruction.Saturate) @@ -1138,20 +1266,32 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base, /* 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 = LLVMBuildBitCast(gallivm->builder, value, ctx->i32, ""); + value = ac_to_integer(&ctx->ac, value); values[chan_index] = value; - if (inst->Dst[0].Register.WriteMask != 0xF && !is_tess_factor) { + if (reg->Register.WriteMask != 0xF && !is_tess_factor) { ac_build_buffer_store_dword(&ctx->ac, buffer, value, 1, buf_addr, base, 4 * chan_index, 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_index]); + } else if (chan_index < 2) { + LLVMBuildStore(ctx->ac.builder, value, /* inner */ + ctx->invoc0_tess_factors[4 + chan_index]); + } + } } - if (inst->Dst[0].Register.WriteMask == 0xF && !is_tess_factor) { - LLVMValueRef value = lp_build_gather_values(gallivm, + if (reg->Register.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, buf_addr, base, 0, 1, 0, true, false); @@ -1167,7 +1307,6 @@ static LLVMValueRef fetch_input_gs( 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]; @@ -1184,7 +1323,7 @@ static LLVMValueRef fetch_input_gs( 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) { @@ -1205,7 +1344,7 @@ static LLVMValueRef fetch_input_gs( return NULL; } - vtx_offset = LLVMBuildAdd(gallivm->builder, vtx_offset, + vtx_offset = LLVMBuildAdd(ctx->ac.builder, vtx_offset, LLVMConstInt(ctx->i32, param * 4, 0), ""); return lds_load(bld_base, type, swizzle, vtx_offset); } @@ -1217,22 +1356,15 @@ static LLVMValueRef fetch_input_gs( 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); @@ -1245,12 +1377,10 @@ static LLVMValueRef fetch_input_gs( 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 LLVMBuildBitCast(gallivm->builder, - value, - tgsi2llvmtype(bld_base, type), ""); + return bitcast(bld_base, type, value); } static int lookup_interp_param_index(unsigned interpolate, unsigned location) @@ -1325,7 +1455,6 @@ static void interp_fs_input(struct si_shader_context *ctx, LLVMValueRef face, LLVMValueRef result[4]) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef i = NULL, j = NULL; unsigned chan; @@ -1345,12 +1474,12 @@ static void interp_fs_input(struct si_shader_context *ctx, bool interp = interp_param != NULL; if (interp) { - interp_param = LLVMBuildBitCast(gallivm->builder, interp_param, + interp_param = LLVMBuildBitCast(ctx->ac.builder, interp_param, LLVMVectorType(ctx->f32, 2), ""); - i = LLVMBuildExtractElement(gallivm->builder, interp_param, + i = LLVMBuildExtractElement(ctx->ac.builder, interp_param, ctx->i32_0, ""); - j = LLVMBuildExtractElement(gallivm->builder, interp_param, + j = LLVMBuildExtractElement(ctx->ac.builder, interp_param, ctx->i32_1, ""); } @@ -1365,7 +1494,7 @@ static void interp_fs_input(struct si_shader_context *ctx, if (semantic_index == 1 && colors_read_mask & 0xf) back_attr_offset += 1; - is_face_positive = LLVMBuildICmp(gallivm->builder, LLVMIntNE, + is_face_positive = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE, face, ctx->i32_0, ""); for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) { @@ -1378,7 +1507,7 @@ static void interp_fs_input(struct si_shader_context *ctx, back_attr_offset, chan, prim_mask, i, j); - result[chan] = LLVMBuildSelect(gallivm->builder, + result[chan] = LLVMBuildSelect(ctx->ac.builder, is_face_positive, front, back, @@ -1473,15 +1602,13 @@ static LLVMValueRef buffer_load_const(struct si_shader_context *ctx, 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; - LLVMBuilderRef builder = gallivm->builder; 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); - LLVMValueRef offset1 = LLVMBuildAdd(builder, offset0, LLVMConstInt(ctx->i32, 4, 0), ""); + LLVMValueRef offset1 = LLVMBuildAdd(ctx->ac.builder, offset0, LLVMConstInt(ctx->i32, 4, 0), ""); LLVMValueRef pos[4] = { buffer_load_const(ctx, resource, offset0), @@ -1490,15 +1617,14 @@ static LLVMValueRef load_sample_position(struct si_shader_context *ctx, LLVMValu LLVMConstReal(ctx->f32, 0) }; - return lp_build_gather_values(gallivm, pos, 4); + return lp_build_gather_values(&ctx->gallivm, pos, 4); } -static void declare_system_value(struct si_shader_context *ctx, - unsigned index, - const struct tgsi_full_declaration *decl) +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; - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef value = 0; assert(index < RADEON_LLVM_MAX_SYSTEM_VALUES); @@ -1509,7 +1635,7 @@ static void declare_system_value(struct si_shader_context *ctx, break; case TGSI_SEMANTIC_VERTEXID: - value = LLVMBuildAdd(gallivm->builder, + value = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id, ctx->abi.base_vertex, ""); break; @@ -1529,10 +1655,10 @@ static void declare_system_value(struct si_shader_context *ctx, LLVMValueRef vs_state = LLVMGetParam(ctx->main_fn, ctx->param_vs_state_bits); LLVMValueRef indexed; - indexed = LLVMBuildLShr(gallivm->builder, vs_state, ctx->i32_1, ""); - indexed = LLVMBuildTrunc(gallivm->builder, indexed, ctx->i1, ""); + indexed = LLVMBuildLShr(ctx->ac.builder, vs_state, ctx->i32_1, ""); + indexed = LLVMBuildTrunc(ctx->ac.builder, indexed, ctx->i1, ""); - value = LLVMBuildSelect(gallivm->builder, indexed, + value = LLVMBuildSelect(ctx->ac.builder, indexed, ctx->abi.base_vertex, ctx->i32_0, ""); break; } @@ -1549,8 +1675,7 @@ static void declare_system_value(struct si_shader_context *ctx, 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; @@ -1565,7 +1690,7 @@ static void declare_system_value(struct si_shader_context *ctx, 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; } @@ -1588,7 +1713,7 @@ static void declare_system_value(struct si_shader_context *ctx, 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; } @@ -1604,17 +1729,17 @@ static void declare_system_value(struct si_shader_context *ctx, 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; } @@ -1622,7 +1747,7 @@ static void declare_system_value(struct si_shader_context *ctx, 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 = unpack_param(ctx, ctx->param_tcs_offchip_layout, 6, 6); + value = get_num_tcs_out_vertices(ctx); else assert(!"invalid shader stage for TGSI_SEMANTIC_VERTICESIN"); break; @@ -1653,13 +1778,13 @@ static void declare_system_value(struct si_shader_context *ctx, 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; } @@ -1687,7 +1812,7 @@ static void declare_system_value(struct si_shader_context *ctx, 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); } @@ -1705,7 +1830,7 @@ static void declare_system_value(struct si_shader_context *ctx, ctx->param_block_id[i]); } } - value = lp_build_gather_values(gallivm, values, 3); + value = lp_build_gather_values(&ctx->gallivm, values, 3); break; } @@ -1714,12 +1839,12 @@ static void declare_system_value(struct si_shader_context *ctx, break; case TGSI_SEMANTIC_HELPER_INVOCATION: - value = lp_build_intrinsic(gallivm->builder, + value = lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.ps.live", ctx->i1, NULL, 0, LP_FUNC_ATTR_READNONE); - value = LLVMBuildNot(gallivm->builder, value, ""); - value = LLVMBuildSExt(gallivm->builder, value, ctx->i32, ""); + value = LLVMBuildNot(ctx->ac.builder, value, ""); + value = LLVMBuildSExt(ctx->ac.builder, value, ctx->i32, ""); break; case TGSI_SEMANTIC_SUBGROUP_SIZE: @@ -1733,9 +1858,9 @@ static void declare_system_value(struct si_shader_context *ctx, case TGSI_SEMANTIC_SUBGROUP_EQ_MASK: { LLVMValueRef id = ac_get_thread_id(&ctx->ac); - id = LLVMBuildZExt(gallivm->builder, id, ctx->i64, ""); - value = LLVMBuildShl(gallivm->builder, LLVMConstInt(ctx->i64, 1, 0), id, ""); - value = LLVMBuildBitCast(gallivm->builder, value, ctx->v2i32, ""); + id = LLVMBuildZExt(ctx->ac.builder, id, ctx->i64, ""); + value = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->i64, 1, 0), id, ""); + value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->v2i32, ""); break; } @@ -1753,12 +1878,12 @@ static void declare_system_value(struct si_shader_context *ctx, /* All bits set */ value = LLVMConstInt(ctx->i64, -1, 0); } - id = LLVMBuildZExt(gallivm->builder, id, ctx->i64, ""); - value = LLVMBuildShl(gallivm->builder, value, id, ""); + id = LLVMBuildZExt(ctx->ac.builder, id, ctx->i64, ""); + value = LLVMBuildShl(ctx->ac.builder, value, id, ""); if (decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LE_MASK || decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LT_MASK) - value = LLVMBuildNot(gallivm->builder, value, ""); - value = LLVMBuildBitCast(gallivm->builder, value, ctx->v2i32, ""); + value = LLVMBuildNot(ctx->ac.builder, value, ""); + value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->v2i32, ""); break; } @@ -1770,26 +1895,25 @@ static void declare_system_value(struct si_shader_context *ctx, ctx->system_values[index] = value; } -static void declare_compute_memory(struct si_shader_context *ctx, - const struct tgsi_full_declaration *decl) +void si_declare_compute_memory(struct si_shader_context *ctx, + 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(gallivm->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) @@ -1797,8 +1921,8 @@ 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) @@ -1807,10 +1931,10 @@ static LLVMValueRef load_ubo(struct ac_shader_abi *abi, LLVMValueRef index) LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers); index = si_llvm_bound_index(ctx, index, ctx->num_const_buffers); - index = LLVMBuildAdd(ctx->gallivm.builder, 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 @@ -1821,11 +1945,11 @@ load_ssbo(struct ac_shader_abi *abi, LLVMValueRef index, bool write) ctx->param_const_and_shader_buffers); index = si_llvm_bound_index(ctx, index, ctx->num_shader_buffers); - index = LLVMBuildSub(ctx->gallivm.builder, + index = LLVMBuildSub(ctx->ac.builder, 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( @@ -1835,12 +1959,11 @@ static LLVMValueRef fetch_constant( unsigned swizzle) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct lp_build_context *base = &bld_base->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; @@ -1851,54 +1974,99 @@ static LLVMValueRef fetch_constant( return lp_build_gather_values(&ctx->gallivm, values, 4); } - buf = reg->Register.Dimension ? reg->Dimension.Index : 0; + /* 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); + } - if (reg->Register.Dimension && reg->Dimension.Indirect) { + /* 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; + + if (reg->Dimension.Indirect) { LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers); LLVMValueRef index; index = si_get_bounded_indirect_index(ctx, ®->DimIndirect, reg->Dimension.Index, ctx->num_const_buffers); - index = LLVMBuildAdd(ctx->gallivm.builder, index, + 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 = ctx->addrs[ireg->Index][ireg->Swizzle]; - addr = LLVMBuildLoad(base->gallivm->builder, addr, "load addr reg"); - addr = lp_build_mul_imm(&bld_base->uint_bld, addr, 16); - addr = lp_build_add(&bld_base->uint_bld, addr, - LLVMConstInt(ctx->i32, idx * 4, 0)); - } 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. */ static LLVMValueRef si_llvm_pack_two_int16(struct si_shader_context *ctx, LLVMValueRef val[2]) { - return LLVMBuildOr(ctx->gallivm.builder, val[0], - LLVMBuildShl(ctx->gallivm.builder, val[1], + return LLVMBuildOr(ctx->ac.builder, val[0], + LLVMBuildShl(ctx->ac.builder, val[1], LLVMConstInt(ctx->i32, 16, 0), ""), ""); } @@ -1908,7 +2076,7 @@ static LLVMValueRef si_llvm_pack_two_int32_as_int16(struct si_shader_context *ct LLVMValueRef val[2]) { LLVMValueRef v[2] = { - LLVMBuildAnd(ctx->gallivm.builder, val[0], + LLVMBuildAnd(ctx->ac.builder, val[0], LLVMConstInt(ctx->i32, 0xffff, 0), ""), val[1], }; @@ -1916,14 +2084,13 @@ static LLVMValueRef si_llvm_pack_two_int32_as_int16(struct si_shader_context *ct } /* 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; - LLVMBuilderRef builder = ctx->gallivm.builder; + 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; unsigned chan; @@ -1953,10 +2120,10 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, } 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: @@ -1992,9 +2159,7 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, LLVMValueRef packed; packed = ac_build_cvt_pkrtz_f16(&ctx->ac, pack_args); - args->out[chan] = - LLVMBuildBitCast(ctx->gallivm.builder, - packed, ctx->f32, ""); + args->out[chan] = ac_to_float(&ctx->ac, packed); } break; @@ -2010,19 +2175,17 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, } args->compr = 1; /* COMPR flag */ - args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int16(ctx, val)); - args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int16(ctx, val+2)); + args->out[0] = ac_to_float(&ctx->ac, si_llvm_pack_two_int16(ctx, val)); + args->out[1] = ac_to_float(&ctx->ac, si_llvm_pack_two_int16(ctx, val+2)); break; 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]. */ @@ -2032,17 +2195,15 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, 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, ""); } args->compr = 1; /* COMPR flag */ - args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int32_as_int16(ctx, val)); - args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int32_as_int16(ctx, val+2)); + args->out[0] = ac_to_float(&ctx->ac, si_llvm_pack_two_int32_as_int16(ctx, val)); + args->out[1] = ac_to_float(&ctx->ac, si_llvm_pack_two_int32_as_int16(ctx, val+2)); break; case V_028714_SPI_SHADER_UINT16_ABGR: { @@ -2053,17 +2214,15 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, /* Clamp. */ for (chan = 0; chan < 4; chan++) { - val[chan] = bitcast(bld_base, TGSI_TYPE_UNSIGNED, values[chan]); - val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_UMIN, + val[chan] = ac_to_integer(&ctx->ac, values[chan]); + val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_UMIN, val[chan], chan == 3 ? max_alpha : max_rgb); } args->compr = 1; /* COMPR flag */ - args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int16(ctx, val)); - args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int16(ctx, val+2)); + args->out[0] = ac_to_float(&ctx->ac, si_llvm_pack_two_int16(ctx, val)); + args->out[1] = ac_to_float(&ctx->ac, si_llvm_pack_two_int16(ctx, val+2)); break; } @@ -2079,20 +2238,18 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, /* Clamp. */ for (chan = 0; chan < 4; chan++) { - val[chan] = bitcast(bld_base, TGSI_TYPE_UNSIGNED, values[chan]); - val[chan] = lp_build_emit_llvm_binary(bld_base, + val[chan] = ac_to_integer(&ctx->ac, values[chan]); + 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); } args->compr = 1; /* COMPR flag */ - args->out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int32_as_int16(ctx, val)); - args->out[1] = bitcast(bld_base, TGSI_TYPE_FLOAT, - si_llvm_pack_two_int32_as_int16(ctx, val+2)); + args->out[0] = ac_to_float(&ctx->ac, si_llvm_pack_two_int32_as_int16(ctx, val)); + args->out[1] = ac_to_float(&ctx->ac, si_llvm_pack_two_int32_as_int16(ctx, val+2)); break; } @@ -2108,22 +2265,24 @@ static void si_alpha_test(struct lp_build_tgsi_context *bld_base, 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)); } } @@ -2132,33 +2291,30 @@ static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context * unsigned samplemask_param) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef coverage; /* alpha = alpha * popcount(coverage) / SI_NUM_SMOOTH_AA_SAMPLES */ coverage = LLVMGetParam(ctx->main_fn, samplemask_param); - coverage = bitcast(bld_base, TGSI_TYPE_SIGNED, coverage); + coverage = ac_to_integer(&ctx->ac, coverage); - coverage = lp_build_intrinsic(gallivm->builder, "llvm.ctpop.i32", + coverage = lp_build_intrinsic(ctx->ac.builder, "llvm.ctpop.i32", ctx->i32, &coverage, 1, LP_FUNC_ATTR_READNONE); - coverage = LLVMBuildUIToFP(gallivm->builder, coverage, + coverage = LLVMBuildUIToFP(ctx->ac.builder, coverage, ctx->f32, ""); - coverage = LLVMBuildFMul(gallivm->builder, coverage, + coverage = LLVMBuildFMul(ctx->ac.builder, coverage, LLVMConstReal(ctx->f32, 1.0 / SI_NUM_SMOOTH_AA_SAMPLES), ""); - return LLVMBuildFMul(gallivm->builder, alpha, coverage, ""); + 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; @@ -2166,7 +2322,7 @@ static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base, 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]; @@ -2185,8 +2341,8 @@ static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base, 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])); } } @@ -2226,8 +2382,6 @@ static void emit_streamout_output(struct si_shader_context *ctx, struct pipe_stream_output *stream_out, struct si_shader_output_values *shader_out) { - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMBuilderRef builder = gallivm->builder; unsigned buf_idx = stream_out->output_buffer; unsigned start = stream_out->start_component; unsigned num_comps = stream_out->num_components; @@ -2241,9 +2395,7 @@ static void emit_streamout_output(struct si_shader_context *ctx, for (int j = 0; j < num_comps; j++) { assert(stream_out->stream == shader_out->vertex_stream[start + j]); - out[j] = LLVMBuildBitCast(builder, - shader_out->values[start + j], - ctx->i32, ""); + out[j] = ac_to_integer(&ctx->ac, shader_out->values[start + j]); } /* Pack the output. */ @@ -2258,7 +2410,7 @@ static void emit_streamout_output(struct si_shader_context *ctx, case 4: /* as v4i32 */ vdata = LLVMGetUndef(LLVMVectorType(ctx->i32, util_next_power_of_two(num_comps))); for (int j = 0; j < num_comps; j++) { - vdata = LLVMBuildInsertElement(builder, vdata, out[j], + vdata = LLVMBuildInsertElement(ctx->ac.builder, vdata, out[j], LLVMConstInt(ctx->i32, j, 0), ""); } break; @@ -2281,8 +2433,7 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx, { struct si_shader_selector *sel = ctx->shader->selector; struct pipe_stream_output_info *so = &sel->so; - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMBuilderRef builder = gallivm->builder; + LLVMBuilderRef builder = ctx->ac.builder; int i; struct lp_build_if_state if_ctx; @@ -2299,7 +2450,7 @@ static void si_llvm_emit_streamout(struct si_shader_context *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 + @@ -2328,7 +2479,7 @@ static void si_llvm_emit_streamout(struct si_shader_context *ctx, 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]); @@ -2361,7 +2512,7 @@ static void si_export_param(struct si_shader_context *ctx, unsigned index, { 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); } @@ -2414,13 +2565,11 @@ static void si_build_param_exports(struct si_shader_context *ctx, } /* 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; @@ -2430,7 +2579,7 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, 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: @@ -2448,14 +2597,14 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, 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; @@ -2469,10 +2618,10 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, 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). */ @@ -2488,10 +2637,10 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, 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; @@ -2499,7 +2648,7 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, if (shader->selector->info.writes_edgeflag) { /* The output is a float, but the hw expects an integer * with the first bit containing the edge flag. */ - edgeflag_value = LLVMBuildFPToUI(ctx->gallivm.builder, + edgeflag_value = LLVMBuildFPToUI(ctx->ac.builder, edgeflag_value, ctx->i32, ""); edgeflag_value = ac_build_umin(&ctx->ac, @@ -2507,12 +2656,10 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, ctx->i32_1); /* The LLVM intrinsic expects a float. */ - pos_args[1].out[1] = LLVMBuildBitCast(ctx->gallivm.builder, - edgeflag_value, - ctx->f32, ""); + 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]. */ @@ -2522,13 +2669,12 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, if (shader->selector->info.writes_viewport_index) { LLVMValueRef v = viewport_index_value; - v = bitcast(bld_base, TGSI_TYPE_UNSIGNED, v); - v = LLVMBuildShl(ctx->gallivm.builder, v, + v = ac_to_integer(&ctx->ac, v); + v = LLVMBuildShl(ctx->ac.builder, v, LLVMConstInt(ctx->i32, 16, 0), ""); - v = LLVMBuildOr(ctx->gallivm.builder, v, - bitcast(bld_base, TGSI_TYPE_UNSIGNED, - pos_args[1].out[2]), ""); - pos_args[1].out[2] = bitcast(bld_base, TGSI_TYPE_FLOAT, v); + v = LLVMBuildOr(ctx->ac.builder, v, + ac_to_integer(&ctx->ac, pos_args[1].out[2]), ""); + pos_args[1].out[2] = ac_to_float(&ctx->ac, v); pos_args[1].enabled_channels |= 1 << 2; } } else { @@ -2572,7 +2718,6 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef invocation_id, buffer, buffer_offset; LLVMValueRef lds_vertex_stride, lds_vertex_offset, lds_base; uint64_t inputs; @@ -2581,17 +2726,17 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base) buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k); buffer_offset = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset); - lds_vertex_stride = unpack_param(ctx, ctx->param_vs_state_bits, 24, 8); - lds_vertex_offset = LLVMBuildMul(gallivm->builder, invocation_id, + lds_vertex_stride = get_tcs_in_vertex_dw_stride(ctx); + lds_vertex_offset = LLVMBuildMul(ctx->ac.builder, invocation_id, lds_vertex_stride, ""); lds_base = get_tcs_in_current_patch_offset(ctx); - lds_base = LLVMBuildAdd(gallivm->builder, lds_base, lds_vertex_offset, ""); + lds_base = LLVMBuildAdd(ctx->ac.builder, lds_base, lds_vertex_offset, ""); inputs = ctx->shader->key.mono.u.ff_tcs_inputs_to_copy; while (inputs) { unsigned i = u_bit_scan64(&inputs); - LLVMValueRef lds_ptr = LLVMBuildAdd(gallivm->builder, lds_base, + LLVMValueRef lds_ptr = LLVMBuildAdd(ctx->ac.builder, lds_base, LLVMConstInt(ctx->i32, 4 * i, 0), ""); @@ -2611,10 +2756,11 @@ static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base) static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, LLVMValueRef rel_patch_id, LLVMValueRef invocation_id, - LLVMValueRef tcs_out_current_patch_data_offset) + LLVMValueRef tcs_out_current_patch_data_offset, + LLVMValueRef invoc0_tf_outer[4], + 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; @@ -2622,7 +2768,9 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, unsigned stride, outer_comps, inner_comps, i, offset; struct lp_build_if_state if_ctx, inner_if_ctx; - si_llvm_emit_barrier(NULL, bld_base, NULL); + /* Add a barrier before loading tess factors from LDS. */ + if (!shader->key.part.tcs.epilog.invoc0_tess_factors_are_def) + si_llvm_emit_barrier(NULL, bld_base, NULL); /* Do this only for invocation 0, because the tess levels are per-patch, * not per-vertex. @@ -2630,8 +2778,8 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, * This can't jump, because invocation 0 executes this. It should * at least mask out the loads and stores for other invocations. */ - lp_build_if(&if_ctx, gallivm, - LLVMBuildICmp(gallivm->builder, LLVMIntEQ, + lp_build_if(&if_ctx, &ctx->gallivm, + LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, invocation_id, ctx->i32_0, "")); /* Determine the layout of one tess factor element in the buffer. */ @@ -2656,32 +2804,32 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, return; } - /* Load tess_inner and tess_outer from LDS. - * Any invocation can write them, so we can't get them from a temporary. - */ - tess_inner_index = si_shader_io_get_unique_index_patch(TGSI_SEMANTIC_TESSINNER, 0); - tess_outer_index = si_shader_io_get_unique_index_patch(TGSI_SEMANTIC_TESSOUTER, 0); - - lds_base = tcs_out_current_patch_data_offset; - lds_inner = LLVMBuildAdd(gallivm->builder, lds_base, - LLVMConstInt(ctx->i32, - tess_inner_index * 4, 0), ""); - lds_outer = LLVMBuildAdd(gallivm->builder, lds_base, - LLVMConstInt(ctx->i32, - tess_outer_index * 4, 0), ""); - for (i = 0; i < 4; i++) { inner[i] = LLVMGetUndef(ctx->i32); outer[i] = LLVMGetUndef(ctx->i32); } - if (shader->key.part.tcs.epilog.prim_mode == PIPE_PRIM_LINES) { - /* For isolines, the hardware expects tess factors in the - * reverse order from what GLSL / TGSI specify. - */ - outer[0] = out[1] = lds_load(bld_base, TGSI_TYPE_SIGNED, 0, lds_outer); - outer[1] = out[0] = lds_load(bld_base, TGSI_TYPE_SIGNED, 1, lds_outer); + if (shader->key.part.tcs.epilog.invoc0_tess_factors_are_def) { + /* Tess factors are in VGPRs. */ + for (i = 0; i < outer_comps; i++) + outer[i] = out[i] = invoc0_tf_outer[i]; + for (i = 0; i < inner_comps; i++) + inner[i] = out[outer_comps+i] = invoc0_tf_inner[i]; } else { + /* Load tess_inner and tess_outer from LDS. + * Any invocation can write them, so we can't get them from a temporary. + */ + tess_inner_index = si_shader_io_get_unique_index_patch(TGSI_SEMANTIC_TESSINNER, 0); + tess_outer_index = si_shader_io_get_unique_index_patch(TGSI_SEMANTIC_TESSOUTER, 0); + + lds_base = tcs_out_current_patch_data_offset; + lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base, + LLVMConstInt(ctx->i32, + tess_inner_index * 4, 0), ""); + lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base, + LLVMConstInt(ctx->i32, + tess_outer_index * 4, 0), ""); + for (i = 0; i < outer_comps; i++) { outer[i] = out[i] = lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_outer); @@ -2692,12 +2840,21 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, } } + if (shader->key.part.tcs.epilog.prim_mode == PIPE_PRIM_LINES) { + /* For isolines, the hardware expects tess factors in the + * reverse order from what GLSL / TGSI specify. + */ + LLVMValueRef tmp = out[0]; + out[0] = out[1]; + out[1] = tmp; + } + /* 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); @@ -2705,16 +2862,16 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, /* Get the offset. */ tf_base = LLVMGetParam(ctx->main_fn, ctx->param_tcs_factor_offset); - byteoffset = LLVMBuildMul(gallivm->builder, rel_patch_id, + byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id, LLVMConstInt(ctx->i32, 4 * stride, 0), ""); - lp_build_if(&inner_if_ctx, gallivm, - LLVMBuildICmp(gallivm->builder, LLVMIntEQ, + 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, @@ -2748,7 +2905,7 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_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, @@ -2761,7 +2918,7 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, 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); @@ -2775,7 +2932,7 @@ static LLVMValueRef si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret, unsigned param, unsigned return_index) { - return LLVMBuildInsertValue(ctx->gallivm.builder, ret, + return LLVMBuildInsertValue(ctx->ac.builder, ret, LLVMGetParam(ctx->main_fn, param), return_index, ""); } @@ -2784,11 +2941,11 @@ static LLVMValueRef si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret, unsigned param, unsigned return_index) { - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMValueRef p = LLVMGetParam(ctx->main_fn, param); return LLVMBuildInsertValue(builder, ret, - LLVMBuildBitCast(builder, p, ctx->f32, ""), + ac_to_float(&ctx->ac, p), return_index, ""); } @@ -2796,7 +2953,7 @@ static LLVMValueRef si_insert_input_ptr_as_2xi32(struct si_shader_context *ctx, LLVMValueRef ret, unsigned param, unsigned return_index) { - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMValueRef ptr, lo, hi; ptr = LLVMGetParam(ctx->main_fn, param); @@ -2812,7 +2969,7 @@ si_insert_input_ptr_as_2xi32(struct si_shader_context *ctx, LLVMValueRef ret, static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base) { struct si_shader_context *ctx = si_shader_context(bld_base); - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMValueRef rel_patch_id, invocation_id, tf_lds_offset; si_copy_tcs_inputs(bld_base); @@ -2821,7 +2978,7 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base) 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 @@ -2832,22 +2989,22 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base) values[0] = rel_patch_id; values[1] = LLVMGetUndef(ctx->i32); - rel_patch_id = build_phi(&ctx->ac, ctx->i32, 2, values, blocks); + rel_patch_id = ac_build_phi(&ctx->ac, ctx->i32, 2, values, blocks); values[0] = tf_lds_offset; values[1] = LLVMGetUndef(ctx->i32); - tf_lds_offset = build_phi(&ctx->ac, ctx->i32, 2, values, blocks); + tf_lds_offset = ac_build_phi(&ctx->ac, ctx->i32, 2, values, blocks); values[0] = invocation_id; values[1] = ctx->i32_1; /* cause the epilog to skip threads */ - invocation_id = build_phi(&ctx->ac, ctx->i32, 2, values, blocks); + invocation_id = ac_build_phi(&ctx->ac, ctx->i32, 2, values, blocks); } /* Return epilog parameters from this function. */ 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, @@ -2874,9 +3031,9 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base) } /* VGPRs */ - rel_patch_id = bitcast(bld_base, TGSI_TYPE_FLOAT, rel_patch_id); - invocation_id = bitcast(bld_base, TGSI_TYPE_FLOAT, invocation_id); - tf_lds_offset = bitcast(bld_base, TGSI_TYPE_FLOAT, tf_lds_offset); + rel_patch_id = ac_to_float(&ctx->ac, rel_patch_id); + invocation_id = ac_to_float(&ctx->ac, invocation_id); + 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, @@ -2886,7 +3043,18 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base) ret = LLVMBuildInsertValue(builder, ret, rel_patch_id, vgpr++, ""); ret = LLVMBuildInsertValue(builder, ret, invocation_id, vgpr++, ""); - ret = LLVMBuildInsertValue(builder, ret, tf_lds_offset, vgpr++, ""); + + if (ctx->shader->selector->tcs_info.tessfactors_are_def_in_all_invocs) { + vgpr++; /* skip the tess factor LDS offset */ + for (unsigned i = 0; i < 6; i++) { + LLVMValueRef value = + LLVMBuildLoad(builder, ctx->invoc0_tess_factors[i], ""); + value = ac_to_float(&ctx->ac, value); + ret = LLVMBuildInsertValue(builder, ret, value, vgpr++, ""); + } + } else { + ret = LLVMBuildInsertValue(builder, ret, tf_lds_offset, vgpr++, ""); + } ctx->return_value = ret; } @@ -2895,12 +3063,17 @@ static void si_set_ls_return_value_for_tcs(struct si_shader_context *ctx) { 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); + ret = si_insert_input_ret(ctx, ret, ctx->param_vs_state_bits, 8 + SI_SGPR_VS_STATE_BITS); ret = si_insert_input_ret(ctx, ret, ctx->param_tcs_offchip_layout, @@ -2933,12 +3106,16 @@ static void si_set_es_return_value_for_gs(struct si_shader_context *ctx) { 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); + unsigned desc_param = ctx->param_vs_state_bits + 1; ret = si_insert_input_ptr_as_2xi32(ctx, ret, desc_param, 8 + GFX9_SGPR_GS_CONST_AND_SHADER_BUFFERS); @@ -2953,24 +3130,23 @@ static void si_set_es_return_value_for_gs(struct si_shader_context *ctx) 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; - struct gallivm_state *gallivm = &ctx->gallivm; unsigned i, chan; LLVMValueRef vertex_id = LLVMGetParam(ctx->main_fn, ctx->param_rel_auto_id); - LLVMValueRef vertex_dw_stride = - unpack_param(ctx, ctx->param_vs_state_bits, 24, 8); - LLVMValueRef base_dw_addr = LLVMBuildMul(gallivm->builder, vertex_id, + LLVMValueRef vertex_dw_stride = get_tcs_in_vertex_dw_stride(ctx); + LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id, vertex_dw_stride, ""); /* 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]; @@ -2994,23 +3170,27 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base) continue; int param = si_shader_io_get_unique_index(name, index); - LLVMValueRef dw_addr = LLVMBuildAdd(gallivm->builder, base_dw_addr, + LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr, LLVMConstInt(ctx->i32, param * 4, 0), ""); for (chan = 0; chan < 4; chan++) { - lds_store(bld_base, chan, dw_addr, - LLVMBuildLoad(gallivm->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 gallivm_state *gallivm = &ctx->gallivm; + 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, @@ -3019,19 +3199,18 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base) 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); - vertex_idx = LLVMBuildOr(gallivm->builder, vertex_idx, - LLVMBuildMul(gallivm->builder, wave_idx, + vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx, + LLVMBuildMul(ctx->ac.builder, wave_idx, LLVMConstInt(ctx->i32, 64, false), ""), ""); - lds_base = LLVMBuildMul(gallivm->builder, vertex_idx, + lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx, LLVMConstInt(ctx->i32, itemsize_dw, 0), ""); } 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 || @@ -3042,12 +3221,12 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base) info->output_semantic_index[i]); for (chan = 0; chan < 4; chan++) { - LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], ""); - out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, ""); + 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; } @@ -3059,35 +3238,50 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base) } } - 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; @@ -3118,16 +3312,16 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi, /* The state is in the first bit of the user SGPR. */ cond = LLVMGetParam(ctx->main_fn, ctx->param_vs_state_bits); - cond = LLVMBuildTrunc(gallivm->builder, cond, + 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++) { addr = addrs[4 * i + j]; - val = LLVMBuildLoad(gallivm->builder, addr, ""); + val = LLVMBuildLoad(ctx->ac.builder, addr, ""); val = ac_build_clamp(&ctx->ac, val); - LLVMBuildStore(gallivm->builder, val, addr); + LLVMBuildStore(ctx->ac.builder, val, addr); } } @@ -3141,7 +3335,7 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi, for (j = 0; j < 4; j++) { outputs[i].values[j] = - LLVMBuildLoad(gallivm->builder, + LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + j], ""); outputs[i].vertex_stream[j] = @@ -3156,8 +3350,7 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi, if (ctx->shader->key.mono.u.vs_export_prim_id) { outputs[i].semantic_name = TGSI_SEMANTIC_PRIMID; outputs[i].semantic_index = 0; - outputs[i].values[0] = LLVMBuildBitCast(gallivm->builder, - get_primitive_id(ctx, 0), ctx->f32, ""); + outputs[i].values[0] = ac_to_float(&ctx->ac, get_primitive_id(ctx, 0)); for (j = 1; j < 4; j++) outputs[i].values[j] = LLVMConstReal(ctx->f32, 0); @@ -3166,7 +3359,7 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi, i++; } - si_llvm_export_vs(&ctx->bld_base, outputs, i); + si_llvm_export_vs(ctx, outputs, i); FREE(outputs); } @@ -3234,10 +3427,10 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base, if (stencil) { /* Stencil should be in X[23:16]. */ - stencil = bitcast(bld_base, TGSI_TYPE_UNSIGNED, stencil); - stencil = LLVMBuildShl(ctx->gallivm.builder, stencil, + stencil = ac_to_integer(&ctx->ac, stencil); + stencil = LLVMBuildShl(ctx->ac.builder, stencil, LLVMConstInt(ctx->i32, 16, 0), ""); - args.out[0] = bitcast(bld_base, TGSI_TYPE_FLOAT, stencil); + args.out[0] = ac_to_float(&ctx->ac, stencil); mask |= 0x3; } if (samplemask) { @@ -3262,9 +3455,9 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base, /* 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 */ @@ -3279,7 +3472,6 @@ static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base, 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 */ @@ -3289,7 +3481,7 @@ static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base, /* 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 && @@ -3308,7 +3500,7 @@ static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base, /* 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; @@ -3328,7 +3520,7 @@ static void si_export_mrt_color(struct lp_build_tgsi_context *bld_base, 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 */ @@ -3386,7 +3578,7 @@ static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi, 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; - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; unsigned i, j, first_vgpr, vgpr; LLVMValueRef color[8][4] = {}; @@ -3394,7 +3586,7 @@ static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi, 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++) { @@ -3433,10 +3625,9 @@ static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi, /* Set SGPRs. */ ret = LLVMBuildInsertValue(builder, ret, - LLVMBuildBitCast(ctx->ac.builder, - LLVMGetParam(ctx->main_fn, - SI_PARAM_ALPHA_REF), - ctx->i32, ""), + ac_to_integer(&ctx->ac, + LLVMGetParam(ctx->main_fn, + SI_PARAM_ALPHA_REF)), SI_SGPR_ALPHA_REF, ""); /* Set VGPRs */ @@ -3465,55 +3656,12 @@ static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi, ctx->return_value = ret; } -/* Prevent optimizations (at least of memory accesses) across the current - * point in the program by emitting empty inline assembly that is marked as - * having side effects. - * - * Optionally, a value can be passed through the inline assembly to prevent - * LLVM from hoisting calls to ReadNone functions. - */ -static void emit_optimization_barrier(struct si_shader_context *ctx, - LLVMValueRef *pvgpr) -{ - static int counter = 0; - - LLVMBuilderRef builder = ctx->gallivm.builder; - char code[16]; - - snprintf(code, sizeof(code), "; %d", p_atomic_inc_return(&counter)); - - if (!pvgpr) { - LLVMTypeRef ftype = LLVMFunctionType(ctx->voidt, NULL, 0, false); - LLVMValueRef inlineasm = LLVMConstInlineAsm(ftype, code, "", true, false); - LLVMBuildCall(builder, inlineasm, NULL, 0, ""); - } else { - LLVMTypeRef ftype = LLVMFunctionType(ctx->i32, &ctx->i32, 1, false); - LLVMValueRef inlineasm = LLVMConstInlineAsm(ftype, code, "=v,0", true, false); - LLVMValueRef vgpr = *pvgpr; - LLVMTypeRef vgpr_type = LLVMTypeOf(vgpr); - unsigned vgpr_size = llvm_get_type_size(vgpr_type); - LLVMValueRef vgpr0; - - assert(vgpr_size % 4 == 0); - - vgpr = LLVMBuildBitCast(builder, vgpr, LLVMVectorType(ctx->i32, vgpr_size / 4), ""); - vgpr0 = LLVMBuildExtractElement(builder, vgpr, ctx->i32_0, ""); - vgpr0 = LLVMBuildCall(builder, inlineasm, &vgpr0, 1, ""); - vgpr = LLVMBuildInsertElement(builder, vgpr, vgpr0, ctx->i32_0, ""); - vgpr = LLVMBuildBitCast(builder, vgpr, vgpr_type, ""); - - *pvgpr = vgpr; - } -} - void si_emit_waitcnt(struct si_shader_context *ctx, unsigned simm16) { - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMBuilderRef builder = gallivm->builder; LLVMValueRef args[1] = { LLVMConstInt(ctx->i32, simm16, 0) }; - lp_build_intrinsic(builder, "llvm.amdgcn.s.waitcnt", + lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.s.waitcnt", ctx->voidt, args, 1, 0); } @@ -3548,17 +3696,16 @@ static void clock_emit( struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef tmp; - tmp = lp_build_intrinsic(gallivm->builder, "llvm.readcyclecounter", + tmp = lp_build_intrinsic(ctx->ac.builder, "llvm.readcyclecounter", ctx->i64, NULL, 0, 0); - tmp = LLVMBuildBitCast(gallivm->builder, tmp, ctx->v2i32, ""); + tmp = LLVMBuildBitCast(ctx->ac.builder, tmp, ctx->v2i32, ""); emit_data->output[0] = - LLVMBuildExtractElement(gallivm->builder, tmp, ctx->i32_0, ""); + LLVMBuildExtractElement(ctx->ac.builder, tmp, ctx->i32_0, ""); emit_data->output[1] = - LLVMBuildExtractElement(gallivm->builder, tmp, ctx->i32_1, ""); + LLVMBuildExtractElement(ctx->ac.builder, tmp, ctx->i32_1, ""); } LLVMTypeRef si_const_array(LLVMTypeRef elem_type, int num_elements) @@ -3573,7 +3720,6 @@ static void si_llvm_emit_ddxy( struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; unsigned opcode = emit_data->info->opcode; LLVMValueRef val; int idx; @@ -3589,9 +3735,8 @@ static void si_llvm_emit_ddxy( /* for DDX we want to next X pixel, DDY next Y pixel. */ idx = (opcode == TGSI_OPCODE_DDX || opcode == TGSI_OPCODE_DDX_FINE) ? 1 : 2; - val = LLVMBuildBitCast(gallivm->builder, emit_data->args[0], ctx->i32, ""); - val = ac_build_ddxy(&ctx->ac, ctx->screen->has_ds_bpermute, - mask, idx, val); + val = ac_to_integer(&ctx->ac, emit_data->args[0]); + val = ac_build_ddxy(&ctx->ac, mask, idx, val); emit_data->output[emit_data->chan] = val; } @@ -3605,18 +3750,17 @@ static LLVMValueRef si_llvm_emit_ddxy_interp( 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; for (i = 0; i < 2; i++) { - a = LLVMBuildExtractElement(gallivm->builder, interp_ij, + a = LLVMBuildExtractElement(ctx->ac.builder, interp_ij, LLVMConstInt(ctx->i32, i, 0), ""); result[i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDX, a); result[2+i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDY, a); } - return lp_build_gather_values(gallivm, result, 4); + return lp_build_gather_values(&ctx->gallivm, result, 4); } static void interp_fetch_args( @@ -3624,7 +3768,6 @@ 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) { @@ -3646,19 +3789,44 @@ static void interp_fetch_args( */ sample_id = lp_build_emit_fetch(bld_base, emit_data->inst, 1, TGSI_CHAN_X); - sample_id = LLVMBuildBitCast(gallivm->builder, sample_id, - ctx->i32, ""); - sample_position = load_sample_position(ctx, sample_id); + sample_id = ac_to_integer(&ctx->ac, sample_id); - emit_data->args[0] = LLVMBuildExtractElement(gallivm->builder, + /* Section 8.13.2 (Interpolation Functions) of the OpenGL Shading + * Language 4.50 spec says about interpolateAtSample: + * + * "Returns the value of the input interpolant variable at + * the location of sample number sample. If multisample + * buffers are not available, the input variable will be + * evaluated at the center of the pixel. If sample sample + * does not exist, the position used to interpolate the + * input variable is undefined." + * + * This means that sample_id values outside of the valid are + * in fact valid input, and the usual mechanism for loading the + * sample position doesn't work. + */ + if (ctx->shader->key.mono.u.ps.interpolate_at_sample_force_center) { + LLVMValueRef center[4] = { + LLVMConstReal(ctx->f32, 0.5), + LLVMConstReal(ctx->f32, 0.5), + ctx->ac.f32_0, + ctx->ac.f32_0, + }; + + sample_position = lp_build_gather_values(&ctx->gallivm, center, 4); + } else { + sample_position = load_sample_position(ctx, sample_id); + } + + emit_data->args[0] = LLVMBuildExtractElement(ctx->ac.builder, sample_position, ctx->i32_0, ""); - emit_data->args[0] = LLVMBuildFSub(gallivm->builder, emit_data->args[0], halfval, ""); - emit_data->args[1] = LLVMBuildExtractElement(gallivm->builder, + emit_data->args[0] = LLVMBuildFSub(ctx->ac.builder, emit_data->args[0], halfval, ""); + emit_data->args[1] = LLVMBuildExtractElement(ctx->ac.builder, sample_position, ctx->i32_1, ""); - emit_data->args[1] = LLVMBuildFSub(gallivm->builder, emit_data->args[1], halfval, ""); + emit_data->args[1] = LLVMBuildFSub(ctx->ac.builder, emit_data->args[1], halfval, ""); emit_data->arg_count = 2; } } @@ -3669,7 +3837,6 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action, { 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; @@ -3697,7 +3864,7 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action, } array_idx = si_get_indirect_index(ctx, &input->Indirect, - input->Register.Index - input_base); + 1, input->Register.Index - input_base); } else { input_base = inst->Src[0].Register.Index; input_array_size = 1; @@ -3736,32 +3903,29 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action, for (i = 0; i < 2; i++) { LLVMValueRef ix_ll = LLVMConstInt(ctx->i32, i, 0); LLVMValueRef iy_ll = LLVMConstInt(ctx->i32, i + 2, 0); - LLVMValueRef ddx_el = LLVMBuildExtractElement(gallivm->builder, + LLVMValueRef ddx_el = LLVMBuildExtractElement(ctx->ac.builder, ddxy_out, ix_ll, ""); - LLVMValueRef ddy_el = LLVMBuildExtractElement(gallivm->builder, + LLVMValueRef ddy_el = LLVMBuildExtractElement(ctx->ac.builder, ddxy_out, iy_ll, ""); - LLVMValueRef interp_el = LLVMBuildExtractElement(gallivm->builder, + LLVMValueRef interp_el = LLVMBuildExtractElement(ctx->ac.builder, interp_param, ix_ll, ""); LLVMValueRef temp1, temp2; - interp_el = LLVMBuildBitCast(gallivm->builder, interp_el, - ctx->f32, ""); + interp_el = ac_to_float(&ctx->ac, interp_el); - temp1 = LLVMBuildFMul(gallivm->builder, ddx_el, emit_data->args[0], ""); + temp1 = LLVMBuildFMul(ctx->ac.builder, ddx_el, emit_data->args[0], ""); - temp1 = LLVMBuildFAdd(gallivm->builder, temp1, interp_el, ""); + temp1 = LLVMBuildFAdd(ctx->ac.builder, temp1, interp_el, ""); - temp2 = LLVMBuildFMul(gallivm->builder, ddy_el, emit_data->args[1], ""); + temp2 = LLVMBuildFMul(ctx->ac.builder, ddy_el, emit_data->args[1], ""); - ij_out[i] = LLVMBuildFAdd(gallivm->builder, temp2, temp1, ""); + 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) { - interp_param = LLVMBuildBitCast(gallivm->builder, - interp_param, LLVMVectorType(ctx->f32, 2), ""); - } + if (interp_param) + interp_param = ac_to_float(&ctx->ac, interp_param); for (chan = 0; chan < 4; chan++) { LLVMValueRef gather = LLVMGetUndef(LLVMVectorType(ctx->f32, input_array_size)); @@ -3771,67 +3935,33 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action, LLVMValueRef v, i = NULL, j = NULL; if (interp_param) { - interp_param = LLVMBuildBitCast(gallivm->builder, - interp_param, LLVMVectorType(ctx->f32, 2), ""); i = LLVMBuildExtractElement( - gallivm->builder, interp_param, ctx->i32_0, ""); + ctx->ac.builder, interp_param, ctx->i32_0, ""); j = LLVMBuildExtractElement( - gallivm->builder, interp_param, ctx->i32_1, ""); + ctx->ac.builder, interp_param, ctx->i32_1, ""); } v = si_build_fs_interp(ctx, input_base + idx, schan, prim_mask, i, j); - gather = LLVMBuildInsertElement(gallivm->builder, + gather = LLVMBuildInsertElement(ctx->ac.builder, gather, v, LLVMConstInt(ctx->i32, idx, false), ""); } emit_data->output[chan] = LLVMBuildExtractElement( - gallivm->builder, gather, array_idx, ""); + ctx->ac.builder, gather, array_idx, ""); } } -static LLVMValueRef si_emit_ballot(struct si_shader_context *ctx, - LLVMValueRef value) -{ - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMValueRef args[3] = { - value, - ctx->i32_0, - LLVMConstInt(ctx->i32, LLVMIntNE, 0) - }; - - /* We currently have no other way to prevent LLVM from lifting the icmp - * calls to a dominating basic block. - */ - emit_optimization_barrier(ctx, &args[0]); - - if (LLVMTypeOf(args[0]) != ctx->i32) - args[0] = LLVMBuildBitCast(gallivm->builder, args[0], ctx->i32, ""); - - return lp_build_intrinsic(gallivm->builder, - "llvm.amdgcn.icmp.i32", - ctx->i64, args, 3, - LP_FUNC_ATTR_NOUNWIND | - LP_FUNC_ATTR_READNONE | - LP_FUNC_ATTR_CONVERGENT); -} - static void vote_all_emit( const struct lp_build_tgsi_action *action, struct lp_build_tgsi_context *bld_base, struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMValueRef active_set, vote_set; - LLVMValueRef tmp; - - active_set = si_emit_ballot(ctx, ctx->i32_1); - vote_set = si_emit_ballot(ctx, emit_data->args[0]); - tmp = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, vote_set, active_set, ""); + LLVMValueRef tmp = ac_build_vote_all(&ctx->ac, emit_data->args[0]); emit_data->output[emit_data->chan] = - LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, ""); + LLVMBuildSExt(ctx->ac.builder, tmp, ctx->i32, ""); } static void vote_any_emit( @@ -3840,16 +3970,10 @@ static void vote_any_emit( struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMValueRef vote_set; - LLVMValueRef tmp; - - vote_set = si_emit_ballot(ctx, emit_data->args[0]); - tmp = LLVMBuildICmp(gallivm->builder, LLVMIntNE, - vote_set, LLVMConstInt(ctx->i64, 0, 0), ""); + LLVMValueRef tmp = ac_build_vote_any(&ctx->ac, emit_data->args[0]); emit_data->output[emit_data->chan] = - LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, ""); + LLVMBuildSExt(ctx->ac.builder, tmp, ctx->i32, ""); } static void vote_eq_emit( @@ -3858,19 +3982,10 @@ static void vote_eq_emit( struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMValueRef active_set, vote_set; - LLVMValueRef all, none, tmp; - active_set = si_emit_ballot(ctx, ctx->i32_1); - vote_set = si_emit_ballot(ctx, emit_data->args[0]); - - all = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, vote_set, active_set, ""); - none = LLVMBuildICmp(gallivm->builder, LLVMIntEQ, - vote_set, LLVMConstInt(ctx->i64, 0, 0), ""); - tmp = LLVMBuildOr(gallivm->builder, all, none, ""); + LLVMValueRef tmp = ac_build_vote_eq(&ctx->ac, emit_data->args[0]); emit_data->output[emit_data->chan] = - LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, ""); + LLVMBuildSExt(ctx->ac.builder, tmp, ctx->i32, ""); } static void ballot_emit( @@ -3879,11 +3994,11 @@ static void ballot_emit( struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMValueRef tmp; tmp = lp_build_emit_fetch(bld_base, emit_data->inst, 0, TGSI_CHAN_X); - tmp = si_emit_ballot(ctx, tmp); + tmp = ac_build_ballot(&ctx->ac, tmp); tmp = LLVMBuildBitCast(builder, tmp, ctx->v2i32, ""); emit_data->output[0] = LLVMBuildExtractElement(builder, tmp, ctx->i32_0, ""); @@ -3909,17 +4024,14 @@ static void read_lane_emit( struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - LLVMBuilderRef builder = ctx->gallivm.builder; /* We currently have no other way to prevent LLVM from lifting the icmp * calls to a dominating basic block. */ - emit_optimization_barrier(ctx, &emit_data->args[0]); + ac_build_optimization_barrier(&ctx->ac, &emit_data->args[0]); - for (unsigned i = 0; i < emit_data->arg_count; ++i) { - emit_data->args[i] = LLVMBuildBitCast(builder, emit_data->args[i], - ctx->i32, ""); - } + for (unsigned i = 0; i < emit_data->arg_count; ++i) + emit_data->args[i] = ac_to_integer(&ctx->ac, emit_data->args[i]); emit_data->output[emit_data->chan] = ac_build_intrinsic(&ctx->ac, action->intr_name, @@ -3944,29 +4056,24 @@ static unsigned si_llvm_get_stream(struct lp_build_tgsi_context *bld_base, } /* 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(gallivm->builder, + gs_next_vertex = LLVMBuildLoad(ctx->ac.builder, ctx->gs_next_vertex[stream], ""); @@ -3978,31 +4085,25 @@ static void si_llvm_emit_vertex( * further memory loads and may allow LLVM to skip to the end * altogether. */ - can_emit = LLVMBuildICmp(gallivm->builder, LLVMIntULT, gs_next_vertex, + can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex, LLVMConstInt(ctx->i32, shader->selector->gs_max_out_vertices, 0), ""); 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(gallivm->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); @@ -4011,7 +4112,7 @@ static void si_llvm_emit_vertex( voffset = lp_build_add(uint, voffset, gs_next_vertex); voffset = lp_build_mul_imm(uint, voffset, 4); - out_val = LLVMBuildBitCast(gallivm->builder, out_val, ctx->i32, ""); + out_val = ac_to_integer(&ctx->ac, out_val); ac_build_buffer_store_dword(&ctx->ac, ctx->gsvs_ring[stream], @@ -4024,7 +4125,7 @@ static void si_llvm_emit_vertex( gs_next_vertex = lp_build_add(uint, gs_next_vertex, ctx->i32_1); - LLVMBuildStore(gallivm->builder, gs_next_vertex, ctx->gs_next_vertex[stream]); + LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]); /* Signal vertex emission */ ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8), @@ -4033,6 +4134,18 @@ static void si_llvm_emit_vertex( 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, @@ -4053,19 +4166,18 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action, struct lp_build_emit_data *emit_data) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; /* SI only (thanks to a hw bug workaround): * 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; } - lp_build_intrinsic(gallivm->builder, + lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.s.barrier", ctx->voidt, NULL, 0, LP_FUNC_ATTR_CONVERGENT); } @@ -4118,7 +4230,7 @@ static void si_create_function(struct si_shader_context *ctx, "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", @@ -4159,49 +4271,16 @@ static void declare_streamout_params(struct si_shader_context *ctx, } } -static unsigned llvm_get_type_size(LLVMTypeRef type) -{ - LLVMTypeKind kind = LLVMGetTypeKind(type); - - switch (kind) { - case LLVMIntegerTypeKind: - return LLVMGetIntTypeWidth(type) / 8; - case LLVMFloatTypeKind: - return 4; - case LLVMPointerTypeKind: - return 8; - case LLVMVectorTypeKind: - return LLVMGetVectorSize(type) * - llvm_get_type_size(LLVMGetElementType(type)); - case LLVMArrayTypeKind: - return LLVMGetArrayLength(type) * - llvm_get_type_size(LLVMGetElementType(type)); - default: - assert(0); - return 0; - } -} - -static void declare_lds_as_pointer(struct si_shader_context *ctx) -{ - struct gallivm_state *gallivm = &ctx->gallivm; - - unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768; - ctx->lds = LLVMBuildIntToPtr(gallivm->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 */ @@ -4229,10 +4308,18 @@ static void declare_per_stage_desc_pointers(struct si_shader_context *ctx, 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, @@ -4244,12 +4331,13 @@ static void declare_per_stage_desc_pointers(struct si_shader_context *ctx, } } -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)); - declare_per_stage_desc_pointers(ctx, fninfo, true); + ctx->param_bindless_samplers_and_images = add_arg(fninfo, ARG_SGPR, + si_const_array(ctx->v8i32, 0)); } static void declare_vs_specific_input_sgprs(struct si_shader_context *ctx, @@ -4312,11 +4400,13 @@ static void create_function(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) @@ -4327,14 +4417,39 @@ static void create_function(struct si_shader_context *ctx) 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) { @@ -4352,7 +4467,8 @@ static void create_function(struct si_shader_context *ctx) 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); @@ -4371,14 +4487,14 @@ static void create_function(struct si_shader_context *ctx) */ for (i = 0; i < GFX6_TCS_NUM_USER_SGPR + 2; i++) returns[num_returns++] = ctx->i32; /* SGPRs */ - for (i = 0; i < 5; i++) + for (i = 0; i < 11; i++) returns[num_returns++] = ctx->f32; /* VGPRs */ break; 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); @@ -4386,8 +4502,7 @@ static void create_function(struct si_shader_context *ctx) 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 */ + 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); @@ -4424,15 +4539,15 @@ static void create_function(struct si_shader_context *ctx) */ for (i = 0; i <= 8 + GFX9_SGPR_TCS_FACTOR_ADDR_BASE64K; i++) returns[num_returns++] = ctx->i32; /* SGPRs */ - for (i = 0; i < 5; i++) + for (i = 0; i < 11; i++) returns[num_returns++] = ctx->f32; /* VGPRs */ } break; 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); @@ -4440,8 +4555,7 @@ static void create_function(struct si_shader_context *ctx) 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 */ + declare_global_desc_pointers(ctx, &fninfo); declare_per_stage_desc_pointers(ctx, &fninfo, (ctx->type == PIPE_SHADER_VERTEX || ctx->type == PIPE_SHADER_TESS_EVAL)); @@ -4464,8 +4578,8 @@ static void create_function(struct si_shader_context *ctx) /* 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) { @@ -4486,7 +4600,8 @@ static void create_function(struct si_shader_context *ctx) 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); @@ -4506,23 +4621,25 @@ static void create_function(struct si_shader_context *ctx) 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); @@ -4547,6 +4664,7 @@ static void create_function(struct si_shader_context *ctx) shader->info.face_vgpr_index = 20; add_arg_assign_checked(&fninfo, ARG_VGPR, ctx->i32, &ctx->abi.ancillary, SI_PARAM_ANCILLARY); + shader->info.ancillary_vgpr_index = 21; add_arg_assign_checked(&fninfo, ARG_VGPR, ctx->f32, &ctx->abi.sample_coverage, SI_PARAM_SAMPLE_COVERAGE); add_arg_checked(&fninfo, ARG_VGPR, ctx->i32, SI_PARAM_POS_FIXED_PT); @@ -4584,7 +4702,8 @@ static void create_function(struct si_shader_context *ctx) 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) @@ -4618,6 +4737,7 @@ static void create_function(struct si_shader_context *ctx) S_0286D0_LINEAR_CENTER_ENA(1) | S_0286D0_LINEAR_CENTROID_ENA(1) | S_0286D0_FRONT_FACE_ENA(1) | + S_0286D0_ANCILLARY_ENA(1) | S_0286D0_POS_FIXED_PT_ENA(1)); } @@ -4625,10 +4745,10 @@ static void create_function(struct si_shader_context *ctx) shader->info.num_input_vgprs = 0; for (i = 0; i < fninfo.num_sgpr_params; ++i) - shader->info.num_input_sgprs += llvm_get_type_size(fninfo.types[i]) / 4; + shader->info.num_input_sgprs += ac_get_type_size(fninfo.types[i]) / 4; for (; i < fninfo.num_params; ++i) - shader->info.num_input_vgprs += llvm_get_type_size(fninfo.types[i]) / 4; + shader->info.num_input_vgprs += ac_get_type_size(fninfo.types[i]) / 4; assert(shader->info.num_input_vgprs >= num_prolog_vgprs); shader->info.num_input_vgprs -= num_prolog_vgprs; @@ -4636,10 +4756,8 @@ static void create_function(struct si_shader_context *ctx) 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); } /** @@ -4648,13 +4766,12 @@ static void create_function(struct si_shader_context *ctx) */ static void preload_ring_buffers(struct si_shader_context *ctx) { - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMBuilderRef builder = gallivm->builder; + LLVMBuilderRef builder = ctx->ac.builder; 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 @@ -4662,20 +4779,20 @@ static void preload_ring_buffers(struct si_shader_context *ctx) 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 .. @@ -4746,8 +4863,7 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx, LLVMValueRef param_rw_buffers, unsigned param_pos_fixed_pt) { - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMBuilderRef builder = gallivm->builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMValueRef slot, desc, offset, row, bit, address[2]; /* Use the fixed-point gl_FragCoord input. @@ -4759,20 +4875,16 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx, /* 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], LLVMConstInt(ctx->i32, 4, 0), ""); row = buffer_load_const(ctx, desc, offset); - row = LLVMBuildBitCast(builder, row, ctx->i32, ""); + 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, @@ -4922,14 +5034,14 @@ int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader) 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); @@ -4956,7 +5068,7 @@ int si_shader_binary_upload(struct si_screen *sscreen, struct si_shader *shader) 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; } @@ -5018,9 +5130,20 @@ static void si_shader_dump_stats(struct si_screen *sscreen, 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 = 10; + unsigned max_simd_waves; + + switch (sscreen->info.family) { + /* These always have 8 waves: */ + case CHIP_POLARIS10: + case CHIP_POLARIS11: + case CHIP_POLARIS12: + max_simd_waves = 8; + break; + default: + max_simd_waves = 10; + } /* Compute LDS usage for PS. */ switch (processor) { @@ -5050,7 +5173,7 @@ static void si_shader_dump_stats(struct si_screen *sscreen, /* 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); @@ -5065,7 +5188,7 @@ static void si_shader_dump_stats(struct si_screen *sscreen, max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave); if (!check_debug_option || - r600_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" @@ -5137,7 +5260,7 @@ void si_shader_dump(struct si_screen *sscreen, const struct si_shader *shader, FILE *file, bool check_debug_option) { if (!check_debug_option || - r600_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) { @@ -5154,8 +5277,8 @@ void si_shader_dump(struct si_screen *sscreen, const struct si_shader *shader, } if (!check_debug_option || - (r600_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) @@ -5190,12 +5313,12 @@ static int si_compile_llvm(struct si_screen *sscreen, 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 (r600_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"); @@ -5253,9 +5376,9 @@ static int si_compile_llvm(struct si_screen *sscreen, static void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret) { if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind) - LLVMBuildRetVoid(ctx->gallivm.builder); + LLVMBuildRetVoid(ctx->ac.builder); else - LLVMBuildRet(ctx->gallivm.builder, ret); + LLVMBuildRet(ctx->ac.builder, ret); } /* Generate code for the hardware VS shader stage to go with a geometry shader */ @@ -5267,7 +5390,6 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, { 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; @@ -5286,6 +5408,9 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, 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; @@ -5294,7 +5419,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, ctx.shader = shader; ctx.type = PIPE_SHADER_VERTEX; - builder = gallivm->builder; + builder = ctx.ac.builder; create_function(&ctx); preload_ring_buffers(&ctx); @@ -5324,7 +5449,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, 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++) { @@ -5337,7 +5462,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, 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); @@ -5372,14 +5497,14 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, } 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); } LLVMPositionBuilderAtEnd(builder, end_bb); - LLVMBuildRetVoid(gallivm->builder); + LLVMBuildRetVoid(ctx.ac.builder); ctx.type = PIPE_SHADER_GEOMETRY; /* override for shader dumping */ si_llvm_optimize_module(&ctx); @@ -5390,7 +5515,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, debug, PIPE_SHADER_GEOMETRY, "GS Copy Shader"); if (!r) { - if (r600_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); @@ -5416,6 +5541,8 @@ static void si_dump_shader_key_vs(const struct si_shader_key *key, prefix, prolog->instance_divisor_is_one); fprintf(f, " %s.instance_divisor_is_fetched = %u\n", prefix, prolog->instance_divisor_is_fetched); + fprintf(f, " %s.ls_vgpr_fix = %u\n", + prefix, prolog->ls_vgpr_fix); fprintf(f, " mono.vs.fix_fetch = {"); for (int i = 0; i < SI_MAX_ATTRIBS; i++) @@ -5441,7 +5568,7 @@ static void si_dump_shader_key(unsigned processor, const struct si_shader *shade 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); } @@ -5459,7 +5586,7 @@ static void si_dump_shader_key(unsigned processor, const struct si_shader *shade 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); @@ -5509,8 +5636,6 @@ static void si_init_shader_ctx(struct si_shader_context *ctx, { struct lp_build_tgsi_context *bld_base; - ctx->abi.chip_class = sscreen->b.chip_class; - si_llvm_context_init(ctx, sscreen, tm); bld_base = &ctx->bld_base; @@ -5539,7 +5664,7 @@ static void si_init_shader_ctx(struct si_shader_context *ctx, 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; } @@ -5581,7 +5706,7 @@ static void si_count_scratch_private_memory(struct si_shader_context *ctx) LLVMTypeRef type = LLVMGetElementType(LLVMTypeOf(inst)); /* No idea why LLVM aligns allocas to 4 elements. */ unsigned alignment = LLVMGetAlignment(inst); - unsigned dw_size = align(llvm_get_type_size(type) / 4, alignment); + unsigned dw_size = align(ac_get_type_size(type) / 4, alignment); ctx->shader->config.private_mem_vgprs += dw_size; } bb = LLVMGetNextBasicBlock(bb); @@ -5591,7 +5716,7 @@ static void si_count_scratch_private_memory(struct si_shader_context *ctx) static void si_init_exec_full_mask(struct si_shader_context *ctx) { LLVMValueRef full_mask = LLVMConstInt(ctx->i64, ~0ull, 0); - lp_build_intrinsic(ctx->gallivm.builder, + lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.init.exec", ctx->voidt, &full_mask, 1, LP_FUNC_ATTR_CONVERGENT); } @@ -5603,11 +5728,19 @@ static void si_init_exec_from_input(struct si_shader_context *ctx, LLVMGetParam(ctx->main_fn, param), LLVMConstInt(ctx->i32, bitoffset, 0), }; - lp_build_intrinsic(ctx->gallivm.builder, + lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.init.exec.from.input", ctx->voidt, args, 2, LP_FUNC_ATTR_CONVERGENT); } +static bool si_vs_needs_prolog(const struct si_shader_selector *sel, + const struct si_vs_prolog_bits *key) +{ + /* VGPR initialization fixup for Vega10 and Raven is always done in the + * VS prolog. */ + return sel->vs_needs_prolog || key->ls_vgpr_fix; +} + static bool si_compile_tgsi_main(struct si_shader_context *ctx, bool is_monolithic) { @@ -5620,13 +5753,12 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx, 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; @@ -5637,15 +5769,16 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx, 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; @@ -5653,7 +5786,6 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx, bld_base->emit_epilogue = si_tgsi_emit_epilogue; break; case PIPE_SHADER_COMPUTE: - ctx->declare_memory_region = declare_compute_memory; break; default: assert(!"Unsupported shader type"); @@ -5678,13 +5810,13 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx, * 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) && (ctx->type == PIPE_SHADER_TESS_EVAL || (ctx->type == PIPE_SHADER_VERTEX && - !sel->vs_needs_prolog))) { + !si_vs_needs_prolog(sel, &shader->key.part.vs.prolog)))) { si_init_exec_from_input(ctx, ctx->param_merged_wave_info, 0); } else if (ctx->type == PIPE_SHADER_TESS_CTRL || @@ -5705,6 +5837,14 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx, } } + if (ctx->type == PIPE_SHADER_TESS_CTRL && + sel->tcs_info.tessfactors_are_def_in_all_invocs) { + for (unsigned i = 0; i < 6; i++) { + ctx->invoc0_tess_factors[i] = + lp_build_alloca_undef(&ctx->gallivm, ctx->i32, ""); + } + } + if (ctx->type == PIPE_SHADER_GEOMETRY) { int i; for (i = 0; i < 4; i++) { @@ -5714,10 +5854,11 @@ static bool si_compile_tgsi_main(struct si_shader_context *ctx, } } - 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) { @@ -5757,11 +5898,13 @@ static void si_get_vs_prolog_key(const struct tgsi_shader_info *info, 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; } @@ -5796,6 +5939,7 @@ static void si_get_ps_prolog_key(struct si_shader *shader, key->ps_prolog.states.force_linear_center_interp || key->ps_prolog.states.bc_optimize_for_persp || key->ps_prolog.states.bc_optimize_for_linear); + key->ps_prolog.ancillary_vgpr_index = shader->info.ancillary_vgpr_index; if (info->colors_read) { unsigned *color = shader->selector->color_attr_index; @@ -5905,7 +6049,8 @@ static bool si_need_ps_prolog(const union si_shader_part_key *key) key->ps_prolog.states.force_linear_center_interp || key->ps_prolog.states.bc_optimize_for_persp || key->ps_prolog.states.bc_optimize_for_linear || - key->ps_prolog.states.poly_stipple; + key->ps_prolog.states.poly_stipple || + key->ps_prolog.states.samplemask_log_ps_iter; } /** @@ -5932,15 +6077,14 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, union si_shader_part_key *key) { unsigned num_sgprs, num_vgprs; - struct gallivm_state *gallivm = &ctx->gallivm; struct si_function_info fninfo; - LLVMBuilderRef builder = gallivm->builder; + LLVMBuilderRef builder = ctx->ac.builder; LLVMTypeRef returns[48]; LLVMValueRef func, ret; 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 { @@ -5967,7 +6111,7 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, * 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, @@ -5980,7 +6124,7 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, } for (unsigned i = 0; i < num_vgprs; i++) { LLVMValueRef p = LLVMGetParam(func, num_sgprs + i); - p = LLVMBuildBitCast(builder, p, ctx->f32, ""); + p = ac_to_float(&ctx->ac, p); ret = LLVMBuildInsertValue(builder, ret, p, num_sgprs + i, ""); } @@ -6002,7 +6146,7 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, 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); @@ -6022,14 +6166,14 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, 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; hi = LLVMBuildShl(builder, vtx_out[i*2+1], LLVMConstInt(ctx->i32, 16, 0), ""); out = LLVMBuildOr(builder, vtx_out[i*2], hi, ""); - out = LLVMBuildBitCast(builder, out, ctx->f32, ""); + out = ac_to_float(&ctx->ac, out); ret = LLVMBuildInsertValue(builder, ret, out, gfx9_vtx_params[i], ""); } @@ -6037,7 +6181,7 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, for (unsigned i = 0; i < 6; i++) { LLVMValueRef out; - out = LLVMBuildBitCast(builder, vtx_out[i], ctx->f32, ""); + out = ac_to_float(&ctx->ac, vtx_out[i]); ret = LLVMBuildInsertValue(builder, ret, out, gfx6_vtx_params[i], ""); } @@ -6057,8 +6201,7 @@ static void si_build_wrapper_function(struct si_shader_context *ctx, unsigned main_part, unsigned next_shader_first_part) { - struct gallivm_state *gallivm = &ctx->gallivm; - LLVMBuilderRef builder = ctx->gallivm.builder; + LLVMBuilderRef builder = ctx->ac.builder; /* PS epilog has one arg per color component; gfx9 merged shader * prologs need to forward 32 user SGPRs. */ @@ -6096,9 +6239,9 @@ static void si_build_wrapper_function(struct si_shader_context *ctx, if (ac_is_sgpr_param(param)) { assert(num_vgprs == 0); - num_sgprs += llvm_get_type_size(LLVMTypeOf(param)) / 4; + num_sgprs += ac_get_type_size(LLVMTypeOf(param)) / 4; } else { - num_vgprs += llvm_get_type_size(LLVMTypeOf(param)) / 4; + num_vgprs += ac_get_type_size(LLVMTypeOf(param)) / 4; } } @@ -6106,7 +6249,7 @@ static void si_build_wrapper_function(struct si_shader_context *ctx, while (gprs < num_sgprs + num_vgprs) { LLVMValueRef param = LLVMGetParam(parts[main_part], fninfo.num_params); LLVMTypeRef type = LLVMTypeOf(param); - unsigned size = llvm_get_type_size(type) / 4; + unsigned size = ac_get_type_size(type) / 4; add_arg(&fninfo, gprs < num_sgprs ? ARG_SGPR : ARG_VGPR, type); @@ -6133,7 +6276,7 @@ static void si_build_wrapper_function(struct si_shader_context *ctx, LLVMValueRef param = LLVMGetParam(ctx->main_fn, i); LLVMTypeRef param_type = LLVMTypeOf(param); LLVMTypeRef out_type = i < fninfo.num_sgpr_params ? ctx->i32 : ctx->f32; - unsigned size = llvm_get_type_size(param_type) / 4; + unsigned size = ac_get_type_size(param_type) / 4; if (size == 1) { if (param_type != out_type) @@ -6195,7 +6338,7 @@ static void si_build_wrapper_function(struct si_shader_context *ctx, param = LLVMGetParam(parts[part], param_idx); param_type = LLVMTypeOf(param); - param_size = llvm_get_type_size(param_type) / 4; + param_size = ac_get_type_size(param_type) / 4; is_sgpr = ac_is_sgpr_param(param); if (is_sgpr) { @@ -6214,7 +6357,7 @@ static void si_build_wrapper_function(struct si_shader_context *ctx, 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) { @@ -6289,8 +6432,8 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, /* Dump TGSI code before doing TGSI->LLVM conversion in case the * conversion fails. */ - if (r600_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 @@ -6307,8 +6450,6 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, shader->info.uses_instanceid = sel->info.uses_instanceid; - ctx.load_system_value = declare_system_value; - if (!si_compile_tgsi_main(&ctx, is_monolithic)) { si_llvm_dispose(&ctx); return -1; @@ -6333,9 +6474,11 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, 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_vs_needs_prolog(ls, &shader->key.part.tcs.ls_prolog); /* TCS main part */ parts[2] = ctx.main_fn; @@ -6348,7 +6491,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, parts[3] = ctx.main_fn; /* VS prolog */ - if (ls->vs_needs_prolog) { + if (vs_needs_prolog) { union si_shader_part_key vs_prolog_key; si_get_vs_prolog_key(&ls->info, shader->info.num_input_sgprs, @@ -6379,9 +6522,9 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, ctx.type = PIPE_SHADER_TESS_CTRL; si_build_wrapper_function(&ctx, - parts + !ls->vs_needs_prolog, - 4 - !ls->vs_needs_prolog, 0, - ls->vs_needs_prolog ? 2 : 1); + parts + !vs_needs_prolog, + 4 - !vs_needs_prolog, 0, + vs_needs_prolog ? 2 : 1); } else { LLVMValueRef parts[2]; union si_shader_part_key epilog_key; @@ -6396,7 +6539,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, 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; @@ -6416,7 +6559,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, 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); @@ -6498,7 +6641,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, si_optimize_vs_outputs(&ctx); if ((debug && debug->debug_message) || - r600_can_dump_shader(&sscreen->b, ctx.type)) + si_can_dump_shader(sscreen, ctx.type)) si_count_scratch_private_memory(&ctx); /* Compile to bytecode. */ @@ -6516,7 +6659,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, 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); @@ -6549,6 +6692,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, if (ctx.type == PIPE_SHADER_FRAGMENT) { shader->info.num_input_vgprs = 0; shader->info.face_vgpr_index = -1; + shader->info.ancillary_vgpr_index = -1; if (G_0286CC_PERSP_SAMPLE_ENA(shader->config.spi_ps_input_addr)) shader->info.num_input_vgprs += 2; @@ -6578,8 +6722,10 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, shader->info.face_vgpr_index = shader->info.num_input_vgprs; shader->info.num_input_vgprs += 1; } - if (G_0286CC_ANCILLARY_ENA(shader->config.spi_ps_input_addr)) + if (G_0286CC_ANCILLARY_ENA(shader->config.spi_ps_input_addr)) { + shader->info.ancillary_vgpr_index = shader->info.num_input_vgprs; shader->info.num_input_vgprs += 1; + } if (G_0286CC_SAMPLE_COVERAGE_ENA(shader->config.spi_ps_input_addr)) shader->info.num_input_vgprs += 1; if (G_0286CC_POS_FIXED_PT_ENA(shader->config.spi_ps_input_addr)) @@ -6632,7 +6778,6 @@ si_get_shader_part(struct si_screen *sscreen, 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; @@ -6640,6 +6785,8 @@ si_get_shader_part(struct si_screen *sscreen, 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); @@ -6664,7 +6811,7 @@ si_get_shader_part(struct si_screen *sscreen, 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; @@ -6681,15 +6828,19 @@ 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); - list = LLVMBuildBitCast(gallivm->builder, list, ctx->i64, ""); - list = LLVMBuildIntToPtr(gallivm->builder, list, + 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), ""); return list; } @@ -6713,14 +6864,13 @@ static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx) static void si_build_vs_prolog_function(struct si_shader_context *ctx, union si_shader_part_key *key) { - struct gallivm_state *gallivm = &ctx->gallivm; struct si_function_info fninfo; LLVMTypeRef *returns; LLVMValueRef ret, func; int num_returns, i; - unsigned first_vs_vgpr = key->vs_prolog.num_input_sgprs + - key->vs_prolog.num_merged_next_stage_vgprs; + unsigned first_vs_vgpr = key->vs_prolog.num_merged_next_stage_vgprs; unsigned num_input_vgprs = key->vs_prolog.num_merged_next_stage_vgprs + 4; + LLVMValueRef input_vgprs[9]; unsigned num_all_input_regs = key->vs_prolog.num_input_sgprs + num_input_vgprs; unsigned user_sgpr_base = key->vs_prolog.num_merged_next_stage_vgprs ? 8 : 0; @@ -6740,13 +6890,10 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, /* Preloaded VGPRs (outputs must be floats) */ for (i = 0; i < num_input_vgprs; i++) { - add_arg(&fninfo, ARG_VGPR, ctx->i32); + add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &input_vgprs[i]); returns[num_returns++] = ctx->f32; } - fninfo.assign[first_vs_vgpr] = &ctx->abi.vertex_id; - fninfo.assign[first_vs_vgpr + (key->vs_prolog.as_ls ? 2 : 1)] = &ctx->abi.instance_id; - /* Vertex load indices. */ for (i = 0; i <= key->vs_prolog.last_input; i++) returns[num_returns++] = ctx->f32; @@ -6755,9 +6902,32 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, si_create_function(ctx, "vs_prolog", returns, num_returns, &fninfo, 0); func = ctx->main_fn; - if (key->vs_prolog.num_merged_next_stage_vgprs && - !key->vs_prolog.is_monolithic) - si_init_exec_from_input(ctx, 3, 0); + if (key->vs_prolog.num_merged_next_stage_vgprs) { + if (!key->vs_prolog.is_monolithic) + si_init_exec_from_input(ctx, 3, 0); + + if (key->vs_prolog.as_ls && + 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 has_hs_threads = + LLVMBuildICmp(ctx->ac.builder, LLVMIntNE, + unpack_param(ctx, 3, 8, 8), + ctx->i32_0, ""); + + for (i = 4; i > 0; --i) { + input_vgprs[i + 1] = + LLVMBuildSelect(ctx->ac.builder, has_hs_threads, + input_vgprs[i + 1], + input_vgprs[i - 1], ""); + } + } + } + + ctx->abi.vertex_id = input_vgprs[first_vs_vgpr]; + ctx->abi.instance_id = input_vgprs[first_vs_vgpr + (key->vs_prolog.as_ls ? 2 : 1)]; /* Copy inputs to outputs. This should be no-op, as the registers match, * but it will prevent the compiler from overwriting them unintentionally. @@ -6765,12 +6935,13 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, ret = ctx->return_value; for (i = 0; i < key->vs_prolog.num_input_sgprs; i++) { LLVMValueRef p = LLVMGetParam(func, i); - ret = LLVMBuildInsertValue(gallivm->builder, ret, p, i, ""); + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p, i, ""); } - for (; i < fninfo.num_params; i++) { - LLVMValueRef p = LLVMGetParam(func, i); - p = LLVMBuildBitCast(gallivm->builder, p, ctx->f32, ""); - ret = LLVMBuildInsertValue(gallivm->builder, ret, p, i, ""); + for (i = 0; i < num_input_vgprs; i++) { + LLVMValueRef p = input_vgprs[i]; + p = ac_to_float(&ctx->ac, p); + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p, + key->vs_prolog.num_input_sgprs + i, ""); } /* Compute vertex load indices from instance divisors. */ @@ -6781,7 +6952,7 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, 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++) { @@ -6797,8 +6968,7 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, if (divisor_is_fetched) { divisor = buffer_load_const(ctx, instance_divisor_constbuf, LLVMConstInt(ctx->i32, i * 4, 0)); - divisor = LLVMBuildBitCast(gallivm->builder, divisor, - ctx->i32, ""); + divisor = ac_to_integer(&ctx->ac, divisor); } /* InstanceID / Divisor + StartInstance */ @@ -6808,14 +6978,14 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, divisor); } else { /* VertexID + BaseVertex */ - index = LLVMBuildAdd(gallivm->builder, + index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id, LLVMGetParam(func, user_sgpr_base + SI_SGPR_BASE_VERTEX), ""); } - index = LLVMBuildBitCast(gallivm->builder, index, ctx->f32, ""); - ret = LLVMBuildInsertValue(gallivm->builder, ret, index, + index = ac_to_float(&ctx->ac, index); + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, index, fninfo.num_params + i, ""); } @@ -6831,8 +7001,7 @@ static bool si_get_vs_prolog(struct si_screen *sscreen, { struct si_shader_selector *vs = main_part->selector; - /* The prolog is a no-op if there are no inputs. */ - if (!vs->vs_needs_prolog) + if (!si_vs_needs_prolog(vs, key)) return true; /* Get the prolog. */ @@ -6867,14 +7036,13 @@ static bool si_shader_select_vs_parts(struct si_screen *sscreen, static void si_build_tcs_epilog_function(struct si_shader_context *ctx, union si_shader_part_key *key) { - struct gallivm_state *gallivm = &ctx->gallivm; struct lp_build_tgsi_context *bld_base = &ctx->bld_base; struct si_function_info fninfo; LLVMValueRef func; 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 */ @@ -6886,6 +7054,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, add_arg(&fninfo, ARG_SGPR, ctx->i64); add_arg(&fninfo, ARG_SGPR, ctx->i64); add_arg(&fninfo, ARG_SGPR, ctx->i64); + add_arg(&fninfo, ARG_SGPR, ctx->i64); add_arg(&fninfo, ARG_SGPR, ctx->i32); add_arg(&fninfo, ARG_SGPR, ctx->i32); add_arg(&fninfo, ARG_SGPR, ctx->i32); @@ -6899,6 +7068,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, add_arg(&fninfo, ARG_SGPR, ctx->i64); add_arg(&fninfo, ARG_SGPR, ctx->i64); add_arg(&fninfo, ARG_SGPR, ctx->i64); + add_arg(&fninfo, ARG_SGPR, ctx->i64); ctx->param_tcs_offchip_layout = add_arg(&fninfo, ARG_SGPR, ctx->i32); add_arg(&fninfo, ARG_SGPR, ctx->i32); add_arg(&fninfo, ARG_SGPR, ctx->i32); @@ -6916,18 +7086,26 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, add_arg(&fninfo, ARG_VGPR, ctx->i32); /* invocation ID within the patch */ add_arg(&fninfo, ARG_VGPR, ctx->i32); /* LDS offset where tess factors should be loaded from */ + for (unsigned i = 0; i < 6; i++) + add_arg(&fninfo, ARG_VGPR, ctx->i32); /* tess factors */ + /* 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]; + for (unsigned i = 0; i < 6; i++) + invoc0_tess_factors[i] = LLVMGetParam(func, tess_factors_idx + 3 + i); + si_write_tess_factors(bld_base, LLVMGetParam(func, tess_factors_idx), LLVMGetParam(func, tess_factors_idx + 1), - LLVMGetParam(func, tess_factors_idx + 2)); + LLVMGetParam(func, tess_factors_idx + 2), + invoc0_tess_factors, invoc0_tess_factors + 4); - LLVMBuildRetVoid(gallivm->builder); + LLVMBuildRetVoid(ctx->ac.builder); } /** @@ -6938,7 +7116,7 @@ static bool si_shader_select_tcs_parts(struct si_screen *sscreen, 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; @@ -6970,7 +7148,7 @@ static bool si_shader_select_gs_parts(struct si_screen *sscreen, 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; @@ -7010,7 +7188,6 @@ static bool si_shader_select_gs_parts(struct si_screen *sscreen, 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; @@ -7043,7 +7220,7 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, ret = ctx->return_value; for (i = 0; i < fninfo.num_params; i++) { LLVMValueRef p = LLVMGetParam(func, i); - ret = LLVMBuildInsertValue(gallivm->builder, ret, p, i, ""); + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p, i, ""); } /* Polygon stippling. */ @@ -7068,9 +7245,9 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, * PRIM_MASK is after user SGPRs. */ bc_optimize = LLVMGetParam(func, SI_PS_NUM_USER_SGPR); - bc_optimize = LLVMBuildLShr(gallivm->builder, bc_optimize, + bc_optimize = LLVMBuildLShr(ctx->ac.builder, bc_optimize, LLVMConstInt(ctx->i32, 31, 0), ""); - bc_optimize = LLVMBuildTrunc(gallivm->builder, bc_optimize, + bc_optimize = LLVMBuildTrunc(ctx->ac.builder, bc_optimize, ctx->i1, ""); if (key->ps_prolog.states.bc_optimize_for_persp) { @@ -7082,9 +7259,9 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, centroid[i] = LLVMGetParam(func, base + 4 + i); /* Select PERSP_CENTROID. */ for (i = 0; i < 2; i++) { - tmp = LLVMBuildSelect(gallivm->builder, bc_optimize, + tmp = LLVMBuildSelect(ctx->ac.builder, bc_optimize, center[i], centroid[i], ""); - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, tmp, base + 4 + i, ""); } } @@ -7097,9 +7274,9 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, centroid[i] = LLVMGetParam(func, base + 10 + i); /* Select LINEAR_CENTROID. */ for (i = 0; i < 2; i++) { - tmp = LLVMBuildSelect(gallivm->builder, bc_optimize, + tmp = LLVMBuildSelect(ctx->ac.builder, bc_optimize, center[i], centroid[i], ""); - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, tmp, base + 10 + i, ""); } } @@ -7115,11 +7292,11 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, persp_sample[i] = LLVMGetParam(func, base + i); /* Overwrite PERSP_CENTER. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, persp_sample[i], base + 2 + i, ""); /* Overwrite PERSP_CENTROID. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, persp_sample[i], base + 4 + i, ""); } if (key->ps_prolog.states.force_linear_sample_interp) { @@ -7131,11 +7308,11 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, linear_sample[i] = LLVMGetParam(func, base + 6 + i); /* Overwrite LINEAR_CENTER. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, linear_sample[i], base + 8 + i, ""); /* Overwrite LINEAR_CENTROID. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, linear_sample[i], base + 10 + i, ""); } @@ -7149,11 +7326,11 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, persp_center[i] = LLVMGetParam(func, base + 2 + i); /* Overwrite PERSP_SAMPLE. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, persp_center[i], base + i, ""); /* Overwrite PERSP_CENTROID. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, persp_center[i], base + 4 + i, ""); } if (key->ps_prolog.states.force_linear_center_interp) { @@ -7165,11 +7342,11 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, linear_center[i] = LLVMGetParam(func, base + 8 + i); /* Overwrite LINEAR_SAMPLE. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, linear_center[i], base + 6 + i, ""); /* Overwrite LINEAR_CENTROID. */ for (i = 0; i < 2; i++) - ret = LLVMBuildInsertValue(gallivm->builder, ret, + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, linear_center[i], base + 10 + i, ""); } @@ -7191,11 +7368,11 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, key->ps_prolog.color_interp_vgpr_index[i]; /* Get the (i,j) updated by bc_optimize handling. */ - interp[0] = LLVMBuildExtractValue(gallivm->builder, ret, + interp[0] = LLVMBuildExtractValue(ctx->ac.builder, ret, interp_vgpr, ""); - interp[1] = LLVMBuildExtractValue(gallivm->builder, ret, + 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. */ @@ -7203,7 +7380,7 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, if (key->ps_prolog.states.color_two_side) { face = LLVMGetParam(func, face_vgpr); - face = LLVMBuildBitCast(gallivm->builder, face, ctx->i32, ""); + face = ac_to_integer(&ctx->ac, face); } interp_fs_input(ctx, @@ -7215,11 +7392,59 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, while (writemask) { unsigned chan = u_bit_scan(&writemask); - ret = LLVMBuildInsertValue(gallivm->builder, ret, color[chan], + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, color[chan], fninfo.num_params + color_out_idx++, ""); } } + /* Section 15.2.2 (Shader Inputs) of the OpenGL 4.5 (Core Profile) spec + * says: + * + * "When per-sample shading is active due to the use of a fragment + * input qualified by sample or due to the use of the gl_SampleID + * or gl_SamplePosition variables, only the bit for the current + * sample is set in gl_SampleMaskIn. When state specifies multiple + * fragment shader invocations for a given fragment, the sample + * mask for any single fragment shader invocation may specify a + * subset of the covered samples for the fragment. In this case, + * the bit corresponding to each covered sample will be set in + * exactly one fragment shader invocation." + * + * The samplemask loaded by hardware is always the coverage of the + * entire pixel/fragment, so mask bits out based on the sample ID. + */ + if (key->ps_prolog.states.samplemask_log_ps_iter) { + /* The bit pattern matches that used by fixed function fragment + * processing. */ + static const uint16_t ps_iter_masks[] = { + 0xffff, /* not used */ + 0x5555, + 0x1111, + 0x0101, + 0x0001, + }; + assert(key->ps_prolog.states.samplemask_log_ps_iter < ARRAY_SIZE(ps_iter_masks)); + + uint32_t ps_iter_mask = ps_iter_masks[key->ps_prolog.states.samplemask_log_ps_iter]; + unsigned ancillary_vgpr = key->ps_prolog.num_input_sgprs + + key->ps_prolog.ancillary_vgpr_index; + LLVMValueRef sampleid = unpack_param(ctx, ancillary_vgpr, 8, 4); + LLVMValueRef samplemask = LLVMGetParam(func, ancillary_vgpr + 1); + + samplemask = ac_to_integer(&ctx->ac, samplemask); + samplemask = LLVMBuildAnd( + ctx->ac.builder, + samplemask, + LLVMBuildShl(ctx->ac.builder, + LLVMConstInt(ctx->i32, ps_iter_mask, false), + sampleid, ""), + ""); + samplemask = ac_to_float(&ctx->ac, samplemask); + + ret = LLVMBuildInsertValue(ctx->ac.builder, ret, samplemask, + ancillary_vgpr + 1, ""); + } + /* Tell LLVM to insert WQM instruction sequence when needed. */ if (key->ps_prolog.wqm) { LLVMAddTargetDependentFunctionAttr(func, @@ -7236,7 +7461,6 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, static void si_build_ps_epilog_function(struct si_shader_context *ctx, union si_shader_part_key *key) { - struct gallivm_state *gallivm = &ctx->gallivm; struct lp_build_tgsi_context *bld_base = &ctx->bld_base; struct si_function_info fninfo; LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL; @@ -7247,6 +7471,7 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx, /* Declare input SGPRs. */ ctx->param_rw_buffers = add_arg(&fninfo, ARG_SGPR, ctx->i64); + ctx->param_bindless_samplers_and_images = add_arg(&fninfo, ARG_SGPR, ctx->i64); ctx->param_const_and_shader_buffers = add_arg(&fninfo, ARG_SGPR, ctx->i64); ctx->param_samplers_and_images = add_arg(&fninfo, ARG_SGPR, ctx->i64); add_arg_checked(&fninfo, ARG_SGPR, ctx->f32, SI_PARAM_ALPHA_REF); @@ -7325,7 +7550,7 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx, si_emit_ps_exports(ctx, &exp); /* Compile. */ - LLVMBuildRetVoid(gallivm->builder); + LLVMBuildRetVoid(ctx->ac.builder); } /** @@ -7415,6 +7640,12 @@ static bool si_shader_select_ps_parts(struct si_screen *sscreen, assert(G_0286CC_LINEAR_CENTER_ENA(shader->config.spi_ps_input_addr)); } + /* Samplemask fixup requires the sample ID. */ + if (shader->key.part.ps.prolog.samplemask_log_ps_iter) { + shader->config.spi_ps_input_ena |= S_0286CC_ANCILLARY_ENA(1); + assert(G_0286CC_ANCILLARY_ENA(shader->config.spi_ps_input_addr)); + } + /* The sample mask input is always enabled, because the API shader always * passes it through to the epilog. Disable it here if it's unused. */ @@ -7432,9 +7663,9 @@ void si_multiwave_lds_size_workaround(struct si_screen *sscreen, * 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); } @@ -7499,6 +7730,7 @@ int si_shader_create(struct si_screen *sscreen, LLVMTargetMachineRef tm, shader->info.num_input_sgprs = mainp->info.num_input_sgprs; shader->info.num_input_vgprs = mainp->info.num_input_vgprs; shader->info.face_vgpr_index = mainp->info.face_vgpr_index; + shader->info.ancillary_vgpr_index = mainp->info.ancillary_vgpr_index; memcpy(shader->info.vs_output_param_offset, mainp->info.vs_output_param_offset, sizeof(mainp->info.vs_output_param_offset)); @@ -7597,7 +7829,7 @@ void si_shader_destroy(struct si_shader *shader) r600_resource_reference(&shader->bo, NULL); if (!shader->is_binary_shared) - radeon_shader_binary_clean(&shader->binary); + ac_shader_binary_clean(&shader->binary); free(shader->shader_log); }