X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fgallium%2Fdrivers%2Fradeonsi%2Fsi_shader.c;h=c3b5f58cd26a9bc788bdc0aea401f898a76f8ebd;hb=7ef1e42c14fb23592e8e003f7a80db9a43cb9bc9;hp=f0e580943e8715277448ef4577ef89a7ce9ced95;hpb=d0751f6c1ffe146b0fa36d1290b40551d2cc65cd;p=mesa.git diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index f0e580943e8..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" @@ -110,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 || @@ -412,7 +407,7 @@ static LLVMValueRef get_tcs_in_vertex_dw_stride(struct si_shader_context *ctx) return LLVMConstInt(ctx->i32, stride * 4, 0); case PIPE_SHADER_TESS_CTRL: - if (ctx->screen->b.chip_class >= GFX9 && + if (ctx->screen->info.chip_class >= GFX9 && ctx->shader->is_monolithic) { stride = util_last_bit64(ctx->shader->key.part.tcs.ls->outputs_written); return LLVMConstInt(ctx->i32, stride * 4, 0); @@ -446,7 +441,7 @@ static LLVMValueRef extract_double_to_float(struct si_shader_context *ctx, unsigned double_index) { LLVMBuilderRef builder = ctx->ac.builder; - LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->gallivm.context); + LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->ac.context); LLVMValueRef dvec2 = LLVMBuildBitCast(builder, vec4, LLVMVectorType(f64, 2), ""); LLVMValueRef index = LLVMConstInt(ctx->i32, double_index, 0); @@ -454,11 +449,97 @@ 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]) { + unsigned vs_blit_property = + ctx->shader->selector->info.properties[TGSI_PROPERTY_VS_BLIT_SGPRS]; + + if (vs_blit_property) { + LLVMValueRef vertex_id = ctx->abi.vertex_id; + LLVMValueRef sel_x1 = LLVMBuildICmp(ctx->ac.builder, + LLVMIntULE, vertex_id, + ctx->i32_1, ""); + /* Use LLVMIntNE, because we have 3 vertices and only + * the middle one should use y2. + */ + LLVMValueRef sel_y1 = LLVMBuildICmp(ctx->ac.builder, + LLVMIntNE, vertex_id, + ctx->i32_1, ""); + + if (input_index == 0) { + /* Position: */ + LLVMValueRef x1y1 = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs); + LLVMValueRef x2y2 = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 1); + + LLVMValueRef x1 = unpack_sint16(ctx, x1y1, 0); + LLVMValueRef y1 = unpack_sint16(ctx, x1y1, 1); + LLVMValueRef x2 = unpack_sint16(ctx, x2y2, 0); + LLVMValueRef y2 = unpack_sint16(ctx, x2y2, 1); + + LLVMValueRef x = LLVMBuildSelect(ctx->ac.builder, sel_x1, + x1, x2, ""); + LLVMValueRef y = LLVMBuildSelect(ctx->ac.builder, sel_y1, + y1, y2, ""); + + out[0] = LLVMBuildSIToFP(ctx->ac.builder, x, ctx->f32, ""); + out[1] = LLVMBuildSIToFP(ctx->ac.builder, y, ctx->f32, ""); + out[2] = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 2); + out[3] = ctx->ac.f32_1; + return; + } + + /* Color or texture coordinates: */ + assert(input_index == 1); + + if (vs_blit_property == SI_VS_BLIT_SGPRS_POS_COLOR) { + for (int i = 0; i < 4; i++) { + out[i] = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 3 + i); + } + } else { + assert(vs_blit_property == SI_VS_BLIT_SGPRS_POS_TEXCOORD); + LLVMValueRef x1 = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 3); + LLVMValueRef y1 = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 4); + LLVMValueRef x2 = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 5); + LLVMValueRef y2 = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 6); + + out[0] = LLVMBuildSelect(ctx->ac.builder, sel_x1, + x1, x2, ""); + out[1] = LLVMBuildSelect(ctx->ac.builder, sel_y1, + y1, y2, ""); + out[2] = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 7); + out[3] = LLVMGetParam(ctx->main_fn, + ctx->param_vs_blit_inputs + 8); + } + return; + } + unsigned chan; unsigned fix_fetch; unsigned num_fetches; @@ -475,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 + @@ -678,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; @@ -983,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); } /** @@ -998,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) { @@ -1007,21 +1087,24 @@ 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 bitcast(bld_base, type, value); } @@ -1033,18 +1116,14 @@ 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); - - dw_addr = lp_build_add(&bld_base->uint_bld, dw_addr, + dw_addr = lp_build_add(&ctx->bld_base.uint_bld, dw_addr, LLVMConstInt(ctx->i32, dw_offset_imm, 0)); - value = ac_to_integer(&ctx->ac, value); - ac_build_indexed_store(&ctx->ac, ctx->lds, - dw_addr, value); + ac_lds_store(&ctx->ac, dw_addr, value); } static LLVMValueRef desc_from_addr_base64k(struct si_shader_context *ctx, @@ -1129,7 +1208,6 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base, LLVMValueRef dst[4]) { struct si_shader_context *ctx = si_shader_context(bld_base); - struct gallivm_state *gallivm = &ctx->gallivm; const struct tgsi_full_dst_register *reg = &inst->Dst[index]; const struct tgsi_shader_info *sh_info = &ctx->shader->selector->info; unsigned chan_index; @@ -1188,7 +1266,7 @@ 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 = ac_to_integer(&ctx->ac, value); values[chan_index] = value; @@ -1213,7 +1291,7 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base, } if (reg->Register.WriteMask == 0xF && !is_tess_factor) { - LLVMValueRef value = lp_build_gather_values(gallivm, + LLVMValueRef value = lp_build_gather_values(&ctx->gallivm, values, 4); ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr, base, 0, 1, 0, true, false); @@ -1229,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]; @@ -1246,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) { @@ -1279,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); @@ -1307,7 +1377,7 @@ 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 bitcast(bld_base, type, value); @@ -1532,10 +1602,9 @@ 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; 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); @@ -1548,7 +1617,7 @@ 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); } void si_load_system_value(struct si_shader_context *ctx, @@ -1556,7 +1625,6 @@ void si_load_system_value(struct si_shader_context *ctx, const struct tgsi_full_declaration *decl) { struct lp_build_context *bld = &ctx->bld_base.base; - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef value = 0; assert(index < RADEON_LLVM_MAX_SYSTEM_VALUES); @@ -1607,8 +1675,7 @@ void si_load_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; @@ -1623,7 +1690,7 @@ void si_load_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; } @@ -1646,7 +1713,7 @@ void si_load_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; } @@ -1662,17 +1729,17 @@ void si_load_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; } @@ -1711,13 +1778,13 @@ void si_load_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; } @@ -1745,7 +1812,7 @@ void si_load_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); } @@ -1763,7 +1830,7 @@ void si_load_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; } @@ -1832,22 +1899,21 @@ 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(ctx->ac.builder, var, i8p, ""); + ctx->ac.lds = LLVMBuildBitCast(ctx->ac.builder, var, i8p, ""); } static LLVMValueRef load_const_buffer_desc(struct si_shader_context *ctx, int i) @@ -1855,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) @@ -1868,7 +1934,7 @@ static LLVMValueRef load_ubo(struct ac_shader_abi *abi, LLVMValueRef index) index = LLVMBuildAdd(ctx->ac.builder, index, LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS, 0), ""); - return ac_build_indexed_load_const(&ctx->ac, ptr, index); + return ac_build_load_to_sgpr(&ctx->ac, ptr, index); } static LLVMValueRef @@ -1883,7 +1949,7 @@ load_ssbo(struct ac_shader_abi *abi, LLVMValueRef index, bool write) 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( @@ -1893,11 +1959,11 @@ static LLVMValueRef fetch_constant( unsigned swizzle) { struct si_shader_context *ctx = si_shader_context(bld_base); + struct si_shader_selector *sel = ctx->shader->selector; const struct tgsi_ind_register *ireg = ®->Indirect; unsigned buf, idx; LLVMValueRef addr, bufp; - LLVMValueRef result; if (swizzle == LP_CHAN_ALL) { unsigned chan; @@ -1908,9 +1974,77 @@ static LLVMValueRef fetch_constant( return lp_build_gather_values(&ctx->gallivm, values, 4); } + /* Split 64-bit loads. */ + if (tgsi_type_is_64bit(type)) { + LLVMValueRef lo, hi; + + lo = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, swizzle); + hi = fetch_constant(bld_base, reg, TGSI_TYPE_UNSIGNED, swizzle + 1); + return si_llvm_emit_fetch_64bit(bld_base, tgsi2llvmtype(bld_base, type), + lo, hi); + } + + idx = reg->Register.Index * 4 + swizzle; + if (reg->Register.Indirect) { + addr = si_get_indirect_index(ctx, ireg, 16, idx * 4); + } else { + addr = LLVMConstInt(ctx->i32, idx * 4, 0); + } + + /* Fast path when user data SGPRs point to constant buffer 0 directly. */ + if (sel->info.const_buffers_declared == 1 && + sel->info.shader_buffers_declared == 0) { + LLVMValueRef ptr = + LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers); + + /* This enables use of s_load_dword and flat_load_dword for const buffer 0 + * loads, and up to x4 load opcode merging. However, it leads to horrible + * code reducing SIMD wave occupancy from 8 to 2 in many cases. + * + * Using s_buffer_load_dword (x1) seems to be the best option right now. + * + * LLVM 5.0 on SI doesn't insert a required s_nop between SALU setting + * a descriptor and s_buffer_load_dword using it, so we can't expand + * the pointer into a full descriptor like below. We have to use + * s_load_dword instead. The only case when LLVM 5.0 would select + * s_buffer_load_dword (that we have to prevent) is when we use use + * a literal offset where we don't need bounds checking. + */ + if (ctx->screen->info.chip_class == SI && + HAVE_LLVM < 0x0600 && + !reg->Register.Indirect) { + addr = LLVMBuildLShr(ctx->ac.builder, addr, LLVMConstInt(ctx->i32, 2, 0), ""); + LLVMValueRef result = ac_build_load_invariant(&ctx->ac, ptr, addr); + return bitcast(bld_base, type, result); + } + + /* Do the bounds checking with a descriptor, because + * doing computation and manual bounds checking of 64-bit + * addresses generates horrible VALU code with very high + * VGPR usage and very low SIMD occupancy. + */ + ptr = LLVMBuildPtrToInt(ctx->ac.builder, ptr, ctx->i64, ""); + ptr = LLVMBuildBitCast(ctx->ac.builder, ptr, ctx->v2i32, ""); + + LLVMValueRef desc_elems[] = { + LLVMBuildExtractElement(ctx->ac.builder, ptr, ctx->i32_0, ""), + LLVMBuildExtractElement(ctx->ac.builder, ptr, ctx->i32_1, ""), + LLVMConstInt(ctx->i32, (sel->info.const_file_max[0] + 1) * 16, 0), + LLVMConstInt(ctx->i32, + S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) | + S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) | + S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) | + S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) | + S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) | + S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32), 0) + }; + LLVMValueRef desc = ac_build_gather_values(&ctx->ac, desc_elems, 4); + LLVMValueRef result = buffer_load_const(ctx, desc, addr); + return bitcast(bld_base, type, result); + } + assert(reg->Register.Dimension); buf = reg->Dimension.Index; - idx = reg->Register.Index * 4 + swizzle; if (reg->Dimension.Indirect) { LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers); @@ -1920,31 +2054,11 @@ static LLVMValueRef fetch_constant( ctx->num_const_buffers); index = LLVMBuildAdd(ctx->ac.builder, index, LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS, 0), ""); - bufp = ac_build_indexed_load_const(&ctx->ac, ptr, index); + bufp = ac_build_load_to_sgpr(&ctx->ac, ptr, index); } else bufp = load_const_buffer_desc(ctx, buf); - if (reg->Register.Indirect) { - addr = si_get_indirect_index(ctx, ireg, 16, idx * 4); - } else { - addr = LLVMConstInt(ctx->i32, idx * 4, 0); - } - - result = buffer_load_const(ctx, bufp, addr); - - if (!tgsi_type_is_64bit(type)) - result = bitcast(bld_base, type, result); - else { - LLVMValueRef addr2, result2; - - addr2 = lp_build_add(&bld_base->uint_bld, addr, - LLVMConstInt(ctx->i32, 4, 0)); - result2 = buffer_load_const(ctx, bufp, addr2); - - result = si_llvm_emit_fetch_64bit(bld_base, type, - result, result2); - } - return result; + return bitcast(bld_base, type, buffer_load_const(ctx, bufp, addr)); } /* Upper 16 bits must be zero. */ @@ -1970,13 +2084,12 @@ 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; + 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; @@ -2007,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: @@ -2069,10 +2182,10 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, 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]. */ @@ -2082,7 +2195,7 @@ 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, ""); @@ -2102,7 +2215,7 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, /* Clamp. */ for (chan = 0; chan < 4; chan++) { val[chan] = ac_to_integer(&ctx->ac, values[chan]); - val[chan] = lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_UMIN, + val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_UMIN, val[chan], chan == 3 ? max_alpha : max_rgb); } @@ -2126,10 +2239,10 @@ static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base, /* Clamp. */ for (chan = 0; chan < 4; chan++) { val[chan] = ac_to_integer(&ctx->ac, values[chan]); - val[chan] = lp_build_emit_llvm_binary(bld_base, + val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_IMIN, val[chan], chan == 3 ? max_alpha : max_rgb); - val[chan] = lp_build_emit_llvm_binary(bld_base, + val[chan] = lp_build_emit_llvm_binary(&ctx->bld_base, TGSI_OPCODE_IMAX, val[chan], chan == 3 ? min_alpha : min_rgb); } @@ -2152,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)); } } @@ -2197,11 +2312,9 @@ static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context * 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; @@ -2209,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]; @@ -2228,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])); } } @@ -2320,7 +2433,6 @@ 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 = ctx->ac.builder; int i; struct lp_build_if_state if_ctx; @@ -2338,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 + @@ -2367,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]); @@ -2400,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); } @@ -2453,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; @@ -2469,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: @@ -2487,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; @@ -2508,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). */ @@ -2527,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; @@ -2549,7 +2659,7 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base, 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]. */ @@ -2651,7 +2761,6 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, 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; @@ -2669,7 +2778,7 @@ 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, + lp_build_if(&if_ctx, &ctx->gallivm, LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, invocation_id, ctx->i32_0, "")); @@ -2741,11 +2850,11 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, } /* 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); @@ -2756,13 +2865,13 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base, byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id, LLVMConstInt(ctx->i32, 4 * stride, 0), ""); - lp_build_if(&inner_if_ctx, gallivm, + lp_build_if(&inner_if_ctx, &ctx->gallivm, LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, rel_patch_id, ctx->i32_0, "")); /* Store the dynamic HS control word. */ offset = 0; - if (ctx->screen->b.chip_class <= VI) { + if (ctx->screen->info.chip_class <= VI) { ac_build_buffer_store_dword(&ctx->ac, buffer, LLVMConstInt(ctx->i32, 0x80000000, 0), 1, ctx->i32_0, tf_base, @@ -2796,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, @@ -2809,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); @@ -2869,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 @@ -2895,7 +3004,7 @@ static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base) 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, @@ -2954,11 +3063,13 @@ 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); @@ -2995,11 +3106,12 @@ 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); @@ -3018,9 +3130,11 @@ 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; unsigned i, chan; @@ -3033,7 +3147,6 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base) /* 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]; @@ -3061,18 +3174,23 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base) LLVMConstInt(ctx->i32, param * 4, 0), ""); for (chan = 0; chan < 4; chan++) { - lds_store(bld_base, chan, dw_addr, - LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], "")); + if (!(info->output_usagemask[i] & (1 << chan))) + continue; + + lds_store(ctx, chan, dw_addr, + LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], "")); } } - if (ctx->screen->b.chip_class >= GFX9) + if (ctx->screen->info.chip_class >= GFX9) si_set_ls_return_value_for_tcs(ctx); } -static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base) +static void si_llvm_emit_es_epilogue(struct ac_shader_abi *abi, + unsigned max_outputs, + LLVMValueRef *addrs) { - struct si_shader_context *ctx = si_shader_context(bld_base); + struct si_shader_context *ctx = si_shader_context_from_abi(abi); struct si_shader *es = ctx->shader; struct tgsi_shader_info *info = &es->selector->info; LLVMValueRef soffset = LLVMGetParam(ctx->main_fn, @@ -3081,7 +3199,7 @@ 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); @@ -3093,7 +3211,6 @@ static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base) } 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 || @@ -3104,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(ctx->ac.builder, out_ptr[chan], ""); + LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], ""); out_val = ac_to_integer(&ctx->ac, out_val); /* GFX9 has the ESGS ring in LDS. */ - if (ctx->screen->b.chip_class >= GFX9) { - lds_store(bld_base, param * 4 + chan, lds_base, out_val); + if (ctx->screen->info.chip_class >= GFX9) { + lds_store(ctx, param * 4 + chan, lds_base, out_val); continue; } @@ -3121,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; @@ -3182,7 +3314,7 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi, ctx->param_vs_state_bits); cond = LLVMBuildTrunc(ctx->ac.builder, cond, ctx->i1, ""); - lp_build_if(&if_ctx, gallivm, cond); + lp_build_if(&if_ctx, &ctx->gallivm, cond); } for (j = 0; j < 4; j++) { @@ -3227,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); } @@ -3323,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 */ @@ -3340,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 */ @@ -3350,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 && @@ -3369,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; @@ -3389,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 */ @@ -3455,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++) { @@ -3619,7 +3750,6 @@ 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; @@ -3630,7 +3760,7 @@ static LLVMValueRef si_llvm_emit_ddxy_interp( 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( @@ -3638,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) { @@ -3684,7 +3813,7 @@ static void interp_fetch_args( ctx->ac.f32_0, }; - sample_position = lp_build_gather_values(gallivm, center, 4); + sample_position = lp_build_gather_values(&ctx->gallivm, center, 4); } else { sample_position = load_sample_position(ctx, sample_id); } @@ -3708,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; @@ -3793,7 +3921,7 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action, 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) @@ -3928,26 +4056,21 @@ 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(ctx->ac.builder, @@ -3968,25 +4091,19 @@ static void si_llvm_emit_vertex( bool use_kill = !info->writes_memory; if (use_kill) { - kill = lp_build_select(&bld_base->base, can_emit, - LLVMConstReal(ctx->f32, 1.0f), - LLVMConstReal(ctx->f32, -1.0f)); - - ac_build_kill(&ctx->ac, kill); + ac_build_kill_if_false(&ctx->ac, can_emit); } else { - lp_build_if(&if_state, gallivm, can_emit); + lp_build_if(&if_state, &ctx->gallivm, can_emit); } offset = 0; for (i = 0; i < info->num_outputs; i++) { - LLVMValueRef *out_ptr = ctx->outputs[i]; - for (chan = 0; chan < 4; chan++) { if (!(info->output_usagemask[i] & (1 << chan)) || ((info->output_streams[i] >> (2 * chan)) & 3) != stream) continue; - LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], ""); + LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, addrs[4 * i + chan], ""); LLVMValueRef voffset = LLVMConstInt(ctx->i32, offset * shader->selector->gs_max_out_vertices, 0); @@ -4017,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, @@ -4042,7 +4171,7 @@ static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action, * The real barrier instruction isn’t needed, because an entire patch * always fits into a single wave. */ - if (ctx->screen->b.chip_class == SI && + if (ctx->screen->info.chip_class == SI && ctx->type == PIPE_SHADER_TESS_CTRL) { si_emit_waitcnt(ctx, LGKM_CNT & VM_CNT); return; @@ -4101,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", @@ -4142,24 +4271,16 @@ static void declare_streamout_params(struct si_shader_context *ctx, } } -static void declare_lds_as_pointer(struct si_shader_context *ctx) -{ - unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768; - ctx->lds = LLVMBuildIntToPtr(ctx->ac.builder, ctx->i32_0, - LLVMPointerType(LLVMArrayType(ctx->i32, lds_size / 4), LOCAL_ADDR_SPACE), - "lds"); -} - static unsigned si_get_max_workgroup_size(const struct si_shader *shader) { switch (shader->selector->type) { case PIPE_SHADER_TESS_CTRL: /* Return this so that LLVM doesn't remove s_barrier * instructions on chips where we use s_barrier. */ - return shader->selector->screen->b.chip_class >= CIK ? 128 : 64; + return shader->selector->screen->info.chip_class >= CIK ? 128 : 64; case PIPE_SHADER_GEOMETRY: - return shader->selector->screen->b.chip_class >= GFX9 ? 128 : 64; + return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 64; case PIPE_SHADER_COMPUTE: break; /* see below */ @@ -4187,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, @@ -4202,14 +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)); ctx->param_bindless_samplers_and_images = add_arg(fninfo, ARG_SGPR, si_const_array(ctx->v8i32, 0)); - declare_per_stage_desc_pointers(ctx, fninfo, true); } static void declare_vs_specific_input_sgprs(struct si_shader_context *ctx, @@ -4272,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) @@ -4287,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) { @@ -4312,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); @@ -4337,8 +4493,8 @@ static void create_function(struct si_shader_context *ctx) 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); @@ -4346,12 +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 */ - - ctx->param_bindless_samplers_and_images = - add_arg(&fninfo, ARG_SGPR, si_const_array(ctx->v8i32, 0)); - + declare_global_desc_pointers(ctx, &fninfo); declare_per_stage_desc_pointers(ctx, &fninfo, ctx->type == PIPE_SHADER_VERTEX); declare_vs_specific_input_sgprs(ctx, &fninfo); @@ -4395,8 +4546,8 @@ static void create_function(struct si_shader_context *ctx) 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); @@ -4404,12 +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 */ - - ctx->param_bindless_samplers_and_images = - add_arg(&fninfo, ARG_SGPR, si_const_array(ctx->v8i32, 0)); - + declare_global_desc_pointers(ctx, &fninfo); declare_per_stage_desc_pointers(ctx, &fninfo, (ctx->type == PIPE_SHADER_VERTEX || ctx->type == PIPE_SHADER_TESS_EVAL)); @@ -4432,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) { @@ -4454,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); @@ -4474,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); @@ -4553,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) @@ -4606,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); } /** @@ -4623,7 +4771,7 @@ static void preload_ring_buffers(struct si_shader_context *ctx) 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 @@ -4631,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 .. @@ -4727,7 +4875,7 @@ 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], @@ -4736,11 +4884,7 @@ static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx, 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, @@ -4890,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); @@ -4924,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; } @@ -4986,11 +5130,11 @@ 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; - switch (sscreen->b.family) { + switch (sscreen->info.family) { /* These always have 8 waves: */ case CHIP_POLARIS10: case CHIP_POLARIS11: @@ -5029,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); @@ -5044,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 || - si_can_dump_shader(&sscreen->b, processor)) { + si_can_dump_shader(sscreen, processor)) { if (processor == PIPE_SHADER_FRAGMENT) { fprintf(file, "*** SHADER CONFIG ***\n" "SPI_PS_INPUT_ADDR = 0x%04x\n" @@ -5116,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 || - si_can_dump_shader(&sscreen->b, processor)) + si_can_dump_shader(sscreen, processor)) si_dump_shader_key(processor, shader, file); if (!check_debug_option && shader->binary.llvm_ir_string) { @@ -5133,8 +5277,8 @@ void si_shader_dump(struct si_screen *sscreen, const struct si_shader *shader, } if (!check_debug_option || - (si_can_dump_shader(&sscreen->b, processor) && - !(sscreen->b.debug_flags & DBG_NO_ASM))) { + (si_can_dump_shader(sscreen, processor) && + !(sscreen->debug_flags & DBG(NO_ASM)))) { fprintf(file, "\n%s:\n", si_get_shader_name(shader, processor)); if (shader->prolog) @@ -5169,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 (si_can_dump_shader(&sscreen->b, processor)) { + if (si_can_dump_shader(sscreen, processor)) { fprintf(stderr, "radeonsi: Compiling shader %d\n", count); - if (!(sscreen->b.debug_flags & (DBG_NO_IR | DBG_PREOPT_IR))) { + if (!(sscreen->debug_flags & (DBG(NO_IR) | DBG(PREOPT_IR)))) { fprintf(stderr, "%s LLVM IR:\n\n", name); ac_dump_module(mod); fprintf(stderr, "\n"); @@ -5246,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; @@ -5265,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; @@ -5303,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++) { @@ -5316,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); @@ -5351,7 +5497,7 @@ 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); } @@ -5369,7 +5515,7 @@ si_generate_gs_copy_shader(struct si_screen *sscreen, debug, PIPE_SHADER_GEOMETRY, "GS Copy Shader"); if (!r) { - if (si_can_dump_shader(&sscreen->b, PIPE_SHADER_GEOMETRY)) + if (si_can_dump_shader(sscreen, PIPE_SHADER_GEOMETRY)) fprintf(stderr, "GS Copy Shader:\n"); si_shader_dump(sscreen, ctx.shader, debug, PIPE_SHADER_GEOMETRY, stderr, true); @@ -5422,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); } @@ -5440,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); @@ -5518,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; } @@ -5607,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; @@ -5624,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; @@ -5664,7 +5810,7 @@ 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) && @@ -5708,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) { @@ -5751,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; } @@ -5935,7 +6084,7 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx, 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 { @@ -5962,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, @@ -5997,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); @@ -6017,7 +6166,7 @@ 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; @@ -6052,7 +6201,6 @@ 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->ac.builder; /* PS epilog has one arg per color component; gfx9 merged shader * prologs need to forward 32 user SGPRs. @@ -6209,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) { @@ -6284,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 (si_can_dump_shader(&sscreen->b, sel->info.processor) && - !(sscreen->b.debug_flags & DBG_NO_TGSI)) { + if (si_can_dump_shader(sscreen, sel->info.processor) && + !(sscreen->debug_flags & DBG(NO_TGSI))) { if (sel->tokens) tgsi_dump(sel->tokens, 0); else @@ -6326,7 +6474,7 @@ 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 = @@ -6391,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; @@ -6411,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); @@ -6493,7 +6641,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, si_optimize_vs_outputs(&ctx); if ((debug && debug->debug_message) || - si_can_dump_shader(&sscreen->b, ctx.type)) + si_can_dump_shader(sscreen, ctx.type)) si_count_scratch_private_memory(&ctx); /* Compile to bytecode. */ @@ -6511,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); @@ -6630,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; @@ -6638,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); @@ -6662,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; @@ -6679,13 +6828,17 @@ out: static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx) { - struct gallivm_state *gallivm = &ctx->gallivm; LLVMValueRef ptr[2], list; + bool is_merged_shader = + ctx->screen->info.chip_class >= GFX9 && + (ctx->type == PIPE_SHADER_TESS_CTRL || + ctx->type == PIPE_SHADER_GEOMETRY || + ctx->shader->key.as_ls || ctx->shader->key.as_es); /* Get the pointer to rw buffers. */ - ptr[0] = LLVMGetParam(ctx->main_fn, SI_SGPR_RW_BUFFERS); - ptr[1] = LLVMGetParam(ctx->main_fn, SI_SGPR_RW_BUFFERS_HI); - list = lp_build_gather_values(gallivm, ptr, 2); + ptr[0] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS); + ptr[1] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS_HI); + list = lp_build_gather_values(&ctx->gallivm, ptr, 2); list = LLVMBuildBitCast(ctx->ac.builder, list, ctx->i64, ""); list = LLVMBuildIntToPtr(ctx->ac.builder, list, si_const_array(ctx->v4i32, SI_NUM_RW_BUFFERS), ""); @@ -6754,8 +6907,7 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx, si_init_exec_from_input(ctx, 3, 0); if (key->vs_prolog.as_ls && - (ctx->screen->b.family == CHIP_VEGA10 || - ctx->screen->b.family == CHIP_RAVEN)) { + ctx->screen->has_ls_vgpr_init_bug) { /* If there are no HS threads, SPI loads the LS VGPRs * starting at VGPR 0. Shift them back to where they * belong. @@ -6800,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++) { @@ -6890,7 +7042,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, 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 */ @@ -6939,8 +7091,8 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, /* 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]; @@ -6964,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; @@ -6996,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; @@ -7036,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; @@ -7221,7 +7372,7 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx, interp_vgpr, ""); interp[1] = LLVMBuildExtractValue(ctx->ac.builder, ret, interp_vgpr + 1, ""); - interp_ij = lp_build_gather_values(gallivm, interp, 2); + interp_ij = lp_build_gather_values(&ctx->gallivm, interp, 2); } /* Use the absolute location of the input. */ @@ -7512,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); } @@ -7678,7 +7829,7 @@ void si_shader_destroy(struct si_shader *shader) r600_resource_reference(&shader->bo, NULL); if (!shader->is_binary_shared) - si_radeon_shader_binary_clean(&shader->binary); + ac_shader_binary_clean(&shader->binary); free(shader->shader_log); }