RADEON_LLVM_AMDGPU_GS = 88,
RADEON_LLVM_AMDGPU_PS = 89,
RADEON_LLVM_AMDGPU_CS = 90,
+ RADEON_LLVM_AMDGPU_HS = 93,
};
#define CONST_ADDR_SPACE 2
LLVMValueRef descriptor_sets[AC_UD_MAX_SETS];
LLVMValueRef ring_offsets;
LLVMValueRef push_constants;
+ LLVMValueRef view_index;
LLVMValueRef num_work_groups;
LLVMValueRef workgroup_ids;
LLVMValueRef local_invocation_ids;
LLVMValueRef tcs_out_layout;
LLVMValueRef tcs_in_layout;
LLVMValueRef oc_lds;
+ LLVMValueRef merged_wave_info;
LLVMValueRef tess_factor_offset;
LLVMValueRef tcs_patch_id;
LLVMValueRef tcs_rel_ids;
unreachable("illegal slot in get unique index\n");
}
-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));
- default:
- assert(0);
- return 0;
- }
-}
-
static void set_llvm_calling_convention(LLVMValueRef func,
gl_shader_stage stage)
{
switch (stage) {
case MESA_SHADER_VERTEX:
- case MESA_SHADER_TESS_CTRL:
case MESA_SHADER_TESS_EVAL:
calling_conv = RADEON_LLVM_AMDGPU_VS;
break;
case MESA_SHADER_GEOMETRY:
calling_conv = RADEON_LLVM_AMDGPU_GS;
break;
+ case MESA_SHADER_TESS_CTRL:
+ calling_conv = HAVE_LLVM >= 0x0500 ? RADEON_LLVM_AMDGPU_HS : RADEON_LLVM_AMDGPU_VS;
+ break;
case MESA_SHADER_FRAGMENT:
calling_conv = RADEON_LLVM_AMDGPU_PS;
break;
LLVMTypeRef type, LLVMValueRef *param_ptr)
{
add_argument(info, type, param_ptr);
- info->num_sgprs_used += llvm_get_type_size(type) / 4;
+ info->num_sgprs_used += ac_get_type_size(type) / 4;
info->sgpr_count++;
}
LLVMValueRef *param_ptr)
{
add_sgpr_argument(info, type, param_ptr);
- info->num_user_sgprs_used += llvm_get_type_size(type) / 4;
+ info->num_user_sgprs_used += ac_get_type_size(type) / 4;
info->user_sgpr_count++;
}
LLVMValueRef *param_ptr)
{
add_argument(info, type, param_ptr);
- info->num_vgprs_used += llvm_get_type_size(type) / 4;
+ info->num_vgprs_used += ac_get_type_size(type) / 4;
}
static inline void
CONST_ADDR_SPACE);
}
-static LLVMTypeRef to_integer_type_scalar(struct ac_llvm_context *ctx, LLVMTypeRef t)
-{
- if (t == ctx->f16 || t == ctx->i16)
- return ctx->i16;
- else if (t == ctx->f32 || t == ctx->i32)
- return ctx->i32;
- else if (t == ctx->f64 || t == ctx->i64)
- return ctx->i64;
- else
- unreachable("Unhandled integer size");
-}
-
-static LLVMTypeRef to_integer_type(struct ac_llvm_context *ctx, LLVMTypeRef t)
-{
- if (LLVMGetTypeKind(t) == LLVMVectorTypeKind) {
- LLVMTypeRef elem_type = LLVMGetElementType(t);
- return LLVMVectorType(to_integer_type_scalar(ctx, elem_type),
- LLVMGetVectorSize(t));
- }
- return to_integer_type_scalar(ctx, t);
-}
-
-static LLVMValueRef to_integer(struct ac_llvm_context *ctx, LLVMValueRef v)
-{
- LLVMTypeRef type = LLVMTypeOf(v);
- return LLVMBuildBitCast(ctx->builder, v, to_integer_type(ctx, type), "");
-}
-
-static LLVMTypeRef to_float_type_scalar(struct ac_llvm_context *ctx, LLVMTypeRef t)
-{
- if (t == ctx->i16 || t == ctx->f16)
- return ctx->f16;
- else if (t == ctx->i32 || t == ctx->f32)
- return ctx->f32;
- else if (t == ctx->i64 || t == ctx->f64)
- return ctx->f64;
- else
- unreachable("Unhandled float size");
-}
-
-static LLVMTypeRef to_float_type(struct ac_llvm_context *ctx, LLVMTypeRef t)
-{
- if (LLVMGetTypeKind(t) == LLVMVectorTypeKind) {
- LLVMTypeRef elem_type = LLVMGetElementType(t);
- return LLVMVectorType(to_float_type_scalar(ctx, elem_type),
- LLVMGetVectorSize(t));
- }
- return to_float_type_scalar(ctx, t);
-}
-
-static LLVMValueRef to_float(struct ac_llvm_context *ctx, LLVMValueRef v)
-{
- LLVMTypeRef type = LLVMTypeOf(v);
- return LLVMBuildBitCast(ctx->builder, v, to_float_type(ctx, type), "");
-}
-
static int get_elem_bits(struct ac_llvm_context *ctx, LLVMTypeRef type)
{
if (LLVMGetTypeKind(type) == LLVMVectorTypeKind)
}
}
-static void create_function(struct nir_to_llvm_context *ctx)
+static void
+radv_define_common_user_sgprs_phase1(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage,
+ bool has_previous_stage,
+ gl_shader_stage previous_stage,
+ const struct user_sgpr_info *user_sgpr_info,
+ struct arg_info *args,
+ LLVMValueRef *desc_sets)
{
unsigned num_sets = ctx->options->layout ? ctx->options->layout->num_sets : 0;
- uint8_t user_sgpr_idx;
- struct user_sgpr_info user_sgpr_info;
- struct arg_info args = {};
- LLVMValueRef desc_sets;
-
- allocate_user_sgprs(ctx, &user_sgpr_info);
- if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
- add_user_sgpr_argument(&args, const_array(ctx->v4i32, 16), &ctx->ring_offsets); /* address of rings */
- }
+ unsigned stage_mask = 1 << stage;
+ if (has_previous_stage)
+ stage_mask |= 1 << previous_stage;
/* 1 for each descriptor set */
- if (!user_sgpr_info.indirect_all_descriptor_sets) {
+ if (!user_sgpr_info->indirect_all_descriptor_sets) {
for (unsigned i = 0; i < num_sets; ++i) {
- if (ctx->options->layout->set[i].layout->shader_stages & (1 << ctx->stage)) {
- add_user_sgpr_array_argument(&args, const_array(ctx->i8, 1024 * 1024), &ctx->descriptor_sets[i]);
+ if (ctx->options->layout->set[i].layout->shader_stages & stage_mask) {
+ add_user_sgpr_array_argument(args, const_array(ctx->i8, 1024 * 1024), &ctx->descriptor_sets[i]);
}
}
} else
- add_user_sgpr_array_argument(&args, const_array(const_array(ctx->i8, 1024 * 1024), 32), &desc_sets);
+ add_user_sgpr_array_argument(args, const_array(const_array(ctx->i8, 1024 * 1024), 32), desc_sets);
if (ctx->shader_info->info.needs_push_constants) {
/* 1 for push constants and dynamic descriptors */
- add_user_sgpr_array_argument(&args, const_array(ctx->i8, 1024 * 1024), &ctx->push_constants);
+ add_user_sgpr_array_argument(args, const_array(ctx->i8, 1024 * 1024), &ctx->push_constants);
}
+}
- switch (ctx->stage) {
+static void
+radv_define_common_user_sgprs_phase2(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage,
+ bool has_previous_stage,
+ gl_shader_stage previous_stage,
+ const struct user_sgpr_info *user_sgpr_info,
+ LLVMValueRef desc_sets,
+ uint8_t *user_sgpr_idx)
+{
+ unsigned num_sets = ctx->options->layout ? ctx->options->layout->num_sets : 0;
+ unsigned stage_mask = 1 << stage;
+ if (has_previous_stage)
+ stage_mask |= 1 << previous_stage;
+
+ if (!user_sgpr_info->indirect_all_descriptor_sets) {
+ for (unsigned i = 0; i < num_sets; ++i) {
+ if (ctx->options->layout->set[i].layout->shader_stages & stage_mask) {
+ set_userdata_location(&ctx->shader_info->user_sgprs_locs.descriptor_sets[i], user_sgpr_idx, 2);
+ } else
+ ctx->descriptor_sets[i] = NULL;
+ }
+ } else {
+ uint32_t desc_sgpr_idx = *user_sgpr_idx;
+ set_userdata_location_shader(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS, user_sgpr_idx, 2);
+
+ for (unsigned i = 0; i < num_sets; ++i) {
+ if (ctx->options->layout->set[i].layout->shader_stages & stage_mask) {
+ set_userdata_location_indirect(&ctx->shader_info->user_sgprs_locs.descriptor_sets[i], desc_sgpr_idx, 2, i * 8);
+ ctx->descriptor_sets[i] = ac_build_load_to_sgpr(&ctx->ac, desc_sets, LLVMConstInt(ctx->i32, i, false));
+
+ } else
+ ctx->descriptor_sets[i] = NULL;
+ }
+ ctx->shader_info->need_indirect_descriptor_sets = true;
+ }
+
+ if (ctx->shader_info->info.needs_push_constants) {
+ set_userdata_location_shader(ctx, AC_UD_PUSH_CONSTANTS, user_sgpr_idx, 2);
+ }
+}
+
+static void
+radv_define_vs_user_sgprs_phase1(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage,
+ bool has_previous_stage,
+ gl_shader_stage previous_stage,
+ struct arg_info *args)
+{
+ if (!ctx->is_gs_copy_shader && (stage == MESA_SHADER_VERTEX || (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
+ if (ctx->shader_info->info.vs.has_vertex_buffers)
+ add_user_sgpr_argument(args, const_array(ctx->v4i32, 16), &ctx->vertex_buffers); /* vertex buffers */
+ add_user_sgpr_argument(args, ctx->i32, &ctx->abi.base_vertex); // base vertex
+ add_user_sgpr_argument(args, ctx->i32, &ctx->abi.start_instance);// start instance
+ if (ctx->shader_info->info.vs.needs_draw_id)
+ add_user_sgpr_argument(args, ctx->i32, &ctx->abi.draw_id); // draw id
+ }
+}
+
+static void
+radv_define_vs_user_sgprs_phase2(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage,
+ bool has_previous_stage,
+ gl_shader_stage previous_stage,
+ uint8_t *user_sgpr_idx)
+{
+ if (!ctx->is_gs_copy_shader && (stage == MESA_SHADER_VERTEX || (has_previous_stage && previous_stage == MESA_SHADER_VERTEX))) {
+ if (ctx->shader_info->info.vs.has_vertex_buffers) {
+ set_userdata_location_shader(ctx, AC_UD_VS_VERTEX_BUFFERS, user_sgpr_idx, 2);
+ }
+ unsigned vs_num = 2;
+ if (ctx->shader_info->info.vs.needs_draw_id)
+ vs_num++;
+
+ set_userdata_location_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE, user_sgpr_idx, vs_num);
+ }
+}
+
+
+static void create_function(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage,
+ bool has_previous_stage,
+ gl_shader_stage previous_stage)
+{
+ uint8_t user_sgpr_idx;
+ struct user_sgpr_info user_sgpr_info;
+ struct arg_info args = {};
+ LLVMValueRef desc_sets;
+
+ allocate_user_sgprs(ctx, &user_sgpr_info);
+
+ if (user_sgpr_info.need_ring_offsets && !ctx->options->supports_spill) {
+ add_user_sgpr_argument(&args, const_array(ctx->v4i32, 16), &ctx->ring_offsets); /* address of rings */
+ }
+
+ switch (stage) {
case MESA_SHADER_COMPUTE:
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
if (ctx->shader_info->info.cs.grid_components_used)
add_user_sgpr_argument(&args, LLVMVectorType(ctx->i32, ctx->shader_info->info.cs.grid_components_used), &ctx->num_work_groups); /* grid size */
add_sgpr_argument(&args, LLVMVectorType(ctx->i32, 3), &ctx->workgroup_ids);
add_vgpr_argument(&args, LLVMVectorType(ctx->i32, 3), &ctx->local_invocation_ids);
break;
case MESA_SHADER_VERTEX:
- if (!ctx->is_gs_copy_shader) {
- if (ctx->shader_info->info.vs.has_vertex_buffers)
- add_user_sgpr_argument(&args, const_array(ctx->v4i32, 16), &ctx->vertex_buffers); /* vertex buffers */
- add_user_sgpr_argument(&args, ctx->i32, &ctx->abi.base_vertex); // base vertex
- add_user_sgpr_argument(&args, ctx->i32, &ctx->abi.start_instance);// start instance
- if (ctx->shader_info->info.vs.needs_draw_id)
- add_user_sgpr_argument(&args, ctx->i32, &ctx->abi.draw_id); // draw id
- }
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
+ radv_define_vs_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &args);
+ if (ctx->shader_info->info.needs_multiview_view_index || (!ctx->options->key.vs.as_es && !ctx->options->key.vs.as_ls && ctx->options->key.has_multiview_view_index))
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->view_index);
if (ctx->options->key.vs.as_es)
add_sgpr_argument(&args, ctx->i32, &ctx->es2gs_offset); // es2gs offset
else if (ctx->options->key.vs.as_ls)
}
break;
case MESA_SHADER_TESS_CTRL:
- add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
- add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_offsets); // tcs out offsets
- add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_layout); // tcs out layout
- add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_in_layout); // tcs in layout
- add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // param oc lds
- add_sgpr_argument(&args, ctx->i32, &ctx->tess_factor_offset); // tess factor offset
- add_vgpr_argument(&args, ctx->i32, &ctx->tcs_patch_id); // patch id
- add_vgpr_argument(&args, ctx->i32, &ctx->tcs_rel_ids); // rel ids;
+ if (has_previous_stage) {
+ // First 6 system regs
+ add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // param oc lds
+ add_sgpr_argument(&args, ctx->i32, &ctx->merged_wave_info); // merged wave info
+ add_sgpr_argument(&args, ctx->i32, &ctx->tess_factor_offset); // tess factor offset
+
+ add_sgpr_argument(&args, ctx->i32, NULL); // scratch offset
+ add_sgpr_argument(&args, ctx->i32, NULL); // unknown
+ add_sgpr_argument(&args, ctx->i32, NULL); // unknown
+
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
+ radv_define_vs_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &args);
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->ls_out_layout); // ls out layout
+
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_offsets); // tcs out offsets
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_layout); // tcs out layout
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_in_layout); // tcs in layout
+ if (ctx->shader_info->info.needs_multiview_view_index)
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->view_index);
+
+ add_vgpr_argument(&args, ctx->i32, &ctx->tcs_patch_id); // patch id
+ add_vgpr_argument(&args, ctx->i32, &ctx->tcs_rel_ids); // rel ids;
+ add_vgpr_argument(&args, ctx->i32, &ctx->abi.vertex_id); // vertex id
+ add_vgpr_argument(&args, ctx->i32, &ctx->rel_auto_id); // rel auto id
+ add_vgpr_argument(&args, ctx->i32, &ctx->vs_prim_id); // vs prim id
+ add_vgpr_argument(&args, ctx->i32, &ctx->abi.instance_id); // instance id
+ } else {
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_offsets); // tcs out offsets
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_layout); // tcs out layout
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_in_layout); // tcs in layout
+ if (ctx->shader_info->info.needs_multiview_view_index)
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->view_index);
+ add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // param oc lds
+ add_sgpr_argument(&args, ctx->i32, &ctx->tess_factor_offset); // tess factor offset
+ add_vgpr_argument(&args, ctx->i32, &ctx->tcs_patch_id); // patch id
+ add_vgpr_argument(&args, ctx->i32, &ctx->tcs_rel_ids); // rel ids;
+ }
break;
case MESA_SHADER_TESS_EVAL:
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
+ if (ctx->shader_info->info.needs_multiview_view_index || (!ctx->options->key.tes.as_es && ctx->options->key.has_multiview_view_index))
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->view_index);
if (ctx->options->key.tes.as_es) {
add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // OC LDS
add_sgpr_argument(&args, ctx->i32, NULL); //
add_vgpr_argument(&args, ctx->i32, &ctx->tes_patch_id); // tes patch id
break;
case MESA_SHADER_GEOMETRY:
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
+ radv_define_vs_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &args);
add_user_sgpr_argument(&args, ctx->i32, &ctx->gsvs_ring_stride); // gsvs stride
add_user_sgpr_argument(&args, ctx->i32, &ctx->gsvs_num_entries); // gsvs num entires
+ if (ctx->shader_info->info.needs_multiview_view_index)
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->view_index);
add_sgpr_argument(&args, ctx->i32, &ctx->gs2vs_offset); // gs2vs offset
add_sgpr_argument(&args, ctx->i32, &ctx->gs_wave_id); // wave id
add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[0]); // vtx0
add_vgpr_argument(&args, ctx->i32, &ctx->gs_invocation_id);
break;
case MESA_SHADER_FRAGMENT:
+ radv_define_common_user_sgprs_phase1(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, &args, &desc_sets);
if (ctx->shader_info->info.ps.needs_sample_positions)
add_user_sgpr_argument(&args, ctx->i32, &ctx->sample_pos_offset); /* sample position offset */
add_sgpr_argument(&args, ctx->i32, &ctx->prim_mask); /* prim mask */
ctx->context, ctx->module, ctx->builder, NULL, 0, &args,
ctx->max_workgroup_size,
ctx->options->unsafe_math);
- set_llvm_calling_convention(ctx->main_function, ctx->stage);
+ set_llvm_calling_convention(ctx->main_function, stage);
ctx->shader_info->num_input_vgprs = 0;
- ctx->shader_info->num_input_sgprs = ctx->shader_info->num_user_sgprs =
- ctx->options->supports_spill ? 2 : 0;
+ ctx->shader_info->num_input_sgprs = ctx->options->supports_spill ? 2 : 0;
- ctx->shader_info->num_user_sgprs += args.num_user_sgprs_used;
ctx->shader_info->num_input_sgprs += args.num_sgprs_used;
if (ctx->stage != MESA_SHADER_FRAGMENT)
const_array(ctx->v4i32, 16), "");
}
}
+
+ /* For merged shaders the user SGPRs start at 8, with 8 system SGPRs in front (including
+ * the rw_buffers at s0/s1. With user SGPR0 = s8, lets restart the count from 0 */
+ if (has_previous_stage)
+ user_sgpr_idx = 0;
- if (!user_sgpr_info.indirect_all_descriptor_sets) {
- for (unsigned i = 0; i < num_sets; ++i) {
- if (ctx->options->layout->set[i].layout->shader_stages & (1 << ctx->stage)) {
- set_userdata_location(&ctx->shader_info->user_sgprs_locs.descriptor_sets[i], &user_sgpr_idx, 2);
- } else
- ctx->descriptor_sets[i] = NULL;
- }
- } else {
- uint32_t desc_sgpr_idx = user_sgpr_idx;
- set_userdata_location_shader(ctx, AC_UD_INDIRECT_DESCRIPTOR_SETS, &user_sgpr_idx, 2);
-
- for (unsigned i = 0; i < num_sets; ++i) {
- if (ctx->options->layout->set[i].layout->shader_stages & (1 << ctx->stage)) {
- set_userdata_location_indirect(&ctx->shader_info->user_sgprs_locs.descriptor_sets[i], desc_sgpr_idx, 2, i * 8);
- ctx->descriptor_sets[i] = ac_build_indexed_load_const(&ctx->ac, desc_sets, LLVMConstInt(ctx->i32, i, false));
-
- } else
- ctx->descriptor_sets[i] = NULL;
- }
- ctx->shader_info->need_indirect_descriptor_sets = true;
- }
-
- if (ctx->shader_info->info.needs_push_constants) {
- set_userdata_location_shader(ctx, AC_UD_PUSH_CONSTANTS, &user_sgpr_idx, 2);
- }
+ radv_define_common_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, desc_sets, &user_sgpr_idx);
- switch (ctx->stage) {
+ switch (stage) {
case MESA_SHADER_COMPUTE:
if (ctx->shader_info->info.cs.grid_components_used) {
set_userdata_location_shader(ctx, AC_UD_CS_GRID_SIZE, &user_sgpr_idx, ctx->shader_info->info.cs.grid_components_used);
}
break;
case MESA_SHADER_VERTEX:
- if (!ctx->is_gs_copy_shader) {
- if (ctx->shader_info->info.vs.has_vertex_buffers) {
- set_userdata_location_shader(ctx, AC_UD_VS_VERTEX_BUFFERS, &user_sgpr_idx, 2);
- }
- unsigned vs_num = 2;
- if (ctx->shader_info->info.vs.needs_draw_id)
- vs_num++;
-
- set_userdata_location_shader(ctx, AC_UD_VS_BASE_VERTEX_START_INSTANCE, &user_sgpr_idx, vs_num);
- }
+ radv_define_vs_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_idx);
+ if (ctx->view_index)
+ set_userdata_location_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
if (ctx->options->key.vs.as_ls) {
set_userdata_location_shader(ctx, AC_UD_VS_LS_TCS_IN_LAYOUT, &user_sgpr_idx, 1);
}
declare_tess_lds(ctx);
break;
case MESA_SHADER_TESS_CTRL:
+ radv_define_vs_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_idx);
+ if (has_previous_stage)
+ set_userdata_location_shader(ctx, AC_UD_VS_LS_TCS_IN_LAYOUT, &user_sgpr_idx, 1);
set_userdata_location_shader(ctx, AC_UD_TCS_OFFCHIP_LAYOUT, &user_sgpr_idx, 4);
+ if (ctx->view_index)
+ set_userdata_location_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
declare_tess_lds(ctx);
break;
case MESA_SHADER_TESS_EVAL:
set_userdata_location_shader(ctx, AC_UD_TES_OFFCHIP_LAYOUT, &user_sgpr_idx, 1);
+ if (ctx->view_index)
+ set_userdata_location_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
break;
case MESA_SHADER_GEOMETRY:
+ radv_define_vs_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_idx);
set_userdata_location_shader(ctx, AC_UD_GS_VS_RING_STRIDE_ENTRIES, &user_sgpr_idx, 2);
+ if (ctx->view_index)
+ set_userdata_location_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
break;
case MESA_SHADER_FRAGMENT:
if (ctx->shader_info->info.ps.needs_sample_positions) {
default:
unreachable("Shader stage not implemented");
}
+
+ ctx->shader_info->num_user_sgprs = user_sgpr_idx;
}
static void setup_types(struct nir_to_llvm_context *ctx)
LLVMValueRef src1)
{
LLVMValueRef result;
- src0 = to_float(ctx, src0);
- src1 = to_float(ctx, src1);
+ src0 = ac_to_float(ctx, src0);
+ src1 = ac_to_float(ctx, src1);
result = LLVMBuildFCmp(ctx->builder, pred, src0, src1, "");
return LLVMBuildSelect(ctx->builder, result,
LLVMConstInt(ctx->i32, 0xFFFFFFFF, false),
{
char name[64];
LLVMValueRef params[] = {
- to_float(ctx, src0),
+ ac_to_float(ctx, src0),
};
MAYBE_UNUSED const int length = snprintf(name, sizeof(name), "%s.f%d", intrin,
{
char name[64];
LLVMValueRef params[] = {
- to_float(ctx, src0),
- to_float(ctx, src1),
+ ac_to_float(ctx, src0),
+ ac_to_float(ctx, src1),
};
MAYBE_UNUSED const int length = snprintf(name, sizeof(name), "%s.f%d", intrin,
{
char name[64];
LLVMValueRef params[] = {
- to_float(ctx, src0),
- to_float(ctx, src1),
- to_float(ctx, src2),
+ ac_to_float(ctx, src0),
+ ac_to_float(ctx, src1),
+ ac_to_float(ctx, src2),
};
MAYBE_UNUSED const int length = snprintf(name, sizeof(name), "%s.f%d", intrin,
LLVMValueRef src0)
{
const char *intr = "llvm.floor.f32";
- LLVMValueRef fsrc0 = to_float(ctx, src0);
+ LLVMValueRef fsrc0 = ac_to_float(ctx, src0);
LLVMValueRef params[] = {
fsrc0,
};
static LLVMValueRef emit_f2b(struct ac_llvm_context *ctx,
LLVMValueRef src0)
{
- src0 = to_float(ctx, src0);
+ src0 = ac_to_float(ctx, src0);
return LLVMBuildSExt(ctx->builder,
LLVMBuildFCmp(ctx->builder, LLVMRealUNE, src0, ctx->f32_0, ""),
ctx->i32, "");
LLVMValueRef src0)
{
LLVMValueRef result;
- LLVMValueRef cond;
+ LLVMValueRef cond = NULL;
- src0 = to_float(&ctx->ac, src0);
+ src0 = ac_to_float(&ctx->ac, src0);
result = LLVMBuildFPTrunc(ctx->builder, src0, ctx->f16, "");
if (ctx->options->chip_class >= VI) {
int i;
LLVMValueRef comp[2];
- src0 = to_float(ctx, src0);
+ src0 = ac_to_float(ctx, src0);
comp[0] = LLVMBuildExtractElement(ctx->builder, src0, ctx->i32_0, "");
comp[1] = LLVMBuildExtractElement(ctx->builder, src0, ctx->i32_1, "");
for (i = 0; i < 2; i++) {
unsigned mask;
int idx;
LLVMValueRef result;
- bool has_ds_bpermute = ctx->abi->chip_class >= VI;
if (op == nir_op_fddx_fine || op == nir_op_fddx)
mask = AC_TID_MASK_LEFT;
else
idx = 2;
- result = ac_build_ddxy(&ctx->ac, has_ds_bpermute,
- mask, idx,
- src0);
+ result = ac_build_ddxy(&ctx->ac, mask, idx, src0);
return result;
}
result = src[0];
break;
case nir_op_fneg:
- src[0] = to_float(&ctx->ac, src[0]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
result = LLVMBuildFNeg(ctx->ac.builder, src[0], "");
break;
case nir_op_ineg:
result = LLVMBuildAdd(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_fadd:
- src[0] = to_float(&ctx->ac, src[0]);
- src[1] = to_float(&ctx->ac, src[1]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ src[1] = ac_to_float(&ctx->ac, src[1]);
result = LLVMBuildFAdd(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_fsub:
- src[0] = to_float(&ctx->ac, src[0]);
- src[1] = to_float(&ctx->ac, src[1]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ src[1] = ac_to_float(&ctx->ac, src[1]);
result = LLVMBuildFSub(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_isub:
result = LLVMBuildURem(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_fmod:
- src[0] = to_float(&ctx->ac, src[0]);
- src[1] = to_float(&ctx->ac, src[1]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ src[1] = ac_to_float(&ctx->ac, src[1]);
result = ac_build_fdiv(&ctx->ac, src[0], src[1]);
result = emit_intrin_1f_param(&ctx->ac, "llvm.floor",
- to_float_type(&ctx->ac, def_type), result);
+ ac_to_float_type(&ctx->ac, def_type), result);
result = LLVMBuildFMul(ctx->ac.builder, src[1] , result, "");
result = LLVMBuildFSub(ctx->ac.builder, src[0], result, "");
break;
case nir_op_frem:
- src[0] = to_float(&ctx->ac, src[0]);
- src[1] = to_float(&ctx->ac, src[1]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ src[1] = ac_to_float(&ctx->ac, src[1]);
result = LLVMBuildFRem(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_irem:
result = LLVMBuildUDiv(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_fmul:
- src[0] = to_float(&ctx->ac, src[0]);
- src[1] = to_float(&ctx->ac, src[1]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ src[1] = ac_to_float(&ctx->ac, src[1]);
result = LLVMBuildFMul(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_fdiv:
- src[0] = to_float(&ctx->ac, src[0]);
- src[1] = to_float(&ctx->ac, src[1]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ src[1] = ac_to_float(&ctx->ac, src[1]);
result = ac_build_fdiv(&ctx->ac, src[0], src[1]);
break;
case nir_op_frcp:
- src[0] = to_float(&ctx->ac, src[0]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
result = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, src[0]);
break;
case nir_op_iand:
break;
case nir_op_fabs:
result = emit_intrin_1f_param(&ctx->ac, "llvm.fabs",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_iabs:
result = emit_iabs(&ctx->ac, src[0]);
result = emit_isign(&ctx->ac, src[0]);
break;
case nir_op_fsign:
- src[0] = to_float(&ctx->ac, src[0]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
result = emit_fsign(&ctx->ac, src[0]);
break;
case nir_op_ffloor:
result = emit_intrin_1f_param(&ctx->ac, "llvm.floor",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_ftrunc:
result = emit_intrin_1f_param(&ctx->ac, "llvm.trunc",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_fceil:
result = emit_intrin_1f_param(&ctx->ac, "llvm.ceil",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_fround_even:
result = emit_intrin_1f_param(&ctx->ac, "llvm.rint",
- to_float_type(&ctx->ac, def_type),src[0]);
+ ac_to_float_type(&ctx->ac, def_type),src[0]);
break;
case nir_op_ffract:
result = emit_ffract(&ctx->ac, src[0]);
break;
case nir_op_fsin:
result = emit_intrin_1f_param(&ctx->ac, "llvm.sin",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_fcos:
result = emit_intrin_1f_param(&ctx->ac, "llvm.cos",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_fsqrt:
result = emit_intrin_1f_param(&ctx->ac, "llvm.sqrt",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_fexp2:
result = emit_intrin_1f_param(&ctx->ac, "llvm.exp2",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_flog2:
result = emit_intrin_1f_param(&ctx->ac, "llvm.log2",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
break;
case nir_op_frsq:
result = emit_intrin_1f_param(&ctx->ac, "llvm.sqrt",
- to_float_type(&ctx->ac, def_type), src[0]);
+ ac_to_float_type(&ctx->ac, def_type), src[0]);
result = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, result);
break;
case nir_op_fpow:
result = emit_intrin_2f_param(&ctx->ac, "llvm.pow",
- to_float_type(&ctx->ac, def_type), src[0], src[1]);
+ ac_to_float_type(&ctx->ac, def_type), src[0], src[1]);
break;
case nir_op_fmax:
result = emit_intrin_2f_param(&ctx->ac, "llvm.maxnum",
- to_float_type(&ctx->ac, def_type), src[0], src[1]);
+ ac_to_float_type(&ctx->ac, def_type), src[0], src[1]);
if (instr->dest.dest.ssa.bit_size == 32)
result = emit_intrin_1f_param(&ctx->ac, "llvm.canonicalize",
- to_float_type(&ctx->ac, def_type),
+ ac_to_float_type(&ctx->ac, def_type),
result);
break;
case nir_op_fmin:
result = emit_intrin_2f_param(&ctx->ac, "llvm.minnum",
- to_float_type(&ctx->ac, def_type), src[0], src[1]);
+ ac_to_float_type(&ctx->ac, def_type), src[0], src[1]);
if (instr->dest.dest.ssa.bit_size == 32)
result = emit_intrin_1f_param(&ctx->ac, "llvm.canonicalize",
- to_float_type(&ctx->ac, def_type),
+ ac_to_float_type(&ctx->ac, def_type),
result);
break;
case nir_op_ffma:
- result = emit_intrin_3f_param(&ctx->ac, "llvm.fma",
- to_float_type(&ctx->ac, def_type), src[0], src[1], src[2]);
+ result = emit_intrin_3f_param(&ctx->ac, "llvm.fmuladd",
+ ac_to_float_type(&ctx->ac, def_type), src[0], src[1], src[2]);
break;
case nir_op_ibitfield_extract:
result = emit_bitfield_extract(&ctx->ac, true, src);
case nir_op_vec3:
case nir_op_vec4:
for (unsigned i = 0; i < nir_op_infos[instr->op].num_inputs; i++)
- src[i] = to_integer(&ctx->ac, src[i]);
+ src[i] = ac_to_integer(&ctx->ac, src[i]);
result = ac_build_gather_values(&ctx->ac, src, num_components);
break;
case nir_op_f2i32:
case nir_op_f2i64:
- src[0] = to_float(&ctx->ac, src[0]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
result = LLVMBuildFPToSI(ctx->ac.builder, src[0], def_type, "");
break;
case nir_op_f2u32:
case nir_op_f2u64:
- src[0] = to_float(&ctx->ac, src[0]);
+ src[0] = ac_to_float(&ctx->ac, src[0]);
result = LLVMBuildFPToUI(ctx->ac.builder, src[0], def_type, "");
break;
case nir_op_i2f32:
case nir_op_i2f64:
- result = LLVMBuildSIToFP(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), "");
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
+ result = LLVMBuildSIToFP(ctx->ac.builder, src[0], ac_to_float_type(&ctx->ac, def_type), "");
break;
case nir_op_u2f32:
case nir_op_u2f64:
- result = LLVMBuildUIToFP(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), "");
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
+ result = LLVMBuildUIToFP(ctx->ac.builder, src[0], ac_to_float_type(&ctx->ac, def_type), "");
break;
case nir_op_f2f64:
- result = LLVMBuildFPExt(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), "");
+ result = LLVMBuildFPExt(ctx->ac.builder, src[0], ac_to_float_type(&ctx->ac, def_type), "");
break;
case nir_op_f2f32:
- result = LLVMBuildFPTrunc(ctx->ac.builder, src[0], to_float_type(&ctx->ac, def_type), "");
+ result = LLVMBuildFPTrunc(ctx->ac.builder, src[0], ac_to_float_type(&ctx->ac, def_type), "");
break;
case nir_op_u2u32:
case nir_op_u2u64:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
if (get_elem_bits(&ctx->ac, LLVMTypeOf(src[0])) < get_elem_bits(&ctx->ac, def_type))
result = LLVMBuildZExt(ctx->ac.builder, src[0], def_type, "");
else
break;
case nir_op_i2i32:
case nir_op_i2i64:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
if (get_elem_bits(&ctx->ac, LLVMTypeOf(src[0])) < get_elem_bits(&ctx->ac, def_type))
result = LLVMBuildSExt(ctx->ac.builder, src[0], def_type, "");
else
result = emit_bcsel(&ctx->ac, src[0], src[1], src[2]);
break;
case nir_op_find_lsb:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
result = emit_find_lsb(&ctx->ac, src[0]);
break;
case nir_op_ufind_msb:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
result = emit_ufind_msb(&ctx->ac, src[0]);
break;
case nir_op_ifind_msb:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
result = emit_ifind_msb(&ctx->ac, src[0]);
break;
case nir_op_uadd_carry:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
+ src[1] = ac_to_integer(&ctx->ac, src[1]);
result = emit_uint_carry(&ctx->ac, "llvm.uadd.with.overflow.i32", src[0], src[1]);
break;
case nir_op_usub_borrow:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
+ src[1] = ac_to_integer(&ctx->ac, src[1]);
result = emit_uint_carry(&ctx->ac, "llvm.usub.with.overflow.i32", src[0], src[1]);
break;
case nir_op_b2f:
result = emit_b2i(&ctx->ac, src[0]);
break;
case nir_op_i2b:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
result = emit_i2b(&ctx->ac, src[0]);
break;
case nir_op_fquantize2f16:
result = emit_f2f16(ctx->nctx, src[0]);
break;
case nir_op_umul_high:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
+ src[1] = ac_to_integer(&ctx->ac, src[1]);
result = emit_umul_high(&ctx->ac, src[0], src[1]);
break;
case nir_op_imul_high:
+ src[0] = ac_to_integer(&ctx->ac, src[0]);
+ src[1] = ac_to_integer(&ctx->ac, src[1]);
result = emit_imul_high(&ctx->ac, src[0], src[1]);
break;
case nir_op_pack_half_2x16:
if (result) {
assert(instr->dest.dest.is_ssa);
- result = to_integer(&ctx->ac, result);
+ result = ac_to_integer(&ctx->ac, result);
_mesa_hash_table_insert(ctx->defs, &instr->dest.dest.ssa,
result);
}
LLVMConstInt(ctx->ac.i32, 2, false), "");
/* VI only */
- if (ctx->abi->chip_class >= VI && in_elements) {
+ if (ctx->ac.chip_class == VI && in_elements) {
/* On VI, the descriptor contains the size in bytes,
* but TXQ must return the size in elements.
* The stride is always non-zero for resources using TXQ.
enum glsl_base_type stype = glsl_get_sampler_result_type(instr->texture->var->type);
LLVMValueRef coord = args->addr;
LLVMValueRef half_texel[2];
- LLVMValueRef compare_cube_wa;
+ LLVMValueRef compare_cube_wa = NULL;
LLVMValueRef result;
int c;
unsigned coord_vgpr_index = (unsigned)args->offset + (unsigned)args->compare;
break;
}
- if (instr->op == nir_texop_tg4) {
+ if (instr->op == nir_texop_tg4 && ctx->ac.chip_class <= VI) {
enum glsl_base_type stype = glsl_get_sampler_result_type(instr->texture->var->type);
if (stype == GLSL_TYPE_UINT || stype == GLSL_TYPE_INT) {
return radv_lower_gather4_integer(&ctx->ac, args, instr);
if (components_32bit > 1)
data_type = LLVMVectorType(ctx->ac.f32, components_32bit);
- base_data = to_float(&ctx->ac, src_data);
+ base_data = ac_to_float(&ctx->ac, src_data);
base_data = trim_vector(&ctx->ac, base_data, instr->num_components);
base_data = LLVMBuildBitCast(ctx->ac.builder, base_data,
data_type, "");
LLVMValueRef dw_addr)
{
LLVMValueRef value;
- value = ac_build_indexed_load(&ctx->ac, ctx->lds, dw_addr, false);
+ value = ac_build_load(&ctx->ac, ctx->lds, dw_addr);
return value;
}
load_tcs_output(struct nir_to_llvm_context *ctx,
nir_intrinsic_instr *instr)
{
- LLVMValueRef dw_addr, stride;
+ LLVMValueRef dw_addr;
+ LLVMValueRef stride = NULL;
LLVMValueRef value[4], result;
LLVMValueRef vertex_index = NULL;
LLVMValueRef indir_index = NULL;
LLVMValueRef src,
unsigned writemask)
{
- LLVMValueRef stride, dw_addr;
+ LLVMValueRef dw_addr;
+ LLVMValueRef stride = NULL;
LLVMValueRef buf_addr = NULL;
LLVMValueRef vertex_index = NULL;
LLVMValueRef indir_index = NULL;
{
LLVMValueRef temp_ptr, value;
int idx = instr->variables[0]->var->data.driver_location;
- LLVMValueRef src = to_float(&ctx->ac, get_src(ctx, instr->src[0]));
+ LLVMValueRef src = ac_to_float(&ctx->ac, get_src(ctx, instr->src[0]));
int writemask = instr->const_index[0];
LLVMValueRef indir_index;
unsigned const_index;
res = ac_build_image_opcode(ctx, &args);
- res = to_integer(ctx, res);
+ res = ac_to_integer(ctx, res);
LLVMValueRef four = LLVMConstInt(ctx->i32, 4, false);
LLVMValueRef F = LLVMConstInt(ctx->i32, 0xf, false);
int count;
enum glsl_sampler_dim dim = glsl_get_sampler_dim(type);
+ bool is_array = glsl_sampler_type_is_array(type);
bool add_frag_pos = (dim == GLSL_SAMPLER_DIM_SUBPASS ||
dim == GLSL_SAMPLER_DIM_SUBPASS_MS);
bool is_ms = (dim == GLSL_SAMPLER_DIM_MS ||
dim == GLSL_SAMPLER_DIM_SUBPASS_MS);
-
- count = image_type_to_components_count(dim,
- glsl_sampler_type_is_array(type));
+ bool gfx9_1d = ctx->ac.chip_class >= GFX9 && dim == GLSL_SAMPLER_DIM_1D;
+ count = image_type_to_components_count(dim, is_array);
if (is_ms) {
LLVMValueRef fmask_load_address[3];
fmask_load_address[0] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[0], "");
fmask_load_address[1] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[1], "");
- if (glsl_sampler_type_is_array(type))
+ if (is_array)
fmask_load_address[2] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[2], "");
else
fmask_load_address[2] = NULL;
LLVMBuildAdd(ctx->ac.builder, fmask_load_address[chan],
LLVMBuildFPToUI(ctx->ac.builder, ctx->abi->frag_pos[chan],
ctx->ac.i32, ""), "");
+ fmask_load_address[2] = ac_to_integer(&ctx->ac, ctx->abi->inputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)]);
}
sample_index = adjust_sample_index_using_fmask(&ctx->ac,
fmask_load_address[0],
sample_index,
get_sampler_desc(ctx, instr->variables[0], AC_DESC_FMASK, true, false));
}
- if (count == 1) {
+ if (count == 1 && !gfx9_1d) {
if (instr->src[0].ssa->num_components)
res = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[0], "");
else
if (is_ms)
count--;
for (chan = 0; chan < count; ++chan) {
- coords[chan] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[chan], "");
+ coords[chan] = llvm_extract_elem(&ctx->ac, src0, chan);
}
-
if (add_frag_pos) {
- for (chan = 0; chan < count; ++chan)
+ for (chan = 0; chan < 2; ++chan)
coords[chan] = LLVMBuildAdd(ctx->ac.builder, coords[chan], LLVMBuildFPToUI(ctx->ac.builder, ctx->abi->frag_pos[chan],
ctx->ac.i32, ""), "");
+ coords[2] = ac_to_integer(&ctx->ac, ctx->abi->inputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)]);
+ count++;
}
+
+ if (gfx9_1d) {
+ if (is_array) {
+ coords[2] = coords[1];
+ coords[1] = ctx->ac.i32_0;
+ } else
+ coords[1] = ctx->ac.i32_0;
+ count++;
+ }
+
if (is_ms) {
coords[count] = sample_index;
count++;
params, 5, 0);
res = trim_vector(&ctx->ac, res, instr->dest.ssa.num_components);
- res = to_integer(&ctx->ac, res);
+ res = ac_to_integer(&ctx->ac, res);
} else {
bool is_da = glsl_sampler_type_is_array(type) ||
- glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_CUBE;
+ glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_CUBE ||
+ glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_SUBPASS ||
+ glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_SUBPASS_MS;
LLVMValueRef da = is_da ? i1true : i1false;
LLVMValueRef glc = i1false;
LLVMValueRef slc = i1false;
res = ac_build_intrinsic(&ctx->ac, intrinsic_name, ctx->ac.v4f32,
params, 7, AC_FUNC_ATTR_READONLY);
}
- return to_integer(&ctx->ac, res);
+ return ac_to_integer(&ctx->ac, res);
}
static void visit_image_store(struct ac_nir_context *ctx,
LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
LLVMValueRef i1true = LLVMConstInt(ctx->ac.i1, 1, false);
LLVMValueRef glc = i1false;
- bool force_glc = ctx->abi->chip_class == SI;
+ bool force_glc = ctx->ac.chip_class == SI;
if (force_glc)
glc = i1true;
if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) {
- params[0] = to_float(&ctx->ac, get_src(ctx, instr->src[2])); /* data */
+ params[0] = ac_to_float(&ctx->ac, get_src(ctx, instr->src[2])); /* data */
params[1] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_BUFFER, true, true);
params[2] = LLVMBuildExtractElement(ctx->ac.builder, get_src(ctx, instr->src[0]),
ctx->ac.i32_0, ""); /* vindex */
LLVMValueRef da = is_da ? i1true : i1false;
LLVMValueRef slc = i1false;
- params[0] = to_float(&ctx->ac, get_src(ctx, instr->src[2]));
+ params[0] = ac_to_float(&ctx->ac, get_src(ctx, instr->src[2]));
params[1] = get_image_coords(ctx, instr); /* coords */
params[2] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, true);
params[3] = LLVMConstInt(ctx->ac.i32, 15, false); /* dmask */
static LLVMValueRef visit_image_atomic(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
- LLVMValueRef params[6];
+ LLVMValueRef params[7];
int param_count = 0;
const nir_variable *var = instr->variables[0]->var;
res = ac_build_image_opcode(&ctx->ac, &args);
+ LLVMValueRef two = LLVMConstInt(ctx->ac.i32, 2, false);
+
if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_CUBE &&
glsl_sampler_type_is_array(type)) {
- LLVMValueRef two = LLVMConstInt(ctx->ac.i32, 2, false);
LLVMValueRef six = LLVMConstInt(ctx->ac.i32, 6, false);
LLVMValueRef z = LLVMBuildExtractElement(ctx->ac.builder, res, two, "");
z = LLVMBuildSDiv(ctx->ac.builder, z, six, "");
res = LLVMBuildInsertElement(ctx->ac.builder, res, z, two, "");
}
+ if (ctx->ac.chip_class >= GFX9 &&
+ glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_1D &&
+ glsl_sampler_type_is_array(type)) {
+ LLVMValueRef layers = LLVMBuildExtractElement(ctx->ac.builder, res, two, "");
+ res = LLVMBuildInsertElement(ctx->ac.builder, res, layers,
+ ctx->ac.i32_1, "");
+
+ }
return res;
}
return NULL;
}
- result = LLVMBuildAtomicRMW(ctx->builder, op, ptr, to_integer(&ctx->ac, src),
+ result = LLVMBuildAtomicRMW(ctx->builder, op, ptr, ac_to_integer(&ctx->ac, src),
LLVMAtomicOrderingSequentiallyConsistent,
false);
}
const_array(ctx->v2f32, 64), "");
sample_id = LLVMBuildAdd(ctx->builder, sample_id, ctx->sample_pos_offset, "");
- result = ac_build_indexed_load(&ctx->ac, ptr, sample_id, false);
+ result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
return result;
}
LLVMValueRef interp_param, attr_number;
unsigned location;
unsigned chan;
- LLVMValueRef src_c0, src_c1;
- LLVMValueRef src0;
+ LLVMValueRef src_c0 = NULL;
+ LLVMValueRef src_c1 = NULL;
+ LLVMValueRef src0 = NULL;
int input_index = instr->variables[0]->var->data.location - VARYING_SLOT_VAR0;
switch (instr->intrinsic) {
case nir_intrinsic_interp_var_at_centroid:
}
if (instr->intrinsic == nir_intrinsic_interp_var_at_offset) {
- src_c0 = to_float(&ctx->ac, LLVMBuildExtractElement(ctx->builder, src0, ctx->i32zero, ""));
- src_c1 = to_float(&ctx->ac, LLVMBuildExtractElement(ctx->builder, src0, ctx->i32one, ""));
+ src_c0 = ac_to_float(&ctx->ac, LLVMBuildExtractElement(ctx->builder, src0, ctx->i32zero, ""));
+ src_c1 = ac_to_float(&ctx->ac, LLVMBuildExtractElement(ctx->builder, src0, ctx->i32one, ""));
} else if (instr->intrinsic == nir_intrinsic_interp_var_at_sample) {
LLVMValueRef sample_position;
LLVMValueRef halfval = LLVMConstReal(ctx->f32, 0.5f);
interp_param = lookup_interp_param(ctx, instr->variables[0]->var->data.interpolation, location);
attr_number = LLVMConstInt(ctx->i32, input_index, false);
- if (location == INTERP_SAMPLE || location == INTERP_CENTER) {
+ if (location == INTERP_CENTER) {
LLVMValueRef ij_out[2];
LLVMValueRef ddxy_out = emit_ddxy_interp(ctx->nir, interp_param);
case nir_intrinsic_load_draw_id:
result = ctx->abi->draw_id;
break;
+ case nir_intrinsic_load_view_index:
+ result = ctx->nctx->view_index ? ctx->nctx->view_index : ctx->ac.i32_0;
+ break;
case nir_intrinsic_load_invocation_id:
if (ctx->stage == MESA_SHADER_TESS_CTRL)
result = unpack_param(&ctx->ac, ctx->nctx->tcs_rel_ids, 8, 5);
list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->i32, offset, 0));
list = LLVMBuildPointerCast(builder, list, const_array(type, 0), "");
- return ac_build_indexed_load_const(&ctx->ac, list, index);
+ return ac_build_load_to_sgpr(&ctx->ac, list, index);
}
static LLVMValueRef get_sampler_desc(struct ac_nir_context *ctx,
LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef img7, samp0;
- if (ctx->abi->chip_class >= VI)
+ if (ctx->ac.chip_class >= VI)
return samp;
img7 = LLVMBuildExtractElement(builder, res,
static LLVMValueRef apply_round_slice(struct ac_llvm_context *ctx,
LLVMValueRef coord)
{
- coord = to_float(ctx, coord);
+ coord = ac_to_float(ctx, coord);
coord = ac_build_intrinsic(ctx, "llvm.rint.f32", ctx->f32, &coord, 1, 0);
- coord = to_integer(ctx, coord);
+ coord = ac_to_integer(ctx, coord);
return coord;
}
/* Pack depth comparison value */
if (instr->is_shadow && comparator) {
- LLVMValueRef z = to_float(&ctx->ac,
- llvm_extract_elem(&ctx->ac, comparator, 0));
+ LLVMValueRef z = ac_to_float(&ctx->ac,
+ llvm_extract_elem(&ctx->ac, comparator, 0));
/* TC-compatible HTILE promotes Z16 and Z24 to Z32_FLOAT,
* so the depth comparison value isn't clamped for Z16 and
* It's unnecessary if the original texture format was
* Z32_FLOAT, but we don't know that here.
*/
- if (ctx->abi->chip_class == VI)
+ if (ctx->ac.chip_class == VI)
z = ac_build_clamp(&ctx->ac, z);
address[count++] = z;
/* pack derivatives */
if (ddx || ddy) {
+ int num_src_deriv_channels, num_dest_deriv_channels;
switch (instr->sampler_dim) {
case GLSL_SAMPLER_DIM_3D:
case GLSL_SAMPLER_DIM_CUBE:
num_deriv_comp = 3;
+ num_src_deriv_channels = 3;
+ num_dest_deriv_channels = 3;
break;
case GLSL_SAMPLER_DIM_2D:
default:
+ num_src_deriv_channels = 2;
+ num_dest_deriv_channels = 2;
num_deriv_comp = 2;
break;
case GLSL_SAMPLER_DIM_1D:
- num_deriv_comp = 1;
+ num_src_deriv_channels = 1;
+ if (ctx->ac.chip_class >= GFX9) {
+ num_dest_deriv_channels = 2;
+ num_deriv_comp = 2;
+ } else {
+ num_dest_deriv_channels = 1;
+ num_deriv_comp = 1;
+ }
break;
}
- for (unsigned i = 0; i < num_deriv_comp; i++) {
- derivs[i] = to_float(&ctx->ac, llvm_extract_elem(&ctx->ac, ddx, i));
- derivs[num_deriv_comp + i] = to_float(&ctx->ac, llvm_extract_elem(&ctx->ac, ddy, i));
+ for (unsigned i = 0; i < num_src_deriv_channels; i++) {
+ derivs[i] = ac_to_float(&ctx->ac, llvm_extract_elem(&ctx->ac, ddx, i));
+ derivs[num_dest_deriv_channels + i] = ac_to_float(&ctx->ac, llvm_extract_elem(&ctx->ac, ddy, i));
+ }
+ for (unsigned i = num_src_deriv_channels; i < num_dest_deriv_channels; i++) {
+ derivs[i] = ctx->ac.f32_0;
+ derivs[num_dest_deriv_channels + i] = ctx->ac.f32_0;
}
}
if (instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE && coord) {
- if (instr->is_array && instr->op != nir_texop_lod)
- coords[3] = apply_round_slice(&ctx->ac, coords[3]);
for (chan = 0; chan < instr->coord_components; chan++)
- coords[chan] = to_float(&ctx->ac, coords[chan]);
+ coords[chan] = ac_to_float(&ctx->ac, coords[chan]);
if (instr->coord_components == 3)
coords[3] = LLVMGetUndef(ctx->ac.f32);
ac_prepare_cube_coords(&ctx->ac,
instr->op == nir_texop_txd, instr->is_array,
- coords, derivs);
+ instr->op == nir_texop_lod, coords, derivs);
if (num_deriv_comp)
num_deriv_comp--;
}
}
address[count++] = coords[2];
}
+
+ if (ctx->ac.chip_class >= GFX9) {
+ LLVMValueRef filler;
+ if (instr->op == nir_texop_txf)
+ filler = ctx->ac.i32_0;
+ else
+ filler = LLVMConstReal(ctx->ac.f32, 0.5);
+
+ if (instr->sampler_dim == GLSL_SAMPLER_DIM_1D) {
+ if (instr->is_array) {
+ address[count] = address[count - 1];
+ address[count - 1] = filler;
+ count++;
+ } else
+ address[count++] = filler;
+ }
+ }
}
/* Pack LOD */
LLVMValueRef z = LLVMBuildExtractElement(ctx->ac.builder, result, two, "");
z = LLVMBuildSDiv(ctx->ac.builder, z, six, "");
result = LLVMBuildInsertElement(ctx->ac.builder, result, z, two, "");
+ } else if (ctx->ac.chip_class >= GFX9 &&
+ instr->op == nir_texop_txs &&
+ instr->sampler_dim == GLSL_SAMPLER_DIM_1D &&
+ instr->is_array) {
+ LLVMValueRef two = LLVMConstInt(ctx->ac.i32, 2, false);
+ LLVMValueRef layers = LLVMBuildExtractElement(ctx->ac.builder, result, two, "");
+ result = LLVMBuildInsertElement(ctx->ac.builder, result, layers,
+ ctx->ac.i32_1, "");
} else if (instr->dest.ssa.num_components != 4)
result = trim_vector(&ctx->ac, result, instr->dest.ssa.num_components);
write_result:
if (result) {
assert(instr->dest.is_ssa);
- result = to_integer(&ctx->ac, result);
+ result = ac_to_integer(&ctx->ac, result);
_mesa_hash_table_insert(ctx->defs, &instr->dest.ssa, result);
}
}
for (unsigned i = 0; i < attrib_count; ++i, ++idx) {
t_offset = LLVMConstInt(ctx->i32, index + i, false);
- 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);
input = ac_build_buffer_load_format(&ctx->ac, t_list,
buffer_index,
for (unsigned chan = 0; chan < 4; chan++) {
LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, false);
ctx->inputs[radeon_llvm_reg_index_soa(idx, chan)] =
- to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->builder,
+ ac_to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->builder,
input, llvm_chan, ""));
}
}
}
static void
-handle_shader_input_decl(struct nir_to_llvm_context *ctx,
- struct nir_variable *variable)
-{
- switch (ctx->stage) {
- case MESA_SHADER_VERTEX:
+handle_vs_inputs(struct nir_to_llvm_context *ctx,
+ struct nir_shader *nir) {
+ nir_foreach_variable(variable, &nir->inputs)
handle_vs_input_decl(ctx, variable);
- break;
- case MESA_SHADER_FRAGMENT:
- handle_fs_input_decl(ctx, variable);
- break;
- default:
- break;
+}
+
+static void
+prepare_interp_optimize(struct nir_to_llvm_context *ctx,
+ struct nir_shader *nir)
+{
+ if (!ctx->options->key.fs.multisample)
+ return;
+
+ bool uses_center = false;
+ bool uses_centroid = false;
+ nir_foreach_variable(variable, &nir->inputs) {
+ if (glsl_get_base_type(glsl_without_array(variable->type)) != GLSL_TYPE_FLOAT ||
+ variable->data.sample)
+ continue;
+
+ if (variable->data.centroid)
+ uses_centroid = true;
+ else
+ uses_center = true;
}
+ if (uses_center && uses_centroid) {
+ LLVMValueRef sel = LLVMBuildICmp(ctx->builder, LLVMIntSLT, ctx->prim_mask, ctx->ac.i32_0, "");
+ ctx->persp_centroid = LLVMBuildSelect(ctx->builder, sel, ctx->persp_center, ctx->persp_centroid, "");
+ ctx->linear_centroid = LLVMBuildSelect(ctx->builder, sel, ctx->linear_center, ctx->linear_centroid, "");
+ }
}
static void
-handle_fs_inputs_pre(struct nir_to_llvm_context *ctx,
- struct nir_shader *nir)
+handle_fs_inputs(struct nir_to_llvm_context *ctx,
+ struct nir_shader *nir)
{
+ prepare_interp_optimize(ctx, nir);
+
+ nir_foreach_variable(variable, &nir->inputs)
+ handle_fs_input_decl(ctx, variable);
+
unsigned index = 0;
+
+ if (ctx->shader_info->info.ps.uses_input_attachments ||
+ ctx->shader_info->info.needs_multiview_view_index)
+ ctx->input_mask |= 1ull << VARYING_SLOT_LAYER;
+
for (unsigned i = 0; i < RADEON_LLVM_MAX_INPUTS; ++i) {
LLVMValueRef interp_param;
LLVMValueRef *inputs = ctx->inputs +radeon_llvm_reg_index_soa(i, 0);
if (ctx->input_mask & (1 << VARYING_SLOT_LAYER))
ctx->shader_info->fs.layer_input = true;
ctx->shader_info->fs.input_mask = ctx->input_mask >> VARYING_SLOT_VAR0;
+
+ if (ctx->shader_info->info.needs_multiview_view_index)
+ ctx->view_index = ctx->inputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
}
static LLVMValueRef
static void
scan_shader_output_decl(struct nir_to_llvm_context *ctx,
- struct nir_variable *variable)
+ struct nir_variable *variable,
+ struct nir_shader *shader,
+ gl_shader_stage stage)
{
int idx = variable->data.location + variable->data.index;
unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
variable->data.driver_location = idx * 4;
/* tess ctrl has it's own load/store paths for outputs */
- if (ctx->stage == MESA_SHADER_TESS_CTRL)
+ if (stage == MESA_SHADER_TESS_CTRL)
return;
mask_attribs = ((1ull << attrib_count) - 1) << idx;
- if (ctx->stage == MESA_SHADER_VERTEX ||
- ctx->stage == MESA_SHADER_TESS_EVAL ||
- ctx->stage == MESA_SHADER_GEOMETRY) {
+ if (stage == MESA_SHADER_VERTEX ||
+ stage == MESA_SHADER_TESS_EVAL ||
+ stage == MESA_SHADER_GEOMETRY) {
if (idx == VARYING_SLOT_CLIP_DIST0) {
- int length = ctx->num_output_clips + ctx->num_output_culls;
- if (ctx->stage == MESA_SHADER_VERTEX) {
- ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << ctx->num_output_clips) - 1;
- ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << ctx->num_output_culls) - 1;
+ int length = shader->info.clip_distance_array_size +
+ shader->info.cull_distance_array_size;
+ if (stage == MESA_SHADER_VERTEX) {
+ ctx->shader_info->vs.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
+ ctx->shader_info->vs.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
}
- if (ctx->stage == MESA_SHADER_TESS_EVAL) {
- ctx->shader_info->tes.outinfo.clip_dist_mask = (1 << ctx->num_output_clips) - 1;
- ctx->shader_info->tes.outinfo.cull_dist_mask = (1 << ctx->num_output_culls) - 1;
+ if (stage == MESA_SHADER_TESS_EVAL) {
+ ctx->shader_info->tes.outinfo.clip_dist_mask = (1 << shader->info.clip_distance_array_size) - 1;
+ ctx->shader_info->tes.outinfo.cull_dist_mask = (1 << shader->info.cull_distance_array_size) - 1;
}
if (length > 4)
static LLVMValueRef
emit_float_saturate(struct ac_llvm_context *ctx, LLVMValueRef v, float lo, float hi)
{
- v = to_float(ctx, v);
- v = emit_intrin_2f_param(ctx, "llvm.maxnum.f32", ctx->f32, v, LLVMConstReal(ctx->f32, lo));
- return emit_intrin_2f_param(ctx, "llvm.minnum.f32", ctx->f32, v, LLVMConstReal(ctx->f32, hi));
+ v = ac_to_float(ctx, v);
+ v = emit_intrin_2f_param(ctx, "llvm.maxnum", ctx->f32, v, LLVMConstReal(ctx->f32, lo));
+ return emit_intrin_2f_param(ctx, "llvm.minnum", ctx->f32, v, LLVMConstReal(ctx->f32, hi));
}
LLVMValueRef max_alpha = !is_int10 ? max_rgb : LLVMConstInt(ctx->i32, 3, 0);
for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = to_integer(&ctx->ac, values[chan]);
+ val[chan] = ac_to_integer(&ctx->ac, values[chan]);
val[chan] = emit_minmax_int(&ctx->ac, LLVMIntULT, val[chan], chan == 3 ? max_alpha : max_rgb);
}
/* Clamp. */
for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = to_integer(&ctx->ac, values[chan]);
+ val[chan] = ac_to_integer(&ctx->ac, values[chan]);
val[chan] = emit_minmax_int(&ctx->ac, LLVMIntSLT, val[chan], chan == 3 ? max_alpha : max_rgb);
val[chan] = emit_minmax_int(&ctx->ac, LLVMIntSGT, val[chan], chan == 3 ? min_alpha : min_rgb);
}
memcpy(&args->out[0], values, sizeof(values[0]) * 4);
for (unsigned i = 0; i < 4; ++i)
- args->out[i] = to_float(&ctx->ac, args->out[i]);
+ args->out[i] = ac_to_float(&ctx->ac, args->out[i]);
}
static void
LLVMValueRef psize_value = NULL, layer_value = NULL, viewport_index_value = NULL;
int i;
+ if (ctx->options->key.has_multiview_view_index) {
+ LLVMValueRef* tmp_out = &ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
+ if(!*tmp_out) {
+ for(unsigned i = 0; i < 4; ++i)
+ ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, i)] =
+ si_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
+ }
+
+ LLVMBuildStore(ctx->builder, ac_to_float(&ctx->ac, ctx->view_index), *tmp_out);
+ ctx->output_mask |= 1ull << VARYING_SLOT_LAYER;
+ }
+
memset(outinfo->vs_output_param_offset, AC_EXP_PARAM_UNDEFINED,
sizeof(outinfo->vs_output_param_offset));
i = VARYING_SLOT_CLIP_DIST0;
for (j = 0; j < ctx->num_output_clips + ctx->num_output_culls; j++)
- slots[j] = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
+ slots[j] = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
for (i = ctx->num_output_clips + ctx->num_output_culls; i < 8; i++)
ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_VIEWPORT, 0)], "");
}
- uint32_t mask = ((outinfo->writes_pointsize == true ? 1 : 0) |
- (outinfo->writes_layer == true ? 4 : 0) |
- (outinfo->writes_viewport_index == true ? 8 : 0));
- if (mask) {
- pos_args[1].enabled_channels = mask;
+ if (outinfo->writes_pointsize ||
+ outinfo->writes_layer ||
+ outinfo->writes_viewport_index) {
+ pos_args[1].enabled_channels = ((outinfo->writes_pointsize == true ? 1 : 0) |
+ (outinfo->writes_layer == true ? 4 : 0));
pos_args[1].valid_mask = 0;
pos_args[1].done = 0;
pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
pos_args[1].out[0] = psize_value;
if (outinfo->writes_layer == true)
pos_args[1].out[2] = layer_value;
- if (outinfo->writes_viewport_index == true)
- pos_args[1].out[3] = viewport_index_value;
+ if (outinfo->writes_viewport_index == true) {
+ if (ctx->options->chip_class >= GFX9) {
+ /* GFX9 has the layer in out.z[10:0] and the viewport
+ * index in out.z[19:16].
+ */
+ LLVMValueRef v = viewport_index_value;
+ v = ac_to_integer(&ctx->ac, v);
+ v = LLVMBuildShl(ctx->builder, v,
+ LLVMConstInt(ctx->i32, 16, false),
+ "");
+ v = LLVMBuildOr(ctx->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 {
+ pos_args[1].out[3] = viewport_index_value;
+ pos_args[1].enabled_channels |= 1 << 3;
+ }
+ }
}
for (i = 0; i < 4; i++) {
if (pos_args[i].out[0])
continue;
for (unsigned j = 0; j < 4; j++)
- values[j] = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
- ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
+ values[j] = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
+ ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
if (i == VARYING_SLOT_LAYER) {
target = V_008DFC_SQ_EXP_PARAM + param_count;
if (i == FRAG_RESULT_DEPTH) {
ctx->shader_info->fs.writes_z = true;
- depth = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
+ depth = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
ctx->nir->outputs[radeon_llvm_reg_index_soa(i, 0)], ""));
} else if (i == FRAG_RESULT_STENCIL) {
ctx->shader_info->fs.writes_stencil = true;
- stencil = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
+ stencil = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
ctx->nir->outputs[radeon_llvm_reg_index_soa(i, 0)], ""));
} else if (i == FRAG_RESULT_SAMPLE_MASK) {
ctx->shader_info->fs.writes_sample_mask = true;
- samplemask = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
+ samplemask = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
ctx->nir->outputs[radeon_llvm_reg_index_soa(i, 0)], ""));
} else {
bool last = false;
for (unsigned j = 0; j < 4; j++)
- values[j] = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
+ values[j] = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
if (!ctx->shader_info->fs.writes_z && !ctx->shader_info->fs.writes_stencil && !ctx->shader_info->fs.writes_sample_mask)
{
if ((ctx->stage == MESA_SHADER_VERTEX && ctx->options->key.vs.as_es) ||
(ctx->stage == MESA_SHADER_TESS_EVAL && ctx->options->key.tes.as_es)) {
- ctx->esgs_ring = ac_build_indexed_load_const(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_ESGS_VS, false));
+ ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_ESGS_VS, false));
}
if (ctx->is_gs_copy_shader) {
- ctx->gsvs_ring = ac_build_indexed_load_const(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_GSVS_VS, false));
+ ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_GSVS_VS, false));
}
if (ctx->stage == MESA_SHADER_GEOMETRY) {
LLVMValueRef tmp;
- ctx->esgs_ring = ac_build_indexed_load_const(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_ESGS_GS, false));
- ctx->gsvs_ring = ac_build_indexed_load_const(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_GSVS_GS, false));
+ ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_ESGS_GS, false));
+ ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_GSVS_GS, false));
ctx->gsvs_ring = LLVMBuildBitCast(ctx->builder, ctx->gsvs_ring, ctx->v4i32, "");
if (ctx->stage == MESA_SHADER_TESS_CTRL ||
ctx->stage == MESA_SHADER_TESS_EVAL) {
- ctx->hs_ring_tess_offchip = ac_build_indexed_load_const(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_OFFCHIP, false));
- ctx->hs_ring_tess_factor = ac_build_indexed_load_const(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_FACTOR, false));
+ ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_OFFCHIP, false));
+ ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_FACTOR, false));
}
}
ctx.context = LLVMContextCreate();
ctx.module = LLVMModuleCreateWithNameInContext("shader", ctx.context);
- ac_llvm_context_init(&ctx.ac, ctx.context);
+ ac_llvm_context_init(&ctx.ac, ctx.context, options->chip_class);
ctx.ac.module = ctx.module;
memset(shader_info, 0, sizeof(*shader_info));
for (i = 0; i < AC_UD_MAX_UD; i++)
shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
- create_function(&ctx);
+ create_function(&ctx, nir->stage, false, MESA_SHADER_VERTEX);
if (nir->stage == MESA_SHADER_GEOMETRY) {
ctx.gs_next_vertex = ac_build_alloca(&ctx.ac, ctx.i32, "gs_next_vertex");
ctx.num_output_clips = nir->info.clip_distance_array_size;
ctx.num_output_culls = nir->info.cull_distance_array_size;
- nir_foreach_variable(variable, &nir->inputs)
- handle_shader_input_decl(&ctx, variable);
-
if (nir->stage == MESA_SHADER_FRAGMENT)
- handle_fs_inputs_pre(&ctx, nir);
+ handle_fs_inputs(&ctx, nir);
+ else if(nir->stage == MESA_SHADER_VERTEX)
+ handle_vs_inputs(&ctx, nir);
- ctx.abi.chip_class = options->chip_class;
ctx.abi.inputs = &ctx.inputs[0];
ctx.abi.emit_outputs = handle_shader_outputs_post;
ctx.abi.load_ssbo = radv_load_ssbo;
ctx.abi.load_sampler_desc = radv_get_sampler_desc;
nir_foreach_variable(variable, &nir->outputs)
- scan_shader_output_decl(&ctx, variable);
+ scan_shader_output_decl(&ctx, variable, nir, nir->stage);
ac_nir_translate(&ctx.ac, &ctx.abi, nir, &ctx);
shader_info->num_input_sgprs + 3);
}
+static void
+ac_fill_shader_info(struct ac_shader_variant_info *shader_info, struct nir_shader *nir, const struct ac_nir_compiler_options *options)
+{
+ switch (nir->stage) {
+ case MESA_SHADER_COMPUTE:
+ for (int i = 0; i < 3; ++i)
+ shader_info->cs.block_size[i] = nir->info.cs.local_size[i];
+ break;
+ case MESA_SHADER_FRAGMENT:
+ shader_info->fs.early_fragment_test = nir->info.fs.early_fragment_tests;
+ break;
+ case MESA_SHADER_GEOMETRY:
+ shader_info->gs.vertices_in = nir->info.gs.vertices_in;
+ shader_info->gs.vertices_out = nir->info.gs.vertices_out;
+ shader_info->gs.output_prim = nir->info.gs.output_primitive;
+ shader_info->gs.invocations = nir->info.gs.invocations;
+ break;
+ case MESA_SHADER_TESS_EVAL:
+ shader_info->tes.primitive_mode = nir->info.tess.primitive_mode;
+ shader_info->tes.spacing = nir->info.tess.spacing;
+ shader_info->tes.ccw = nir->info.tess.ccw;
+ shader_info->tes.point_mode = nir->info.tess.point_mode;
+ shader_info->tes.as_es = options->key.tes.as_es;
+ break;
+ case MESA_SHADER_TESS_CTRL:
+ shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
+ break;
+ case MESA_SHADER_VERTEX:
+ shader_info->vs.as_es = options->key.vs.as_es;
+ shader_info->vs.as_ls = options->key.vs.as_ls;
+ /* in LS mode we need at least 1, invocation id needs 3, handled elsewhere */
+ if (options->key.vs.as_ls)
+ shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt);
+ break;
+ default:
+ break;
+ }
+}
+
void ac_compile_nir_shader(LLVMTargetMachineRef tm,
struct ac_shader_binary *binary,
struct ac_shader_config *config,
struct ac_shader_variant_info *shader_info,
- struct nir_shader *nir,
+ struct nir_shader *const *nir,
+ int nir_count,
const struct ac_nir_compiler_options *options,
bool dump_shader)
{
- LLVMModuleRef llvm_module = ac_translate_nir_to_llvm(tm, nir, shader_info,
+ LLVMModuleRef llvm_module = ac_translate_nir_to_llvm(tm, nir[0], shader_info,
options);
- ac_compile_llvm_module(tm, llvm_module, binary, config, shader_info, nir->stage, dump_shader, options->supports_spill);
- switch (nir->stage) {
- case MESA_SHADER_COMPUTE:
- for (int i = 0; i < 3; ++i)
- shader_info->cs.block_size[i] = nir->info.cs.local_size[i];
- break;
- case MESA_SHADER_FRAGMENT:
- shader_info->fs.early_fragment_test = nir->info.fs.early_fragment_tests;
- break;
- case MESA_SHADER_GEOMETRY:
- shader_info->gs.vertices_in = nir->info.gs.vertices_in;
- shader_info->gs.vertices_out = nir->info.gs.vertices_out;
- shader_info->gs.output_prim = nir->info.gs.output_primitive;
- shader_info->gs.invocations = nir->info.gs.invocations;
- break;
- case MESA_SHADER_TESS_EVAL:
- shader_info->tes.primitive_mode = nir->info.tess.primitive_mode;
- shader_info->tes.spacing = nir->info.tess.spacing;
- shader_info->tes.ccw = nir->info.tess.ccw;
- shader_info->tes.point_mode = nir->info.tess.point_mode;
- shader_info->tes.as_es = options->key.tes.as_es;
- break;
- case MESA_SHADER_TESS_CTRL:
- shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
- break;
- case MESA_SHADER_VERTEX:
- shader_info->vs.as_es = options->key.vs.as_es;
- shader_info->vs.as_ls = options->key.vs.as_ls;
- /* in LS mode we need at least 1, invocation id needs 3, handled elsewhere */
- if (options->key.vs.as_ls)
- shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt);
- break;
- default:
- break;
- }
+ ac_compile_llvm_module(tm, llvm_module, binary, config, shader_info, nir[0]->stage, dump_shader, options->supports_spill);
+ for (int i = 0; i < nir_count; ++i)
+ ac_fill_shader_info(shader_info, nir[i], options);
}
static void
AC_FUNC_ATTR_LEGACY);
LLVMBuildStore(ctx->builder,
- to_float(&ctx->ac, value), ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]);
+ ac_to_float(&ctx->ac, value), ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]);
}
idx += slot_inc;
}
ctx.options = options;
ctx.shader_info = shader_info;
- ac_llvm_context_init(&ctx.ac, ctx.context);
+ ac_llvm_context_init(&ctx.ac, ctx.context, options->chip_class);
ctx.ac.module = ctx.module;
ctx.is_gs_copy_shader = true;
ctx.ac.builder = ctx.builder;
ctx.stage = MESA_SHADER_VERTEX;
- create_function(&ctx);
+ create_function(&ctx, MESA_SHADER_VERTEX, false, MESA_SHADER_VERTEX);
ctx.gs_max_out_vertices = geom_shader->info.gs.vertices_out;
ac_setup_rings(&ctx);
ctx.nir = &nir_ctx;
nir_foreach_variable(variable, &geom_shader->outputs) {
- scan_shader_output_decl(&ctx, variable);
+ scan_shader_output_decl(&ctx, variable, geom_shader, MESA_SHADER_VERTEX);
handle_shader_output_decl(&nir_ctx, geom_shader, variable);
}
projects / mesa.git / blobdiff