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
#define RADEON_LLVM_MAX_INPUTS (VARYING_SLOT_VAR31 + 1)
#define RADEON_LLVM_MAX_OUTPUTS (VARYING_SLOT_VAR31 + 1)
-enum desc_type {
- DESC_IMAGE,
- DESC_FMASK,
- DESC_SAMPLER,
- DESC_BUFFER,
+struct nir_to_llvm_context;
+
+struct ac_nir_context {
+ struct ac_llvm_context ac;
+ struct ac_shader_abi *abi;
+
+ gl_shader_stage stage;
+
+ struct hash_table *defs;
+ struct hash_table *phis;
+ struct hash_table *vars;
+
+ LLVMValueRef main_function;
+ LLVMBasicBlockRef continue_block;
+ LLVMBasicBlockRef break_block;
+
+ LLVMValueRef outputs[RADEON_LLVM_MAX_OUTPUTS * 4];
+
+ int num_locals;
+ LLVMValueRef *locals;
+
+ struct nir_to_llvm_context *nctx; /* TODO get rid of this */
};
struct nir_to_llvm_context {
const struct ac_nir_compiler_options *options;
struct ac_shader_variant_info *shader_info;
struct ac_shader_abi abi;
+ struct ac_nir_context *nir;
unsigned max_workgroup_size;
LLVMContextRef context;
struct hash_table *defs;
struct hash_table *phis;
- struct hash_table *vars;
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;
LLVMValueRef sample_pos_offset;
LLVMValueRef persp_sample, persp_center, persp_centroid;
LLVMValueRef linear_sample, linear_center, linear_centroid;
- LLVMValueRef front_face;
- LLVMValueRef ancillary;
- LLVMValueRef sample_coverage;
- LLVMValueRef frag_pos[4];
-
- LLVMBasicBlockRef continue_block;
- LLVMBasicBlockRef break_block;
LLVMTypeRef i1;
LLVMTypeRef i8;
LLVMValueRef lds;
LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4];
- LLVMValueRef outputs[RADEON_LLVM_MAX_OUTPUTS * 4];
uint64_t input_mask;
uint64_t output_mask;
- int num_locals;
- LLVMValueRef *locals;
uint8_t num_output_clips;
uint8_t num_output_culls;
- bool has_ds_bpermute;
-
bool is_gs_copy_shader;
LLVMValueRef gs_next_vertex;
unsigned gs_max_out_vertices;
uint64_t tess_patch_outputs_written;
};
-static LLVMValueRef get_sampler_desc(struct nir_to_llvm_context *ctx,
+static inline struct nir_to_llvm_context *
+nir_to_llvm_context_from_abi(struct ac_shader_abi *abi)
+{
+ struct nir_to_llvm_context *ctx = NULL;
+ return container_of(abi, ctx, abi);
+}
+
+static LLVMValueRef get_sampler_desc(struct ac_nir_context *ctx,
const nir_deref_var *deref,
- enum desc_type desc_type);
+ enum ac_descriptor_type desc_type,
+ bool image, bool write);
+
static unsigned radeon_llvm_reg_index_soa(unsigned index, unsigned chan)
{
return (index * 4) + chan;
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)
unreachable("Unhandled type kind in get_elem_bits");
}
-static LLVMValueRef unpack_param(struct nir_to_llvm_context *ctx,
+static LLVMValueRef unpack_param(struct ac_llvm_context *ctx,
LLVMValueRef param, unsigned rshift,
unsigned bitwidth)
{
{
switch (ctx->stage) {
case MESA_SHADER_TESS_CTRL:
- return unpack_param(ctx, ctx->tcs_rel_ids, 0, 8);
+ return unpack_param(&ctx->ac, ctx->tcs_rel_ids, 0, 8);
case MESA_SHADER_TESS_EVAL:
return ctx->tes_rel_patch_id;
break;
get_tcs_in_patch_stride(struct nir_to_llvm_context *ctx)
{
if (ctx->stage == MESA_SHADER_VERTEX)
- return unpack_param(ctx, ctx->ls_out_layout, 0, 13);
+ return unpack_param(&ctx->ac, ctx->ls_out_layout, 0, 13);
else if (ctx->stage == MESA_SHADER_TESS_CTRL)
- return unpack_param(ctx, ctx->tcs_in_layout, 0, 13);
+ return unpack_param(&ctx->ac, ctx->tcs_in_layout, 0, 13);
else {
assert(0);
return NULL;
static LLVMValueRef
get_tcs_out_patch_stride(struct nir_to_llvm_context *ctx)
{
- return unpack_param(ctx, ctx->tcs_out_layout, 0, 13);
+ return unpack_param(&ctx->ac, ctx->tcs_out_layout, 0, 13);
}
static LLVMValueRef
get_tcs_out_patch0_offset(struct nir_to_llvm_context *ctx)
{
return LLVMBuildMul(ctx->builder,
- unpack_param(ctx, ctx->tcs_out_offsets, 0, 16),
+ unpack_param(&ctx->ac, ctx->tcs_out_offsets, 0, 16),
LLVMConstInt(ctx->i32, 4, false), "");
}
get_tcs_out_patch0_patch_data_offset(struct nir_to_llvm_context *ctx)
{
return LLVMBuildMul(ctx->builder,
- unpack_param(ctx, ctx->tcs_out_offsets, 16, 16),
+ unpack_param(&ctx->ac, ctx->tcs_out_offsets, 16, 16),
LLVMConstInt(ctx->i32, 4, false), "");
}
}
}
-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:
- 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
- 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_vtx_offset[1]); // vtx1
- add_vgpr_argument(&args, ctx->i32, &ctx->gs_prim_id); // prim id
- add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[2]);
- add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[3]);
- add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[4]);
- add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[5]);
- add_vgpr_argument(&args, ctx->i32, &ctx->gs_invocation_id);
+ if (has_previous_stage) {
+ // First 6 system regs
+ add_sgpr_argument(&args, ctx->i32, &ctx->gs2vs_offset); // tess factor offset
+ add_sgpr_argument(&args, ctx->i32, &ctx->merged_wave_info); // merged wave info
+ add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // param oc lds
+
+ 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);
+ if (previous_stage == MESA_SHADER_TESS_EVAL)
+ add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_offchip_layout); // tcs offchip layout
+ else
+ 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_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[0]); // vtx01
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[2]); // vtx23
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_prim_id); // prim id
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_invocation_id);
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[4]);
+
+ if (previous_stage == MESA_SHADER_VERTEX) {
+ 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 {
+ add_vgpr_argument(&args, ctx->f32, &ctx->tes_u); // tes_u
+ add_vgpr_argument(&args, ctx->f32, &ctx->tes_v); // tes_v
+ add_vgpr_argument(&args, ctx->i32, &ctx->tes_rel_patch_id); // tes rel patch id
+ add_vgpr_argument(&args, ctx->i32, &ctx->tes_patch_id); // tes patch id
+ }
+ } else {
+ 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_vtx_offset[1]); // vtx1
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_prim_id); // prim id
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[2]);
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[3]);
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[4]);
+ add_vgpr_argument(&args, ctx->i32, &ctx->gs_vtx_offset[5]);
+ 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 */
add_vgpr_argument(&args, ctx->v2i32, &ctx->linear_center); /* linear center */
add_vgpr_argument(&args, ctx->v2i32, &ctx->linear_centroid); /* linear centroid */
add_vgpr_argument(&args, ctx->f32, NULL); /* line stipple tex */
- add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[0]); /* pos x float */
- add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[1]); /* pos y float */
- add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[2]); /* pos z float */
- add_vgpr_argument(&args, ctx->f32, &ctx->frag_pos[3]); /* pos w float */
- add_vgpr_argument(&args, ctx->i32, &ctx->front_face); /* front face */
- add_vgpr_argument(&args, ctx->i32, &ctx->ancillary); /* ancillary */
- add_vgpr_argument(&args, ctx->i32, &ctx->sample_coverage); /* sample coverage */
+ add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[0]); /* pos x float */
+ add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[1]); /* pos y float */
+ add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[2]); /* pos z float */
+ add_vgpr_argument(&args, ctx->f32, &ctx->abi.frag_pos[3]); /* pos w float */
+ add_vgpr_argument(&args, ctx->i32, &ctx->abi.front_face); /* front face */
+ add_vgpr_argument(&args, ctx->i32, &ctx->abi.ancillary); /* ancillary */
+ add_vgpr_argument(&args, ctx->i32, &ctx->abi.sample_coverage); /* sample coverage */
add_vgpr_argument(&args, ctx->i32, NULL); /* fixed pt */
break;
default:
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:
+ if (has_previous_stage) {
+ if (previous_stage == MESA_SHADER_VERTEX)
+ radv_define_vs_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_idx);
+ else
+ set_userdata_location_shader(ctx, AC_UD_TES_OFFCHIP_LAYOUT, &user_sgpr_idx, 1);
+ }
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);
+ if (has_previous_stage)
+ declare_tess_lds(ctx);
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)
return num_components;
}
-static LLVMValueRef llvm_extract_elem(struct nir_to_llvm_context *ctx,
+static LLVMValueRef llvm_extract_elem(struct ac_llvm_context *ac,
LLVMValueRef value,
int index)
{
if (count == 1)
return value;
- return LLVMBuildExtractElement(ctx->builder, value,
- LLVMConstInt(ctx->i32, index, false), "");
+ return LLVMBuildExtractElement(ac->builder, value,
+ LLVMConstInt(ac->i32, index, false), "");
}
-static LLVMValueRef trim_vector(struct nir_to_llvm_context *ctx,
+static LLVMValueRef trim_vector(struct ac_llvm_context *ctx,
LLVMValueRef value, unsigned count)
{
unsigned num_components = get_llvm_num_components(value);
}
static void
-build_store_values_extended(struct nir_to_llvm_context *ctx,
+build_store_values_extended(struct ac_llvm_context *ac,
LLVMValueRef *values,
unsigned value_count,
unsigned value_stride,
LLVMValueRef vec)
{
- LLVMBuilderRef builder = ctx->builder;
+ LLVMBuilderRef builder = ac->builder;
unsigned i;
- if (value_count == 1) {
- LLVMBuildStore(builder, vec, values[0]);
- return;
- }
-
for (i = 0; i < value_count; i++) {
LLVMValueRef ptr = values[i * value_stride];
- LLVMValueRef index = LLVMConstInt(ctx->i32, i, false);
+ LLVMValueRef index = LLVMConstInt(ac->i32, i, false);
LLVMValueRef value = LLVMBuildExtractElement(builder, vec, index, "");
LLVMBuildStore(builder, value, ptr);
}
}
-static LLVMTypeRef get_def_type(struct nir_to_llvm_context *ctx,
+static LLVMTypeRef get_def_type(struct ac_nir_context *ctx,
const nir_ssa_def *def)
{
- LLVMTypeRef type = LLVMIntTypeInContext(ctx->context, def->bit_size);
+ LLVMTypeRef type = LLVMIntTypeInContext(ctx->ac.context, def->bit_size);
if (def->num_components > 1) {
type = LLVMVectorType(type, def->num_components);
}
return type;
}
-static LLVMValueRef get_src(struct nir_to_llvm_context *ctx, nir_src src)
+static LLVMValueRef get_src(struct ac_nir_context *nir, nir_src src)
{
assert(src.is_ssa);
- struct hash_entry *entry = _mesa_hash_table_search(ctx->defs, src.ssa);
+ struct hash_entry *entry = _mesa_hash_table_search(nir->defs, src.ssa);
return (LLVMValueRef)entry->data;
}
-static LLVMBasicBlockRef get_block(struct nir_to_llvm_context *ctx,
+static LLVMBasicBlockRef get_block(struct ac_nir_context *nir,
const struct nir_block *b)
{
- struct hash_entry *entry = _mesa_hash_table_search(ctx->defs, b);
+ struct hash_entry *entry = _mesa_hash_table_search(nir->defs, b);
return (LLVMBasicBlockRef)entry->data;
}
-static LLVMValueRef get_alu_src(struct nir_to_llvm_context *ctx,
+static LLVMValueRef get_alu_src(struct ac_nir_context *ctx,
nir_alu_src src,
unsigned num_components)
{
if (need_swizzle || num_components != src_components) {
LLVMValueRef masks[] = {
- LLVMConstInt(ctx->i32, src.swizzle[0], false),
- LLVMConstInt(ctx->i32, src.swizzle[1], false),
- LLVMConstInt(ctx->i32, src.swizzle[2], false),
- LLVMConstInt(ctx->i32, src.swizzle[3], false)};
+ LLVMConstInt(ctx->ac.i32, src.swizzle[0], false),
+ LLVMConstInt(ctx->ac.i32, src.swizzle[1], false),
+ LLVMConstInt(ctx->ac.i32, src.swizzle[2], false),
+ LLVMConstInt(ctx->ac.i32, src.swizzle[3], false)};
if (src_components > 1 && num_components == 1) {
- value = LLVMBuildExtractElement(ctx->builder, value,
+ value = LLVMBuildExtractElement(ctx->ac.builder, value,
masks[0], "");
} else if (src_components == 1 && num_components > 1) {
LLVMValueRef values[] = {value, value, value, value};
value = ac_build_gather_values(&ctx->ac, values, num_components);
} else {
LLVMValueRef swizzle = LLVMConstVector(masks, num_components);
- value = LLVMBuildShuffleVector(ctx->builder, value, value,
+ value = LLVMBuildShuffleVector(ctx->ac.builder, value, value,
swizzle, "");
}
}
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,
*/
LLVMConstInt(ctx->i1, 1, false),
};
- return ac_build_intrinsic(ctx, "llvm.cttz.i32", ctx->i32, params, 2, AC_FUNC_ATTR_READNONE);
+
+ LLVMValueRef lsb = ac_build_intrinsic(ctx, "llvm.cttz.i32", ctx->i32,
+ params, 2,
+ AC_FUNC_ATTR_READNONE);
+
+ /* TODO: We need an intrinsic to skip this conditional. */
+ /* Check for zero: */
+ return LLVMBuildSelect(ctx->builder, LLVMBuildICmp(ctx->builder,
+ LLVMIntEQ, src0,
+ ctx->i32_0, ""),
+ LLVMConstInt(ctx->i32, -1, 0), lsb, "");
}
static LLVMValueRef emit_ifind_msb(struct ac_llvm_context *ctx,
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, "");
- /* TODO SI/CIK options here */
if (ctx->options->chip_class >= VI) {
LLVMValueRef args[2];
/* Check if the result is a denormal - and flush to 0 if so. */
if (ctx->options->chip_class >= VI)
result = LLVMBuildSelect(ctx->builder, cond, ctx->f32zero, result, "");
-
+ else {
+ /* for SI/CIK */
+ /* 0x38800000 is smallest half float value (2^-14) in 32-bit float,
+ * so compare the result and flush to 0 if it's smaller.
+ */
+ LLVMValueRef temp, cond2;
+ temp = emit_intrin_1f_param(&ctx->ac, "llvm.fabs",
+ ctx->f32, result);
+ cond = LLVMBuildFCmp(ctx->builder, LLVMRealUGT,
+ LLVMBuildBitCast(ctx->builder, LLVMConstInt(ctx->i32, 0x38800000, false), ctx->f32, ""),
+ temp, "");
+ cond2 = LLVMBuildFCmp(ctx->builder, LLVMRealUNE,
+ temp, ctx->f32zero, "");
+ cond = LLVMBuildAnd(ctx->builder, cond, cond2, "");
+ result = LLVMBuildSelect(ctx->builder, cond, ctx->f32zero, result, "");
+ }
return result;
}
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++) {
return result;
}
-static LLVMValueRef emit_ddxy(struct nir_to_llvm_context *ctx,
+static LLVMValueRef emit_ddxy(struct ac_nir_context *ctx,
nir_op op,
LLVMValueRef src0)
{
int idx;
LLVMValueRef result;
- if (!ctx->lds && !ctx->has_ds_bpermute)
- ctx->lds = LLVMAddGlobalInAddressSpace(ctx->module,
- LLVMArrayType(ctx->i32, 64),
- "ddxy_lds", LOCAL_ADDR_SPACE);
-
if (op == nir_op_fddx_fine || op == nir_op_fddx)
mask = AC_TID_MASK_LEFT;
else if (op == nir_op_fddy_fine || op == nir_op_fddy)
else
idx = 2;
- result = ac_build_ddxy(&ctx->ac, ctx->has_ds_bpermute,
- mask, idx, ctx->lds,
- src0);
+ result = ac_build_ddxy(&ctx->ac, mask, idx, src0);
return result;
}
* it returns DDX(I), DDX(J), DDY(I), DDY(J).
*/
static LLVMValueRef emit_ddxy_interp(
- struct nir_to_llvm_context *ctx,
+ struct ac_nir_context *ctx,
LLVMValueRef interp_ij)
{
LLVMValueRef result[4], a;
unsigned i;
for (i = 0; i < 2; i++) {
- a = LLVMBuildExtractElement(ctx->builder, interp_ij,
- LLVMConstInt(ctx->i32, i, false), "");
+ a = LLVMBuildExtractElement(ctx->ac.builder, interp_ij,
+ LLVMConstInt(ctx->ac.i32, i, false), "");
result[i] = emit_ddxy(ctx, nir_op_fddx, a);
result[2+i] = emit_ddxy(ctx, nir_op_fddy, a);
}
return ac_build_gather_values(&ctx->ac, result, 4);
}
-static void visit_alu(struct nir_to_llvm_context *ctx, const nir_alu_instr *instr)
+static void visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
{
LLVMValueRef src[4], result = NULL;
unsigned num_components = instr->dest.dest.ssa.num_components;
result = src[0];
break;
case nir_op_fneg:
- src[0] = to_float(&ctx->ac, src[0]);
- result = LLVMBuildFNeg(ctx->builder, src[0], "");
+ src[0] = ac_to_float(&ctx->ac, src[0]);
+ result = LLVMBuildFNeg(ctx->ac.builder, src[0], "");
break;
case nir_op_ineg:
- result = LLVMBuildNeg(ctx->builder, src[0], "");
+ result = LLVMBuildNeg(ctx->ac.builder, src[0], "");
break;
case nir_op_inot:
- result = LLVMBuildNot(ctx->builder, src[0], "");
+ result = LLVMBuildNot(ctx->ac.builder, src[0], "");
break;
case nir_op_iadd:
- result = LLVMBuildAdd(ctx->builder, src[0], src[1], "");
+ 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]);
- result = LLVMBuildFAdd(ctx->builder, src[0], 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]);
- result = LLVMBuildFSub(ctx->builder, src[0], 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 = LLVMBuildSub(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildSub(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_imul:
- result = LLVMBuildMul(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildMul(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_imod:
- result = LLVMBuildSRem(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildSRem(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_umod:
- result = LLVMBuildURem(ctx->builder, src[0], src[1], "");
+ 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);
- result = LLVMBuildFMul(ctx->builder, src[1] , result, "");
- result = LLVMBuildFSub(ctx->builder, src[0], 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]);
- result = LLVMBuildFRem(ctx->builder, src[0], 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 = LLVMBuildSRem(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildSRem(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_idiv:
- result = LLVMBuildSDiv(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildSDiv(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_udiv:
- result = LLVMBuildUDiv(ctx->builder, src[0], src[1], "");
+ 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]);
- result = LLVMBuildFMul(ctx->builder, src[0], 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]);
- result = ac_build_fdiv(&ctx->ac, ctx->f32one, 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:
- result = LLVMBuildAnd(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildAnd(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_ior:
- result = LLVMBuildOr(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildOr(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_ixor:
- result = LLVMBuildXor(ctx->builder, src[0], src[1], "");
+ result = LLVMBuildXor(ctx->ac.builder, src[0], src[1], "");
break;
case nir_op_ishl:
- result = LLVMBuildShl(ctx->builder, src[0],
- LLVMBuildZExt(ctx->builder, src[1],
+ result = LLVMBuildShl(ctx->ac.builder, src[0],
+ LLVMBuildZExt(ctx->ac.builder, src[1],
LLVMTypeOf(src[0]), ""),
"");
break;
case nir_op_ishr:
- result = LLVMBuildAShr(ctx->builder, src[0],
- LLVMBuildZExt(ctx->builder, src[1],
+ result = LLVMBuildAShr(ctx->ac.builder, src[0],
+ LLVMBuildZExt(ctx->ac.builder, src[1],
LLVMTypeOf(src[0]), ""),
"");
break;
case nir_op_ushr:
- result = LLVMBuildLShr(ctx->builder, src[0],
- LLVMBuildZExt(ctx->builder, src[1],
+ result = LLVMBuildLShr(ctx->ac.builder, src[0],
+ LLVMBuildZExt(ctx->ac.builder, src[1],
LLVMTypeOf(src[0]), ""),
"");
break;
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]);
- result = ac_build_fdiv(&ctx->ac, ctx->f32one, result);
+ 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);
result = emit_bitfield_insert(&ctx->ac, src[0], src[1], src[2], src[3]);
break;
case nir_op_bitfield_reverse:
- result = ac_build_intrinsic(&ctx->ac, "llvm.bitreverse.i32", ctx->i32, src, 1, AC_FUNC_ATTR_READNONE);
+ result = ac_build_intrinsic(&ctx->ac, "llvm.bitreverse.i32", ctx->ac.i32, src, 1, AC_FUNC_ATTR_READNONE);
break;
case nir_op_bit_count:
- result = ac_build_intrinsic(&ctx->ac, "llvm.ctpop.i32", ctx->i32, src, 1, AC_FUNC_ATTR_READNONE);
+ result = ac_build_intrinsic(&ctx->ac, "llvm.ctpop.i32", ctx->ac.i32, src, 1, AC_FUNC_ATTR_READNONE);
break;
case nir_op_vec2:
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]);
- result = LLVMBuildFPToSI(ctx->builder, src[0], def_type, "");
+ 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]);
- result = LLVMBuildFPToUI(ctx->builder, src[0], def_type, "");
+ 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->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->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->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->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->builder, src[0], def_type, "");
+ result = LLVMBuildZExt(ctx->ac.builder, src[0], def_type, "");
else
- result = LLVMBuildTrunc(ctx->builder, src[0], def_type, "");
+ result = LLVMBuildTrunc(ctx->ac.builder, src[0], def_type, "");
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->builder, src[0], def_type, "");
+ result = LLVMBuildSExt(ctx->ac.builder, src[0], def_type, "");
else
- result = LLVMBuildTrunc(ctx->builder, src[0], def_type, "");
+ result = LLVMBuildTrunc(ctx->ac.builder, src[0], def_type, "");
break;
case nir_op_bcsel:
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, src[0]);
+ 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:
case nir_op_unpack_64_2x32_split_x: {
assert(instr->src[0].src.ssa->num_components == 1);
- LLVMValueRef tmp = LLVMBuildBitCast(ctx->builder, src[0],
- LLVMVectorType(ctx->i32, 2),
+ LLVMValueRef tmp = LLVMBuildBitCast(ctx->ac.builder, src[0],
+ LLVMVectorType(ctx->ac.i32, 2),
"");
- result = LLVMBuildExtractElement(ctx->builder, tmp,
- ctx->i32zero, "");
+ result = LLVMBuildExtractElement(ctx->ac.builder, tmp,
+ ctx->ac.i32_0, "");
break;
}
case nir_op_unpack_64_2x32_split_y: {
assert(instr->src[0].src.ssa->num_components == 1);
- LLVMValueRef tmp = LLVMBuildBitCast(ctx->builder, src[0],
- LLVMVectorType(ctx->i32, 2),
+ LLVMValueRef tmp = LLVMBuildBitCast(ctx->ac.builder, src[0],
+ LLVMVectorType(ctx->ac.i32, 2),
"");
- result = LLVMBuildExtractElement(ctx->builder, tmp,
- ctx->i32one, "");
+ result = LLVMBuildExtractElement(ctx->ac.builder, tmp,
+ ctx->ac.i32_1, "");
break;
}
case nir_op_pack_64_2x32_split: {
- LLVMValueRef tmp = LLVMGetUndef(LLVMVectorType(ctx->i32, 2));
- tmp = LLVMBuildInsertElement(ctx->builder, tmp,
- src[0], ctx->i32zero, "");
- tmp = LLVMBuildInsertElement(ctx->builder, tmp,
- src[1], ctx->i32one, "");
- result = LLVMBuildBitCast(ctx->builder, tmp, ctx->i64, "");
+ LLVMValueRef tmp = LLVMGetUndef(LLVMVectorType(ctx->ac.i32, 2));
+ tmp = LLVMBuildInsertElement(ctx->ac.builder, tmp,
+ src[0], ctx->ac.i32_0, "");
+ tmp = LLVMBuildInsertElement(ctx->ac.builder, tmp,
+ src[1], ctx->ac.i32_1, "");
+ result = LLVMBuildBitCast(ctx->ac.builder, tmp, ctx->ac.i64, "");
break;
}
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);
}
}
-static void visit_load_const(struct nir_to_llvm_context *ctx,
+static void visit_load_const(struct ac_nir_context *ctx,
const nir_load_const_instr *instr)
{
LLVMValueRef values[4], value = NULL;
LLVMTypeRef element_type =
- LLVMIntTypeInContext(ctx->context, instr->def.bit_size);
+ LLVMIntTypeInContext(ctx->ac.context, instr->def.bit_size);
for (unsigned i = 0; i < instr->def.num_components; ++i) {
switch (instr->def.bit_size) {
}
static LLVMValueRef
-get_buffer_size(struct nir_to_llvm_context *ctx, LLVMValueRef descriptor, bool in_elements)
+get_buffer_size(struct ac_nir_context *ctx, LLVMValueRef descriptor, bool in_elements)
{
LLVMValueRef size =
- LLVMBuildExtractElement(ctx->builder, descriptor,
- LLVMConstInt(ctx->i32, 2, false), "");
+ LLVMBuildExtractElement(ctx->ac.builder, descriptor,
+ LLVMConstInt(ctx->ac.i32, 2, false), "");
/* VI only */
- if (ctx->options->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.
*/
LLVMValueRef stride =
- LLVMBuildExtractElement(ctx->builder, descriptor,
- LLVMConstInt(ctx->i32, 1, false), "");
- stride = LLVMBuildLShr(ctx->builder, stride,
- LLVMConstInt(ctx->i32, 16, false), "");
- stride = LLVMBuildAnd(ctx->builder, stride,
- LLVMConstInt(ctx->i32, 0x3fff, false), "");
+ LLVMBuildExtractElement(ctx->ac.builder, descriptor,
+ LLVMConstInt(ctx->ac.i32, 1, false), "");
+ stride = LLVMBuildLShr(ctx->ac.builder, stride,
+ LLVMConstInt(ctx->ac.i32, 16, false), "");
+ stride = LLVMBuildAnd(ctx->ac.builder, stride,
+ LLVMConstInt(ctx->ac.i32, 0x3fff, false), "");
- size = LLVMBuildUDiv(ctx->builder, size, stride, "");
+ size = LLVMBuildUDiv(ctx->ac.builder, size, stride, "");
}
return size;
}
strcpy(buf, "i32");
}
-static LLVMValueRef radv_lower_gather4_integer(struct nir_to_llvm_context *ctx,
+static LLVMValueRef radv_lower_gather4_integer(struct ac_llvm_context *ctx,
struct ac_image_args *args,
const nir_tex_instr *instr)
{
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;
txq_args.da = instr->is_array || instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE;
txq_args.opcode = ac_image_get_resinfo;
txq_args.dmask = 0xf;
- txq_args.addr = ctx->i32zero;
+ txq_args.addr = ctx->i32_0;
txq_args.resource = args->resource;
- LLVMValueRef size = ac_build_image_opcode(&ctx->ac, &txq_args);
+ LLVMValueRef size = ac_build_image_opcode(ctx, &txq_args);
for (c = 0; c < 2; c++) {
half_texel[c] = LLVMBuildExtractElement(ctx->builder, size,
LLVMConstInt(ctx->i32, c, false), "");
half_texel[c] = LLVMBuildUIToFP(ctx->builder, half_texel[c], ctx->f32, "");
- half_texel[c] = ac_build_fdiv(&ctx->ac, ctx->f32one, half_texel[c]);
+ half_texel[c] = ac_build_fdiv(ctx, ctx->f32_1, half_texel[c]);
half_texel[c] = LLVMBuildFMul(ctx->builder, half_texel[c],
LLVMConstReal(ctx->f32, -0.5), "");
}
/* workaround 8/8/8/8 uint/sint cube gather bug */
/* first detect it then change to a scaled read and f2i */
- tmp = LLVMBuildExtractElement(ctx->builder, args->resource, ctx->i32one, "");
+ tmp = LLVMBuildExtractElement(ctx->builder, args->resource, ctx->i32_1, "");
tmp2 = tmp;
/* extract the DATA_FORMAT */
- tmp = ac_build_bfe(&ctx->ac, tmp, LLVMConstInt(ctx->i32, 20, false),
+ tmp = ac_build_bfe(ctx, tmp, LLVMConstInt(ctx->i32, 20, false),
LLVMConstInt(ctx->i32, 6, false), false);
/* is the DATA_FORMAT == 8_8_8_8 */
tmp2 = LLVMBuildAnd(ctx->builder, tmp2, LLVMConstInt(ctx->i32, C_008F14_NUM_FORMAT_GFX6, false), "");
tmp2 = LLVMBuildOr(ctx->builder, tmp2, tmp, "");
- args->resource = LLVMBuildInsertElement(ctx->builder, args->resource, tmp2, ctx->i32one, "");
+ args->resource = LLVMBuildInsertElement(ctx->builder, args->resource, tmp2, ctx->i32_1, "");
/* don't modify the coordinates for this case */
coord = LLVMBuildSelect(ctx->builder, compare_cube_wa, orig_coords, coord, "");
}
args->addr = coord;
- result = ac_build_image_opcode(&ctx->ac, args);
+ result = ac_build_image_opcode(ctx, args);
if (instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE) {
LLVMValueRef tmp, tmp2;
return result;
}
-static LLVMValueRef build_tex_intrinsic(struct nir_to_llvm_context *ctx,
+static LLVMValueRef build_tex_intrinsic(struct ac_nir_context *ctx,
const nir_tex_instr *instr,
bool lod_is_zero,
struct ac_image_args *args)
return ac_build_buffer_load_format(&ctx->ac,
args->resource,
args->addr,
- LLVMConstInt(ctx->i32, 0, false),
+ LLVMConstInt(ctx->ac.i32, 0, false),
true);
}
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, args, instr);
+ return radv_lower_gather4_integer(&ctx->ac, args, instr);
}
}
return ac_build_image_opcode(&ctx->ac, args);
static LLVMValueRef visit_vulkan_resource_index(struct nir_to_llvm_context *ctx,
nir_intrinsic_instr *instr)
{
- LLVMValueRef index = get_src(ctx, instr->src[0]);
+ LLVMValueRef index = get_src(ctx->nir, instr->src[0]);
unsigned desc_set = nir_intrinsic_desc_set(instr);
unsigned binding = nir_intrinsic_binding(instr);
LLVMValueRef desc_ptr = ctx->descriptor_sets[desc_set];
LLVMValueRef ptr, addr;
addr = LLVMConstInt(ctx->i32, nir_intrinsic_base(instr), 0);
- addr = LLVMBuildAdd(ctx->builder, addr, get_src(ctx, instr->src[0]), "");
+ addr = LLVMBuildAdd(ctx->builder, addr, get_src(ctx->nir, instr->src[0]), "");
ptr = ac_build_gep0(&ctx->ac, ctx->push_constants, addr);
- ptr = cast_ptr(ctx, ptr, get_def_type(ctx, &instr->dest.ssa));
+ ptr = cast_ptr(ctx, ptr, get_def_type(ctx->nir, &instr->dest.ssa));
return LLVMBuildLoad(ctx->builder, ptr, "");
}
-static LLVMValueRef visit_get_buffer_size(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_get_buffer_size(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
LLVMValueRef desc = get_src(ctx, instr->src[0]);
return get_buffer_size(ctx, desc, false);
}
-static void visit_store_ssbo(struct nir_to_llvm_context *ctx,
+static void visit_store_ssbo(struct ac_nir_context *ctx,
nir_intrinsic_instr *instr)
{
const char *store_name;
LLVMValueRef src_data = get_src(ctx, instr->src[0]);
- LLVMTypeRef data_type = ctx->f32;
+ LLVMTypeRef data_type = ctx->ac.f32;
int elem_size_mult = get_elem_bits(&ctx->ac, LLVMTypeOf(src_data)) / 32;
int components_32bit = elem_size_mult * instr->num_components;
unsigned writemask = nir_intrinsic_write_mask(instr);
LLVMValueRef base_data, base_offset;
LLVMValueRef params[6];
+ LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
- if (ctx->stage == MESA_SHADER_FRAGMENT)
- ctx->shader_info->fs.writes_memory = true;
-
- params[1] = get_src(ctx, instr->src[1]);
- params[2] = LLVMConstInt(ctx->i32, 0, false); /* vindex */
- params[4] = ctx->i1false; /* glc */
- params[5] = ctx->i1false; /* slc */
+ params[1] = ctx->abi->load_ssbo(ctx->abi,
+ get_src(ctx, instr->src[1]), true);
+ params[2] = LLVMConstInt(ctx->ac.i32, 0, false); /* vindex */
+ params[4] = i1false; /* glc */
+ params[5] = i1false; /* slc */
if (components_32bit > 1)
- data_type = LLVMVectorType(ctx->f32, components_32bit);
+ data_type = LLVMVectorType(ctx->ac.f32, components_32bit);
- base_data = to_float(&ctx->ac, src_data);
- base_data = trim_vector(ctx, base_data, instr->num_components);
- base_data = LLVMBuildBitCast(ctx->builder, base_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, "");
base_offset = get_src(ctx, instr->src[2]); /* voffset */
while (writemask) {
store_name = "llvm.amdgcn.buffer.store.v4f32";
data = base_data;
} else if (count == 2) {
- tmp = LLVMBuildExtractElement(ctx->builder,
- base_data, LLVMConstInt(ctx->i32, start, false), "");
- data = LLVMBuildInsertElement(ctx->builder, LLVMGetUndef(ctx->v2f32), tmp,
- ctx->i32zero, "");
-
- tmp = LLVMBuildExtractElement(ctx->builder,
- base_data, LLVMConstInt(ctx->i32, start + 1, false), "");
- data = LLVMBuildInsertElement(ctx->builder, data, tmp,
- ctx->i32one, "");
+ LLVMTypeRef v2f32 = LLVMVectorType(ctx->ac.f32, 2);
+
+ tmp = LLVMBuildExtractElement(ctx->ac.builder,
+ base_data, LLVMConstInt(ctx->ac.i32, start, false), "");
+ data = LLVMBuildInsertElement(ctx->ac.builder, LLVMGetUndef(v2f32), tmp,
+ ctx->ac.i32_0, "");
+
+ tmp = LLVMBuildExtractElement(ctx->ac.builder,
+ base_data, LLVMConstInt(ctx->ac.i32, start + 1, false), "");
+ data = LLVMBuildInsertElement(ctx->ac.builder, data, tmp,
+ ctx->ac.i32_1, "");
store_name = "llvm.amdgcn.buffer.store.v2f32";
} else {
assert(count == 1);
if (get_llvm_num_components(base_data) > 1)
- data = LLVMBuildExtractElement(ctx->builder, base_data,
- LLVMConstInt(ctx->i32, start, false), "");
+ data = LLVMBuildExtractElement(ctx->ac.builder, base_data,
+ LLVMConstInt(ctx->ac.i32, start, false), "");
else
data = base_data;
store_name = "llvm.amdgcn.buffer.store.f32";
offset = base_offset;
if (start != 0) {
- offset = LLVMBuildAdd(ctx->builder, offset, LLVMConstInt(ctx->i32, start * 4, false), "");
+ offset = LLVMBuildAdd(ctx->ac.builder, offset, LLVMConstInt(ctx->ac.i32, start * 4, false), "");
}
params[0] = data;
params[3] = offset;
ac_build_intrinsic(&ctx->ac, store_name,
- ctx->voidt, params, 6, 0);
+ ctx->ac.voidt, params, 6, 0);
}
}
-static LLVMValueRef visit_atomic_ssbo(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_atomic_ssbo(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
const char *name;
LLVMValueRef params[6];
int arg_count = 0;
- if (ctx->stage == MESA_SHADER_FRAGMENT)
- ctx->shader_info->fs.writes_memory = true;
if (instr->intrinsic == nir_intrinsic_ssbo_atomic_comp_swap) {
- params[arg_count++] = llvm_extract_elem(ctx, get_src(ctx, instr->src[3]), 0);
+ params[arg_count++] = llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[3]), 0);
}
- params[arg_count++] = llvm_extract_elem(ctx, get_src(ctx, instr->src[2]), 0);
- params[arg_count++] = get_src(ctx, instr->src[0]);
- params[arg_count++] = LLVMConstInt(ctx->i32, 0, false); /* vindex */
+ params[arg_count++] = llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[2]), 0);
+ params[arg_count++] = ctx->abi->load_ssbo(ctx->abi,
+ get_src(ctx, instr->src[0]),
+ true);
+ params[arg_count++] = LLVMConstInt(ctx->ac.i32, 0, false); /* vindex */
params[arg_count++] = get_src(ctx, instr->src[1]); /* voffset */
- params[arg_count++] = ctx->i1false; /* slc */
+ params[arg_count++] = LLVMConstInt(ctx->ac.i1, 0, false); /* slc */
switch (instr->intrinsic) {
case nir_intrinsic_ssbo_atomic_add:
abort();
}
- return ac_build_intrinsic(&ctx->ac, name, ctx->i32, params, arg_count, 0);
+ return ac_build_intrinsic(&ctx->ac, name, ctx->ac.i32, params, arg_count, 0);
}
-static LLVMValueRef visit_load_buffer(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_load_buffer(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
LLVMValueRef results[2];
for (int i = 0; i < num_components; i += load_components) {
load_components = MIN2(num_components - i, 4);
const char *load_name;
- LLVMTypeRef data_type = ctx->f32;
- LLVMValueRef offset = LLVMConstInt(ctx->i32, i * 4, false);
- offset = LLVMBuildAdd(ctx->builder, get_src(ctx, instr->src[1]), offset, "");
+ LLVMTypeRef data_type = ctx->ac.f32;
+ LLVMValueRef offset = LLVMConstInt(ctx->ac.i32, i * 4, false);
+ offset = LLVMBuildAdd(ctx->ac.builder, get_src(ctx, instr->src[1]), offset, "");
if (load_components == 3)
- data_type = LLVMVectorType(ctx->f32, 4);
+ data_type = LLVMVectorType(ctx->ac.f32, 4);
else if (load_components > 1)
- data_type = LLVMVectorType(ctx->f32, load_components);
+ data_type = LLVMVectorType(ctx->ac.f32, load_components);
if (load_components >= 3)
load_name = "llvm.amdgcn.buffer.load.v4f32";
else
unreachable("unhandled number of components");
+ LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
LLVMValueRef params[] = {
- get_src(ctx, instr->src[0]),
- LLVMConstInt(ctx->i32, 0, false),
+ ctx->abi->load_ssbo(ctx->abi,
+ get_src(ctx, instr->src[0]),
+ false),
+ LLVMConstInt(ctx->ac.i32, 0, false),
offset,
- ctx->i1false,
- ctx->i1false,
+ i1false,
+ i1false,
};
results[i] = ac_build_intrinsic(&ctx->ac, load_name, data_type, params, 5, 0);
LLVMValueRef ret = results[0];
if (num_components > 4 || num_components == 3) {
LLVMValueRef masks[] = {
- LLVMConstInt(ctx->i32, 0, false), LLVMConstInt(ctx->i32, 1, false),
- LLVMConstInt(ctx->i32, 2, false), LLVMConstInt(ctx->i32, 3, false),
- LLVMConstInt(ctx->i32, 4, false), LLVMConstInt(ctx->i32, 5, false),
- LLVMConstInt(ctx->i32, 6, false), LLVMConstInt(ctx->i32, 7, false)
+ LLVMConstInt(ctx->ac.i32, 0, false), LLVMConstInt(ctx->ac.i32, 1, false),
+ LLVMConstInt(ctx->ac.i32, 2, false), LLVMConstInt(ctx->ac.i32, 3, false),
+ LLVMConstInt(ctx->ac.i32, 4, false), LLVMConstInt(ctx->ac.i32, 5, false),
+ LLVMConstInt(ctx->ac.i32, 6, false), LLVMConstInt(ctx->ac.i32, 7, false)
};
LLVMValueRef swizzle = LLVMConstVector(masks, num_components);
- ret = LLVMBuildShuffleVector(ctx->builder, results[0],
+ ret = LLVMBuildShuffleVector(ctx->ac.builder, results[0],
results[num_components > 4 ? 1 : 0], swizzle, "");
}
- return LLVMBuildBitCast(ctx->builder, ret,
+ return LLVMBuildBitCast(ctx->ac.builder, ret,
get_def_type(ctx, &instr->dest.ssa), "");
}
-static LLVMValueRef visit_load_ubo_buffer(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_load_ubo_buffer(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
LLVMValueRef results[8], ret;
LLVMValueRef offset = get_src(ctx, instr->src[1]);
int num_components = instr->num_components;
+ if (ctx->abi->load_ubo)
+ rsrc = ctx->abi->load_ubo(ctx->abi, rsrc);
+
if (instr->dest.ssa.bit_size == 64)
num_components *= 2;
for (unsigned i = 0; i < num_components; ++i) {
LLVMValueRef params[] = {
rsrc,
- LLVMBuildAdd(ctx->builder, LLVMConstInt(ctx->i32, 4 * i, 0),
+ LLVMBuildAdd(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, 4 * i, 0),
offset, "")
};
- results[i] = ac_build_intrinsic(&ctx->ac, "llvm.SI.load.const.v4i32", ctx->f32,
+ results[i] = ac_build_intrinsic(&ctx->ac, "llvm.SI.load.const.v4i32", ctx->ac.f32,
params, 2,
AC_FUNC_ATTR_READNONE |
AC_FUNC_ATTR_LEGACY);
ret = ac_build_gather_values(&ctx->ac, results, instr->num_components);
- return LLVMBuildBitCast(ctx->builder, ret,
+ return LLVMBuildBitCast(ctx->ac.builder, ret,
get_def_type(ctx, &instr->dest.ssa), "");
}
static void
-radv_get_deref_offset(struct nir_to_llvm_context *ctx, nir_deref_var *deref,
- bool vs_in, unsigned *vertex_index_out,
- LLVMValueRef *vertex_index_ref,
- unsigned *const_out, LLVMValueRef *indir_out)
+get_deref_offset(struct ac_nir_context *ctx, nir_deref_var *deref,
+ bool vs_in, unsigned *vertex_index_out,
+ LLVMValueRef *vertex_index_ref,
+ unsigned *const_out, LLVMValueRef *indir_out)
{
unsigned const_offset = 0;
nir_deref *tail = &deref->deref;
*vertex_index_out = deref_array->base_offset;
if (vertex_index_ref) {
- LLVMValueRef vtx = LLVMConstInt(ctx->i32, deref_array->base_offset, false);
+ LLVMValueRef vtx = LLVMConstInt(ctx->ac.i32, deref_array->base_offset, false);
if (deref_array->deref_array_type == nir_deref_array_type_indirect) {
- vtx = LLVMBuildAdd(ctx->builder, vtx, get_src(ctx, deref_array->indirect), "");
+ vtx = LLVMBuildAdd(ctx->ac.builder, vtx, get_src(ctx, deref_array->indirect), "");
}
*vertex_index_ref = vtx;
}
assert(deref_array->deref_array_type == nir_deref_array_type_indirect);
index = get_src(ctx, deref_array->indirect);
- stride = LLVMConstInt(ctx->i32, size, 0);
- local_offset = LLVMBuildMul(ctx->builder, stride, index, "");
+ stride = LLVMConstInt(ctx->ac.i32, size, 0);
+ local_offset = LLVMBuildMul(ctx->ac.builder, stride, index, "");
if (offset)
- offset = LLVMBuildAdd(ctx->builder, offset, local_offset, "");
+ offset = LLVMBuildAdd(ctx->ac.builder, offset, local_offset, "");
else
offset = local_offset;
} else if (tail->deref_type == nir_deref_type_struct) {
}
out:
if (const_offset && offset)
- offset = LLVMBuildAdd(ctx->builder, offset,
- LLVMConstInt(ctx->i32, const_offset, 0),
+ offset = LLVMBuildAdd(ctx->ac.builder, offset,
+ LLVMConstInt(ctx->ac.i32, const_offset, 0),
"");
*const_out = const_offset;
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;
}
LLVMValueRef param_stride, constant16;
LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
- vertices_per_patch = unpack_param(ctx, ctx->tcs_offchip_layout, 9, 6);
- num_patches = unpack_param(ctx, ctx->tcs_offchip_layout, 0, 9);
+ vertices_per_patch = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 9, 6);
+ num_patches = unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 0, 9);
total_vertices = LLVMBuildMul(ctx->builder, vertices_per_patch,
num_patches, "");
if (!vertex_index) {
LLVMValueRef patch_data_offset =
- unpack_param(ctx, ctx->tcs_offchip_layout, 16, 16);
+ unpack_param(&ctx->ac, ctx->tcs_offchip_layout, 16, 16);
base_addr = LLVMBuildAdd(ctx->builder, base_addr,
patch_data_offset, "");
const bool per_vertex = nir_is_per_vertex_io(instr->variables[0]->var, ctx->stage);
const bool is_compact = instr->variables[0]->var->data.compact;
param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
- radv_get_deref_offset(ctx, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
+ get_deref_offset(ctx->nir, instr->variables[0],
+ false, NULL, per_vertex ? &vertex_index : NULL,
+ &const_index, &indir_index);
- stride = unpack_param(ctx, ctx->tcs_in_layout, 13, 8);
+ stride = unpack_param(&ctx->ac, ctx->tcs_in_layout, 13, 8);
dw_addr = get_tcs_in_current_patch_offset(ctx);
dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
indir_index);
ctx->i32one, "");
}
result = ac_build_gather_values(&ctx->ac, value, instr->num_components);
- result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), "");
+ result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), "");
return result;
}
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;
const bool per_vertex = nir_is_per_vertex_io(instr->variables[0]->var, ctx->stage);
const bool is_compact = instr->variables[0]->var->data.compact;
param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
- radv_get_deref_offset(ctx, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
+ get_deref_offset(ctx->nir, instr->variables[0],
+ false, NULL, per_vertex ? &vertex_index : NULL,
+ &const_index, &indir_index);
if (!instr->variables[0]->var->data.patch) {
- stride = unpack_param(ctx, ctx->tcs_out_layout, 13, 8);
+ stride = unpack_param(&ctx->ac, ctx->tcs_out_layout, 13, 8);
dw_addr = get_tcs_out_current_patch_offset(ctx);
} else {
dw_addr = get_tcs_out_current_patch_data_offset(ctx);
ctx->i32one, "");
}
result = ac_build_gather_values(&ctx->ac, value, instr->num_components);
- result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), "");
+ result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), "");
return result;
}
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;
const bool per_vertex = nir_is_per_vertex_io(instr->variables[0]->var, ctx->stage);
const bool is_compact = instr->variables[0]->var->data.compact;
- radv_get_deref_offset(ctx, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
+ get_deref_offset(ctx->nir, instr->variables[0],
+ false, NULL, per_vertex ? &vertex_index : NULL,
+ &const_index, &indir_index);
param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
if (instr->variables[0]->var->data.location == VARYING_SLOT_CLIP_DIST0 &&
}
if (!instr->variables[0]->var->data.patch) {
- stride = unpack_param(ctx, ctx->tcs_out_layout, 13, 8);
+ stride = unpack_param(&ctx->ac, ctx->tcs_out_layout, 13, 8);
dw_addr = get_tcs_out_current_patch_offset(ctx);
} else {
dw_addr = get_tcs_out_current_patch_data_offset(ctx);
buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, is_compact,
vertex_index, indir_index);
+ bool is_tess_factor = false;
+ if (instr->variables[0]->var->data.location == VARYING_SLOT_TESS_LEVEL_INNER ||
+ instr->variables[0]->var->data.location == VARYING_SLOT_TESS_LEVEL_OUTER)
+ is_tess_factor = true;
+
unsigned base = is_compact ? const_index : 0;
for (unsigned chan = 0; chan < 8; chan++) {
- bool is_tess_factor = false;
if (!(writemask & (1 << chan)))
continue;
- LLVMValueRef value = llvm_extract_elem(ctx, src, chan);
+ LLVMValueRef value = llvm_extract_elem(&ctx->ac, src, chan);
lds_store(ctx, dw_addr, value);
- if (instr->variables[0]->var->data.location == VARYING_SLOT_TESS_LEVEL_INNER ||
- instr->variables[0]->var->data.location == VARYING_SLOT_TESS_LEVEL_OUTER)
- is_tess_factor = true;
-
if (!is_tess_factor && writemask != 0xF)
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, value, 1,
buf_addr, ctx->oc_lds,
const bool per_vertex = nir_is_per_vertex_io(instr->variables[0]->var, ctx->stage);
const bool is_compact = instr->variables[0]->var->data.compact;
- radv_get_deref_offset(ctx, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
+ get_deref_offset(ctx->nir, instr->variables[0],
+ false, NULL, per_vertex ? &vertex_index : NULL,
+ &const_index, &indir_index);
param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
if (instr->variables[0]->var->data.location == VARYING_SLOT_CLIP_DIST0 &&
is_compact && const_index > 3) {
result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, instr->num_components, NULL,
buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false);
- result = trim_vector(ctx, result, instr->num_components);
- result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), "");
+ result = trim_vector(&ctx->ac, result, instr->num_components);
+ result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), "");
return result;
}
unsigned param, vtx_offset_param;
LLVMValueRef value[4], result;
unsigned vertex_index;
- radv_get_deref_offset(ctx, instr->variables[0],
- false, &vertex_index, NULL,
- &const_index, &indir_index);
+ get_deref_offset(ctx->nir, instr->variables[0],
+ false, &vertex_index, NULL,
+ &const_index, &indir_index);
vtx_offset_param = vertex_index;
assert(vtx_offset_param < 6);
vtx_offset = LLVMBuildMul(ctx->builder, ctx->gs_vtx_offset[vtx_offset_param],
param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
for (unsigned i = 0; i < instr->num_components; i++) {
-
- args[0] = ctx->esgs_ring;
- args[1] = vtx_offset;
- args[2] = LLVMConstInt(ctx->i32, (param * 4 + i + const_index) * 256, false);
- args[3] = ctx->i32zero;
- args[4] = ctx->i32one; /* OFFEN */
- args[5] = ctx->i32zero; /* IDXEN */
- args[6] = ctx->i32one; /* GLC */
- args[7] = ctx->i32zero; /* SLC */
- args[8] = ctx->i32zero; /* TFE */
-
- value[i] = ac_build_intrinsic(&ctx->ac, "llvm.SI.buffer.load.dword.i32.i32",
- ctx->i32, args, 9,
- AC_FUNC_ATTR_READONLY |
- AC_FUNC_ATTR_LEGACY);
+ if (ctx->ac.chip_class >= GFX9) {
+ LLVMValueRef dw_addr = ctx->gs_vtx_offset[vtx_offset_param];
+ dw_addr = LLVMBuildAdd(ctx->ac.builder, dw_addr,
+ LLVMConstInt(ctx->ac.i32, param * 4 + i, 0), "");
+ value[i] = lds_load(ctx, dw_addr);
+ } else {
+ args[0] = ctx->esgs_ring;
+ args[1] = vtx_offset;
+ args[2] = LLVMConstInt(ctx->i32, (param * 4 + i + const_index) * 256, false);
+ args[3] = ctx->i32zero;
+ args[4] = ctx->i32one; /* OFFEN */
+ args[5] = ctx->i32zero; /* IDXEN */
+ args[6] = ctx->i32one; /* GLC */
+ args[7] = ctx->i32zero; /* SLC */
+ args[8] = ctx->i32zero; /* TFE */
+
+ value[i] = ac_build_intrinsic(&ctx->ac, "llvm.SI.buffer.load.dword.i32.i32",
+ ctx->i32, args, 9,
+ AC_FUNC_ATTR_READONLY |
+ AC_FUNC_ATTR_LEGACY);
+ }
}
result = ac_build_gather_values(&ctx->ac, value, instr->num_components);
}
static LLVMValueRef
-build_gep_for_deref(struct nir_to_llvm_context *ctx,
+build_gep_for_deref(struct ac_nir_context *ctx,
nir_deref_var *deref)
{
struct hash_entry *entry = _mesa_hash_table_search(ctx->vars, deref->var);
switch (tail->deref_type) {
case nir_deref_type_array: {
nir_deref_array *array = nir_deref_as_array(tail);
- offset = LLVMConstInt(ctx->i32, array->base_offset, 0);
+ offset = LLVMConstInt(ctx->ac.i32, array->base_offset, 0);
if (array->deref_array_type ==
nir_deref_array_type_indirect) {
- offset = LLVMBuildAdd(ctx->builder, offset,
+ offset = LLVMBuildAdd(ctx->ac.builder, offset,
get_src(ctx,
array->indirect),
"");
case nir_deref_type_struct: {
nir_deref_struct *deref_struct =
nir_deref_as_struct(tail);
- offset = LLVMConstInt(ctx->i32,
+ offset = LLVMConstInt(ctx->ac.i32,
deref_struct->index, 0);
break;
}
return val;
}
-static LLVMValueRef visit_load_var(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_load_var(struct ac_nir_context *ctx,
nir_intrinsic_instr *instr)
{
LLVMValueRef values[8];
unsigned const_index;
bool vs_in = ctx->stage == MESA_SHADER_VERTEX &&
instr->variables[0]->var->data.mode == nir_var_shader_in;
- radv_get_deref_offset(ctx, instr->variables[0], vs_in, NULL, NULL,
+ get_deref_offset(ctx, instr->variables[0], vs_in, NULL, NULL,
&const_index, &indir_index);
if (instr->dest.ssa.bit_size == 64)
switch (instr->variables[0]->var->data.mode) {
case nir_var_shader_in:
if (ctx->stage == MESA_SHADER_TESS_CTRL)
- return load_tcs_input(ctx, instr);
+ return load_tcs_input(ctx->nctx, instr);
if (ctx->stage == MESA_SHADER_TESS_EVAL)
- return load_tes_input(ctx, instr);
+ return load_tes_input(ctx->nctx, instr);
if (ctx->stage == MESA_SHADER_GEOMETRY) {
- return load_gs_input(ctx, instr);
+ return load_gs_input(ctx->nctx, instr);
}
for (unsigned chan = 0; chan < ve; chan++) {
if (indir_index) {
ctx->stage == MESA_SHADER_VERTEX);
count -= chan / 4;
LLVMValueRef tmp_vec = ac_build_gather_values_extended(
- &ctx->ac, ctx->inputs + idx + chan, count,
- 4, false);
+ &ctx->ac, ctx->abi->inputs + idx + chan, count,
+ 4, false, true);
- values[chan] = LLVMBuildExtractElement(ctx->builder,
+ values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
tmp_vec,
indir_index, "");
} else
- values[chan] = ctx->inputs[idx + chan + const_index * 4];
+ values[chan] = ctx->abi->inputs[idx + chan + const_index * 4];
}
break;
case nir_var_local:
count -= chan / 4;
LLVMValueRef tmp_vec = ac_build_gather_values_extended(
&ctx->ac, ctx->locals + idx + chan, count,
- 4, true);
+ 4, true, true);
- values[chan] = LLVMBuildExtractElement(ctx->builder,
+ values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
tmp_vec,
indir_index, "");
} else {
- values[chan] = LLVMBuildLoad(ctx->builder, ctx->locals[idx + chan + const_index * 4], "");
+ values[chan] = LLVMBuildLoad(ctx->ac.builder, ctx->locals[idx + chan + const_index * 4], "");
}
}
break;
case nir_var_shared: {
LLVMValueRef address = build_gep_for_deref(ctx,
instr->variables[0]);
- LLVMValueRef val = LLVMBuildLoad(ctx->builder, address, "");
- return LLVMBuildBitCast(ctx->builder, val,
+ LLVMValueRef val = LLVMBuildLoad(ctx->ac.builder, address, "");
+ return LLVMBuildBitCast(ctx->ac.builder, val,
get_def_type(ctx, &instr->dest.ssa),
"");
}
case nir_var_shader_out:
if (ctx->stage == MESA_SHADER_TESS_CTRL)
- return load_tcs_output(ctx, instr);
+ return load_tcs_output(ctx->nctx, instr);
for (unsigned chan = 0; chan < ve; chan++) {
if (indir_index) {
unsigned count = glsl_count_attribute_slots(
count -= chan / 4;
LLVMValueRef tmp_vec = ac_build_gather_values_extended(
&ctx->ac, ctx->outputs + idx + chan, count,
- 4, true);
+ 4, true, true);
- values[chan] = LLVMBuildExtractElement(ctx->builder,
+ values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
tmp_vec,
indir_index, "");
} else {
- values[chan] = LLVMBuildLoad(ctx->builder,
+ values[chan] = LLVMBuildLoad(ctx->ac.builder,
ctx->outputs[idx + chan + const_index * 4],
"");
}
unreachable("unhandle variable mode");
}
ret = ac_build_gather_values(&ctx->ac, values, ve);
- return LLVMBuildBitCast(ctx->builder, ret, get_def_type(ctx, &instr->dest.ssa), "");
+ return LLVMBuildBitCast(ctx->ac.builder, ret, get_def_type(ctx, &instr->dest.ssa), "");
}
static void
-visit_store_var(struct nir_to_llvm_context *ctx,
- nir_intrinsic_instr *instr)
+visit_store_var(struct ac_nir_context *ctx,
+ nir_intrinsic_instr *instr)
{
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;
- radv_get_deref_offset(ctx, instr->variables[0], false,
- NULL, NULL, &const_index, &indir_index);
+ get_deref_offset(ctx, instr->variables[0], false,
+ NULL, NULL, &const_index, &indir_index);
if (get_elem_bits(&ctx->ac, LLVMTypeOf(src)) == 64) {
int old_writemask = writemask;
- src = LLVMBuildBitCast(ctx->builder, src,
- LLVMVectorType(ctx->f32, get_llvm_num_components(src) * 2),
+ src = LLVMBuildBitCast(ctx->ac.builder, src,
+ LLVMVectorType(ctx->ac.f32, get_llvm_num_components(src) * 2),
"");
writemask = 0;
case nir_var_shader_out:
if (ctx->stage == MESA_SHADER_TESS_CTRL) {
- store_tcs_output(ctx, instr, src, writemask);
+ store_tcs_output(ctx->nctx, instr, src, writemask);
return;
}
if (!(writemask & (1 << chan)))
continue;
- value = llvm_extract_elem(ctx, src, chan);
+ value = llvm_extract_elem(&ctx->ac, src, chan);
if (instr->variables[0]->var->data.compact)
stride = 1;
count -= chan / 4;
LLVMValueRef tmp_vec = ac_build_gather_values_extended(
&ctx->ac, ctx->outputs + idx + chan, count,
- stride, true);
+ stride, true, true);
- if (get_llvm_num_components(tmp_vec) > 1) {
- tmp_vec = LLVMBuildInsertElement(ctx->builder, tmp_vec,
- value, indir_index, "");
- } else
- tmp_vec = value;
- build_store_values_extended(ctx, ctx->outputs + idx + chan,
+ tmp_vec = LLVMBuildInsertElement(ctx->ac.builder, tmp_vec,
+ value, indir_index, "");
+ build_store_values_extended(&ctx->ac, ctx->outputs + idx + chan,
count, stride, tmp_vec);
} else {
temp_ptr = ctx->outputs[idx + chan + const_index * stride];
- LLVMBuildStore(ctx->builder, value, temp_ptr);
+ LLVMBuildStore(ctx->ac.builder, value, temp_ptr);
}
}
break;
if (!(writemask & (1 << chan)))
continue;
- value = llvm_extract_elem(ctx, src, chan);
+ value = llvm_extract_elem(&ctx->ac, src, chan);
if (indir_index) {
unsigned count = glsl_count_attribute_slots(
instr->variables[0]->var->type, false);
count -= chan / 4;
LLVMValueRef tmp_vec = ac_build_gather_values_extended(
&ctx->ac, ctx->locals + idx + chan, count,
- 4, true);
+ 4, true, true);
- tmp_vec = LLVMBuildInsertElement(ctx->builder, tmp_vec,
+ tmp_vec = LLVMBuildInsertElement(ctx->ac.builder, tmp_vec,
value, indir_index, "");
- build_store_values_extended(ctx, ctx->locals + idx + chan,
+ build_store_values_extended(&ctx->ac, ctx->locals + idx + chan,
count, 4, tmp_vec);
} else {
temp_ptr = ctx->locals[idx + chan + const_index * 4];
- LLVMBuildStore(ctx->builder, value, temp_ptr);
+ LLVMBuildStore(ctx->ac.builder, value, temp_ptr);
}
}
break;
nir_deref_tail(&instr->variables[0]->deref)->type);
if (writemask == (1 << components) - 1) {
val = LLVMBuildBitCast(
- ctx->builder, val,
+ ctx->ac.builder, val,
LLVMGetElementType(LLVMTypeOf(address)), "");
- LLVMBuildStore(ctx->builder, val, address);
+ LLVMBuildStore(ctx->ac.builder, val, address);
} else {
for (unsigned chan = 0; chan < 4; chan++) {
if (!(writemask & (1 << chan)))
continue;
LLVMValueRef ptr =
- LLVMBuildStructGEP(ctx->builder,
+ LLVMBuildStructGEP(ctx->ac.builder,
address, chan, "");
- LLVMValueRef src = llvm_extract_elem(ctx, val,
+ LLVMValueRef src = llvm_extract_elem(&ctx->ac, val,
chan);
src = LLVMBuildBitCast(
- ctx->builder, src,
+ ctx->ac.builder, src,
LLVMGetElementType(LLVMTypeOf(ptr)), "");
- LLVMBuildStore(ctx->builder, src, ptr);
+ LLVMBuildStore(ctx->ac.builder, src, ptr);
}
}
break;
* The sample index should be adjusted as follows:
* sample_index = (fmask >> (sample_index * 4)) & 0xF;
*/
-static LLVMValueRef adjust_sample_index_using_fmask(struct nir_to_llvm_context *ctx,
+static LLVMValueRef adjust_sample_index_using_fmask(struct ac_llvm_context *ctx,
LLVMValueRef coord_x, LLVMValueRef coord_y,
LLVMValueRef coord_z,
LLVMValueRef sample_index,
args.da = coord_z ? true : false;
args.resource = fmask_desc_ptr;
args.dmask = 0xf;
- args.addr = ac_build_gather_values(&ctx->ac, fmask_load_address, coord_z ? 4 : 2);
+ args.addr = ac_build_gather_values(ctx, fmask_load_address, coord_z ? 4 : 2);
- res = ac_build_image_opcode(&ctx->ac, &args);
+ res = ac_build_image_opcode(ctx, &args);
- res = to_integer(&ctx->ac, res);
+ res = ac_to_integer(ctx, res);
LLVMValueRef four = LLVMConstInt(ctx->i32, 4, false);
LLVMValueRef F = LLVMConstInt(ctx->i32, 0xf, false);
LLVMValueRef fmask = LLVMBuildExtractElement(ctx->builder,
res,
- ctx->i32zero, "");
+ ctx->i32_0, "");
LLVMValueRef sample_index4 =
LLVMBuildMul(ctx->builder, sample_index, four, "");
LLVMValueRef fmask_word1 =
LLVMBuildExtractElement(ctx->builder, fmask_desc,
- ctx->i32one, "");
+ ctx->i32_1, "");
LLVMValueRef word1_is_nonzero =
LLVMBuildICmp(ctx->builder, LLVMIntNE,
- fmask_word1, ctx->i32zero, "");
+ fmask_word1, ctx->i32_0, "");
/* Replace the MSAA sample index. */
sample_index =
return sample_index;
}
-static LLVMValueRef get_image_coords(struct nir_to_llvm_context *ctx,
+static LLVMValueRef get_image_coords(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
const struct glsl_type *type = instr->variables[0]->var->type;
LLVMValueRef src0 = get_src(ctx, instr->src[0]);
LLVMValueRef coords[4];
LLVMValueRef masks[] = {
- LLVMConstInt(ctx->i32, 0, false), LLVMConstInt(ctx->i32, 1, false),
- LLVMConstInt(ctx->i32, 2, false), LLVMConstInt(ctx->i32, 3, false),
+ LLVMConstInt(ctx->ac.i32, 0, false), LLVMConstInt(ctx->ac.i32, 1, false),
+ LLVMConstInt(ctx->ac.i32, 2, false), LLVMConstInt(ctx->ac.i32, 3, false),
};
LLVMValueRef res;
- LLVMValueRef sample_index = llvm_extract_elem(ctx, get_src(ctx, instr->src[1]), 0);
+ LLVMValueRef sample_index = llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[1]), 0);
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];
int chan;
- fmask_load_address[0] = LLVMBuildExtractElement(ctx->builder, src0, masks[0], "");
- fmask_load_address[1] = LLVMBuildExtractElement(ctx->builder, src0, masks[1], "");
- if (glsl_sampler_type_is_array(type))
- fmask_load_address[2] = LLVMBuildExtractElement(ctx->builder, src0, masks[2], "");
+ fmask_load_address[0] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[0], "");
+ fmask_load_address[1] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[1], "");
+ if (is_array)
+ fmask_load_address[2] = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[2], "");
else
fmask_load_address[2] = NULL;
if (add_frag_pos) {
for (chan = 0; chan < 2; ++chan)
- fmask_load_address[chan] = LLVMBuildAdd(ctx->builder, fmask_load_address[chan], LLVMBuildFPToUI(ctx->builder, ctx->frag_pos[chan], ctx->i32, ""), "");
+ fmask_load_address[chan] =
+ 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,
+ sample_index = adjust_sample_index_using_fmask(&ctx->ac,
fmask_load_address[0],
fmask_load_address[1],
fmask_load_address[2],
sample_index,
- get_sampler_desc(ctx, instr->variables[0], DESC_FMASK));
+ 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->builder, src0, masks[0], "");
+ res = LLVMBuildExtractElement(ctx->ac.builder, src0, masks[0], "");
else
res = src0;
} else {
if (is_ms)
count--;
for (chan = 0; chan < count; ++chan) {
- coords[chan] = LLVMBuildExtractElement(ctx->builder, src0, masks[chan], "");
+ coords[chan] = llvm_extract_elem(&ctx->ac, src0, chan);
}
-
if (add_frag_pos) {
- for (chan = 0; chan < count; ++chan)
- coords[chan] = LLVMBuildAdd(ctx->builder, coords[chan], LLVMBuildFPToUI(ctx->builder, ctx->frag_pos[chan], ctx->i32, ""), "");
+ 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++;
}
if (count == 3) {
- coords[3] = LLVMGetUndef(ctx->i32);
+ coords[3] = LLVMGetUndef(ctx->ac.i32);
count = 4;
}
res = ac_build_gather_values(&ctx->ac, coords, count);
return res;
}
-static LLVMValueRef visit_image_load(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_image_load(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
LLVMValueRef params[7];
char intrinsic_name[64];
const nir_variable *var = instr->variables[0]->var;
const struct glsl_type *type = var->type;
+ LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
+ LLVMValueRef i1true = LLVMConstInt(ctx->ac.i1, 1, false);
+
if(instr->variables[0]->deref.child)
type = instr->variables[0]->deref.child->type;
type = glsl_without_array(type);
if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) {
- params[0] = get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER);
- params[1] = LLVMBuildExtractElement(ctx->builder, get_src(ctx, instr->src[0]),
- LLVMConstInt(ctx->i32, 0, false), ""); /* vindex */
- params[2] = LLVMConstInt(ctx->i32, 0, false); /* voffset */
- params[3] = ctx->i1false; /* glc */
- params[4] = ctx->i1false; /* slc */
- res = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.buffer.load.format.v4f32", ctx->v4f32,
+ params[0] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_BUFFER, true, false);
+ params[1] = LLVMBuildExtractElement(ctx->ac.builder, get_src(ctx, instr->src[0]),
+ ctx->ac.i32_0, ""); /* vindex */
+ params[2] = ctx->ac.i32_0; /* voffset */
+ params[3] = i1false; /* glc */
+ params[4] = i1false; /* slc */
+ res = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.buffer.load.format.v4f32", ctx->ac.v4f32,
params, 5, 0);
- res = trim_vector(ctx, res, instr->dest.ssa.num_components);
- res = to_integer(&ctx->ac, res);
+ res = trim_vector(&ctx->ac, res, instr->dest.ssa.num_components);
+ 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;
- LLVMValueRef da = is_da ? ctx->i1true : ctx->i1false;
- LLVMValueRef glc = ctx->i1false;
- LLVMValueRef slc = ctx->i1false;
+ 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;
params[0] = get_image_coords(ctx, instr);
- params[1] = get_sampler_desc(ctx, instr->variables[0], DESC_IMAGE);
- params[2] = LLVMConstInt(ctx->i32, 15, false); /* dmask */
+ params[1] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, false);
+ params[2] = LLVMConstInt(ctx->ac.i32, 15, false); /* dmask */
if (HAVE_LLVM <= 0x0309) {
- params[3] = ctx->i1false; /* r128 */
+ params[3] = i1false; /* r128 */
params[4] = da;
params[5] = glc;
params[6] = slc;
} else {
- LLVMValueRef lwe = ctx->i1false;
+ LLVMValueRef lwe = i1false;
params[3] = glc;
params[4] = slc;
params[5] = lwe;
}
ac_get_image_intr_name("llvm.amdgcn.image.load",
- ctx->v4f32, /* vdata */
+ ctx->ac.v4f32, /* vdata */
LLVMTypeOf(params[0]), /* coords */
LLVMTypeOf(params[1]), /* rsrc */
intrinsic_name, sizeof(intrinsic_name));
- res = ac_build_intrinsic(&ctx->ac, intrinsic_name, ctx->v4f32,
+ 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 nir_to_llvm_context *ctx,
+static void visit_image_store(struct ac_nir_context *ctx,
nir_intrinsic_instr *instr)
{
LLVMValueRef params[8];
char intrinsic_name[64];
const nir_variable *var = instr->variables[0]->var;
const struct glsl_type *type = glsl_without_array(var->type);
- LLVMValueRef glc = ctx->i1false;
- bool force_glc = ctx->options->chip_class == SI;
+ LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
+ LLVMValueRef i1true = LLVMConstInt(ctx->ac.i1, 1, false);
+ LLVMValueRef glc = i1false;
+ bool force_glc = ctx->ac.chip_class == SI;
if (force_glc)
- glc = ctx->i1true;
- if (ctx->stage == MESA_SHADER_FRAGMENT)
- ctx->shader_info->fs.writes_memory = true;
+ 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[1] = get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER);
- params[2] = LLVMBuildExtractElement(ctx->builder, get_src(ctx, instr->src[0]),
- LLVMConstInt(ctx->i32, 0, false), ""); /* vindex */
- params[3] = LLVMConstInt(ctx->i32, 0, false); /* voffset */
+ 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 */
+ params[3] = ctx->ac.i32_0; /* voffset */
params[4] = glc; /* glc */
- params[5] = ctx->i1false; /* slc */
- ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.buffer.store.format.v4f32", ctx->voidt,
+ params[5] = i1false; /* slc */
+ ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.buffer.store.format.v4f32", ctx->ac.voidt,
params, 6, 0);
} else {
bool is_da = glsl_sampler_type_is_array(type) ||
glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_CUBE;
- LLVMValueRef da = is_da ? ctx->i1true : ctx->i1false;
- LLVMValueRef slc = ctx->i1false;
+ 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], DESC_IMAGE);
- params[3] = LLVMConstInt(ctx->i32, 15, false); /* dmask */
+ params[2] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, true);
+ params[3] = LLVMConstInt(ctx->ac.i32, 15, false); /* dmask */
if (HAVE_LLVM <= 0x0309) {
- params[4] = ctx->i1false; /* r128 */
+ params[4] = i1false; /* r128 */
params[5] = da;
params[6] = glc;
params[7] = slc;
} else {
- LLVMValueRef lwe = ctx->i1false;
+ LLVMValueRef lwe = i1false;
params[4] = glc;
params[5] = slc;
params[6] = lwe;
LLVMTypeOf(params[2]), /* rsrc */
intrinsic_name, sizeof(intrinsic_name));
- ac_build_intrinsic(&ctx->ac, intrinsic_name, ctx->voidt,
+ ac_build_intrinsic(&ctx->ac, intrinsic_name, ctx->ac.voidt,
params, 8, 0);
}
}
-static LLVMValueRef visit_image_atomic(struct nir_to_llvm_context *ctx,
+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;
const char *atomic_name;
char intrinsic_name[41];
const struct glsl_type *type = glsl_without_array(var->type);
+ LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
+ LLVMValueRef i1true = LLVMConstInt(ctx->ac.i1, 1, false);
MAYBE_UNUSED int length;
- if (ctx->stage == MESA_SHADER_FRAGMENT)
- ctx->shader_info->fs.writes_memory = true;
-
switch (instr->intrinsic) {
case nir_intrinsic_image_atomic_add:
atomic_name = "add";
params[param_count++] = get_src(ctx, instr->src[2]);
if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) {
- params[param_count++] = get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER);
- params[param_count++] = LLVMBuildExtractElement(ctx->builder, get_src(ctx, instr->src[0]),
- LLVMConstInt(ctx->i32, 0, false), ""); /* vindex */
- params[param_count++] = ctx->i32zero; /* voffset */
- params[param_count++] = ctx->i1false; /* slc */
+ params[param_count++] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_BUFFER,
+ true, true);
+ params[param_count++] = LLVMBuildExtractElement(ctx->ac.builder, get_src(ctx, instr->src[0]),
+ ctx->ac.i32_0, ""); /* vindex */
+ params[param_count++] = ctx->ac.i32_0; /* voffset */
+ params[param_count++] = i1false; /* slc */
length = snprintf(intrinsic_name, sizeof(intrinsic_name),
"llvm.amdgcn.buffer.atomic.%s", atomic_name);
glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_CUBE;
LLVMValueRef coords = params[param_count++] = get_image_coords(ctx, instr);
- params[param_count++] = get_sampler_desc(ctx, instr->variables[0], DESC_IMAGE);
- params[param_count++] = ctx->i1false; /* r128 */
- params[param_count++] = da ? ctx->i1true : ctx->i1false; /* da */
- params[param_count++] = ctx->i1false; /* slc */
+ params[param_count++] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE,
+ true, true);
+ params[param_count++] = i1false; /* r128 */
+ params[param_count++] = da ? i1true : i1false; /* da */
+ params[param_count++] = i1false; /* slc */
build_int_type_name(LLVMTypeOf(coords),
coords_type, sizeof(coords_type));
}
assert(length < sizeof(intrinsic_name));
- return ac_build_intrinsic(&ctx->ac, intrinsic_name, ctx->i32, params, param_count, 0);
+ return ac_build_intrinsic(&ctx->ac, intrinsic_name, ctx->ac.i32, params, param_count, 0);
}
-static LLVMValueRef visit_image_size(struct nir_to_llvm_context *ctx,
+static LLVMValueRef visit_image_size(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
LLVMValueRef res;
type = instr->variables[0]->deref.child->type;
if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF)
- return get_buffer_size(ctx, get_sampler_desc(ctx, instr->variables[0], DESC_BUFFER), true);
+ return get_buffer_size(ctx,
+ get_sampler_desc(ctx, instr->variables[0],
+ AC_DESC_BUFFER, true, false), true);
struct ac_image_args args = { 0 };
args.da = da;
args.dmask = 0xf;
- args.resource = get_sampler_desc(ctx, instr->variables[0], DESC_IMAGE);
+ args.resource = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, false);
args.opcode = ac_image_get_resinfo;
- args.addr = ctx->i32zero;
+ args.addr = ctx->ac.i32_0;
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->i32, 2, false);
- LLVMValueRef six = LLVMConstInt(ctx->i32, 6, false);
- LLVMValueRef z = LLVMBuildExtractElement(ctx->builder, res, two, "");
- z = LLVMBuildSDiv(ctx->builder, z, six, "");
- res = LLVMBuildInsertElement(ctx->builder, res, z, two, "");
+ 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;
}
ctx->voidt, NULL, 0, AC_FUNC_ATTR_CONVERGENT);
}
-static void emit_discard_if(struct nir_to_llvm_context *ctx,
+static void emit_discard_if(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
LLVMValueRef cond;
- ctx->shader_info->fs.can_discard = true;
- cond = LLVMBuildICmp(ctx->builder, LLVMIntNE,
+ cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE,
get_src(ctx, instr->src[0]),
- ctx->i32zero, "");
+ ctx->ac.i32_0, "");
- cond = LLVMBuildSelect(ctx->builder, cond,
- LLVMConstReal(ctx->f32, -1.0f),
- ctx->f32zero, "");
+ cond = LLVMBuildSelect(ctx->ac.builder, cond,
+ LLVMConstReal(ctx->ac.f32, -1.0f),
+ ctx->ac.f32_0, "");
ac_build_kill(&ctx->ac, cond);
}
const nir_intrinsic_instr *instr)
{
LLVMValueRef ptr, result;
- LLVMValueRef src = get_src(ctx, instr->src[0]);
- ptr = build_gep_for_deref(ctx, instr->variables[0]);
+ LLVMValueRef src = get_src(ctx->nir, instr->src[0]);
+ ptr = build_gep_for_deref(ctx->nir, instr->variables[0]);
if (instr->intrinsic == nir_intrinsic_var_atomic_comp_swap) {
- LLVMValueRef src1 = get_src(ctx, instr->src[1]);
+ LLVMValueRef src1 = get_src(ctx->nir, instr->src[1]);
result = LLVMBuildAtomicCmpXchg(ctx->builder,
ptr, src, src1,
LLVMAtomicOrderingSequentiallyConsistent,
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;
}
-static LLVMValueRef load_sample_pos(struct nir_to_llvm_context *ctx)
+static LLVMValueRef load_sample_pos(struct ac_nir_context *ctx)
{
LLVMValueRef values[2];
- values[0] = emit_ffract(&ctx->ac, ctx->frag_pos[0]);
- values[1] = emit_ffract(&ctx->ac, ctx->frag_pos[1]);
+ values[0] = emit_ffract(&ctx->ac, ctx->abi->frag_pos[0]);
+ values[1] = emit_ffract(&ctx->ac, ctx->abi->frag_pos[1]);
return ac_build_gather_values(&ctx->ac, values, 2);
}
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:
case nir_intrinsic_interp_var_at_sample:
case nir_intrinsic_interp_var_at_offset:
location = INTERP_CENTER;
- src0 = get_src(ctx, instr->src[0]);
+ src0 = get_src(ctx->nir, instr->src[0]);
break;
default:
break;
}
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, interp_param);
+ LLVMValueRef ddxy_out = emit_ddxy_interp(ctx->nir, interp_param);
/*
* take the I then J parameters, and the DDX/Y for it, and
/* loop num outputs */
idx = 0;
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
- LLVMValueRef *out_ptr = &ctx->outputs[i * 4];
+ LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4];
int length = 4;
int slot = idx;
int slot_inc = 1;
LLVMValueRef result = ac_build_gather_values(&ctx->ac, coord, instr->num_components);
return LLVMBuildBitCast(ctx->builder, result,
- get_def_type(ctx, &instr->dest.ssa), "");
+ get_def_type(ctx->nir, &instr->dest.ssa), "");
}
-static void visit_intrinsic(struct nir_to_llvm_context *ctx,
+static void visit_intrinsic(struct ac_nir_context *ctx,
nir_intrinsic_instr *instr)
{
LLVMValueRef result = NULL;
switch (instr->intrinsic) {
case nir_intrinsic_load_work_group_id: {
- result = ctx->workgroup_ids;
+ result = ctx->nctx->workgroup_ids;
break;
}
case nir_intrinsic_load_base_vertex: {
- result = ctx->abi.base_vertex;
+ result = ctx->abi->base_vertex;
break;
}
case nir_intrinsic_load_vertex_id_zero_base: {
- result = ctx->abi.vertex_id;
+ result = ctx->abi->vertex_id;
break;
}
case nir_intrinsic_load_local_invocation_id: {
- result = ctx->local_invocation_ids;
+ result = ctx->nctx->local_invocation_ids;
break;
}
case nir_intrinsic_load_base_instance:
- result = ctx->abi.start_instance;
+ result = ctx->abi->start_instance;
break;
case nir_intrinsic_load_draw_id:
- result = ctx->abi.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, ctx->tcs_rel_ids, 8, 5);
+ result = unpack_param(&ctx->ac, ctx->nctx->tcs_rel_ids, 8, 5);
else
- result = ctx->gs_invocation_id;
+ result = ctx->nctx->gs_invocation_id;
break;
case nir_intrinsic_load_primitive_id:
if (ctx->stage == MESA_SHADER_GEOMETRY) {
- ctx->shader_info->gs.uses_prim_id = true;
- result = ctx->gs_prim_id;
+ ctx->nctx->shader_info->gs.uses_prim_id = true;
+ result = ctx->nctx->gs_prim_id;
} else if (ctx->stage == MESA_SHADER_TESS_CTRL) {
- ctx->shader_info->tcs.uses_prim_id = true;
- result = ctx->tcs_patch_id;
+ ctx->nctx->shader_info->tcs.uses_prim_id = true;
+ result = ctx->nctx->tcs_patch_id;
} else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
- ctx->shader_info->tcs.uses_prim_id = true;
- result = ctx->tes_patch_id;
+ ctx->nctx->shader_info->tcs.uses_prim_id = true;
+ result = ctx->nctx->tes_patch_id;
} else
fprintf(stderr, "Unknown primitive id intrinsic: %d", ctx->stage);
break;
case nir_intrinsic_load_sample_id:
- ctx->shader_info->fs.force_persample = true;
- result = unpack_param(ctx, ctx->ancillary, 8, 4);
+ result = unpack_param(&ctx->ac, ctx->abi->ancillary, 8, 4);
break;
case nir_intrinsic_load_sample_pos:
- ctx->shader_info->fs.force_persample = true;
result = load_sample_pos(ctx);
break;
case nir_intrinsic_load_sample_mask_in:
- result = ctx->sample_coverage;
+ result = ctx->abi->sample_coverage;
+ break;
+ case nir_intrinsic_load_frag_coord: {
+ LLVMValueRef values[4] = {
+ ctx->abi->frag_pos[0],
+ ctx->abi->frag_pos[1],
+ ctx->abi->frag_pos[2],
+ ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, ctx->abi->frag_pos[3])
+ };
+ result = ac_build_gather_values(&ctx->ac, values, 4);
break;
+ }
case nir_intrinsic_load_front_face:
- result = ctx->front_face;
+ result = ctx->abi->front_face;
break;
case nir_intrinsic_load_instance_id:
- result = ctx->abi.instance_id;
- ctx->shader_info->vs.vgpr_comp_cnt = MAX2(3,
- ctx->shader_info->vs.vgpr_comp_cnt);
+ result = ctx->abi->instance_id;
break;
case nir_intrinsic_load_num_work_groups:
- result = ctx->num_work_groups;
+ result = ctx->nctx->num_work_groups;
break;
case nir_intrinsic_load_local_invocation_index:
- result = visit_load_local_invocation_index(ctx);
+ result = visit_load_local_invocation_index(ctx->nctx);
break;
case nir_intrinsic_load_push_constant:
- result = visit_load_push_constant(ctx, instr);
+ result = visit_load_push_constant(ctx->nctx, instr);
break;
case nir_intrinsic_vulkan_resource_index:
- result = visit_vulkan_resource_index(ctx, instr);
+ result = visit_vulkan_resource_index(ctx->nctx, instr);
break;
case nir_intrinsic_store_ssbo:
visit_store_ssbo(ctx, instr);
result = visit_image_size(ctx, instr);
break;
case nir_intrinsic_discard:
- ctx->shader_info->fs.can_discard = true;
ac_build_intrinsic(&ctx->ac, "llvm.AMDGPU.kilp",
- ctx->voidt,
+ LLVMVoidTypeInContext(ctx->ac.context),
NULL, 0, AC_FUNC_ATTR_LEGACY);
break;
case nir_intrinsic_discard_if:
emit_discard_if(ctx, instr);
break;
case nir_intrinsic_memory_barrier:
- emit_waitcnt(ctx, VM_CNT);
+ emit_waitcnt(ctx->nctx, VM_CNT);
break;
case nir_intrinsic_barrier:
- emit_barrier(ctx);
+ emit_barrier(ctx->nctx);
break;
case nir_intrinsic_var_atomic_add:
case nir_intrinsic_var_atomic_imin:
case nir_intrinsic_var_atomic_xor:
case nir_intrinsic_var_atomic_exchange:
case nir_intrinsic_var_atomic_comp_swap:
- result = visit_var_atomic(ctx, instr);
+ result = visit_var_atomic(ctx->nctx, instr);
break;
case nir_intrinsic_interp_var_at_centroid:
case nir_intrinsic_interp_var_at_sample:
case nir_intrinsic_interp_var_at_offset:
- result = visit_interp(ctx, instr);
+ result = visit_interp(ctx->nctx, instr);
break;
case nir_intrinsic_emit_vertex:
- visit_emit_vertex(ctx, instr);
+ visit_emit_vertex(ctx->nctx, instr);
break;
case nir_intrinsic_end_primitive:
- visit_end_primitive(ctx, instr);
+ visit_end_primitive(ctx->nctx, instr);
break;
case nir_intrinsic_load_tess_coord:
- result = visit_load_tess_coord(ctx, instr);
+ result = visit_load_tess_coord(ctx->nctx, instr);
break;
case nir_intrinsic_load_patch_vertices_in:
- result = LLVMConstInt(ctx->i32, ctx->options->key.tcs.input_vertices, false);
+ result = LLVMConstInt(ctx->ac.i32, ctx->nctx->options->key.tcs.input_vertices, false);
break;
default:
fprintf(stderr, "Unknown intrinsic: ");
}
}
-static LLVMValueRef get_sampler_desc(struct nir_to_llvm_context *ctx,
- const nir_deref_var *deref,
- enum desc_type desc_type)
+static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
+ LLVMValueRef buffer, bool write)
+{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+
+ if (write && ctx->stage == MESA_SHADER_FRAGMENT)
+ ctx->shader_info->fs.writes_memory = true;
+
+ return buffer;
+}
+
+static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
+ unsigned descriptor_set,
+ unsigned base_index,
+ unsigned constant_index,
+ LLVMValueRef index,
+ enum ac_descriptor_type desc_type,
+ bool image, bool write)
{
- unsigned desc_set = deref->var->data.descriptor_set;
- LLVMValueRef list = ctx->descriptor_sets[desc_set];
- struct radv_descriptor_set_layout *layout = ctx->options->layout->set[desc_set].layout;
- struct radv_descriptor_set_binding_layout *binding = layout->binding + deref->var->data.binding;
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+ LLVMValueRef list = ctx->descriptor_sets[descriptor_set];
+ struct radv_descriptor_set_layout *layout = ctx->options->layout->set[descriptor_set].layout;
+ struct radv_descriptor_set_binding_layout *binding = layout->binding + base_index;
unsigned offset = binding->offset;
unsigned stride = binding->size;
unsigned type_size;
LLVMBuilderRef builder = ctx->builder;
LLVMTypeRef type;
- LLVMValueRef index = NULL;
- unsigned constant_index = 0;
- assert(deref->var->data.binding < layout->binding_count);
+ assert(base_index < layout->binding_count);
+
+ if (write && ctx->stage == MESA_SHADER_FRAGMENT)
+ ctx->shader_info->fs.writes_memory = true;
switch (desc_type) {
- case DESC_IMAGE:
+ case AC_DESC_IMAGE:
type = ctx->v8i32;
type_size = 32;
break;
- case DESC_FMASK:
+ case AC_DESC_FMASK:
type = ctx->v8i32;
offset += 32;
type_size = 32;
break;
- case DESC_SAMPLER:
+ case AC_DESC_SAMPLER:
type = ctx->v4i32;
if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER)
offset += 64;
type_size = 16;
break;
- case DESC_BUFFER:
+ case AC_DESC_BUFFER:
type = ctx->v4i32;
type_size = 16;
break;
unreachable("invalid desc_type\n");
}
- if (deref->deref.child) {
- const nir_deref_array *child =
- (const nir_deref_array *)deref->deref.child;
-
- assert(child->deref_array_type != nir_deref_array_type_wildcard);
- offset += child->base_offset * stride;
- if (child->deref_array_type == nir_deref_array_type_indirect) {
- index = get_src(ctx, child->indirect);
- }
+ offset += constant_index * stride;
- constant_index = child->base_offset;
- }
- if (desc_type == DESC_SAMPLER && binding->immutable_samplers_offset &&
+ if (desc_type == AC_DESC_SAMPLER && binding->immutable_samplers_offset &&
(!index || binding->immutable_samplers_equal)) {
if (binding->immutable_samplers_equal)
constant_index = 0;
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,
+ const nir_deref_var *deref,
+ enum ac_descriptor_type desc_type,
+ bool image, bool write)
+{
+ LLVMValueRef index = NULL;
+ unsigned constant_index = 0;
+ const nir_deref *tail = &deref->deref;
+
+ while (tail->child) {
+ const nir_deref_array *child = nir_deref_as_array(tail->child);
+ unsigned array_size = glsl_get_aoa_size(tail->child->type);
+
+ if (!array_size)
+ array_size = 1;
+
+ assert(child->deref_array_type != nir_deref_array_type_wildcard);
+
+ if (child->deref_array_type == nir_deref_array_type_indirect) {
+ LLVMValueRef indirect = get_src(ctx, child->indirect);
+
+ indirect = LLVMBuildMul(ctx->ac.builder, indirect,
+ LLVMConstInt(ctx->ac.i32, array_size, false), "");
+
+ if (!index)
+ index = indirect;
+ else
+ index = LLVMBuildAdd(ctx->ac.builder, index, indirect, "");
+ }
+
+ constant_index += child->base_offset * array_size;
+
+ tail = &child->deref;
+ }
+
+ return ctx->abi->load_sampler_desc(ctx->abi,
+ deref->var->data.descriptor_set,
+ deref->var->data.binding,
+ constant_index, index,
+ desc_type, image, write);
}
-static void set_tex_fetch_args(struct nir_to_llvm_context *ctx,
+static void set_tex_fetch_args(struct ac_llvm_context *ctx,
struct ac_image_args *args,
const nir_tex_instr *instr,
nir_texop op,
param[count++] = LLVMGetUndef(ctx->i32);
if (count > 1)
- args->addr = ac_build_gather_values(&ctx->ac, param, count);
+ args->addr = ac_build_gather_values(ctx, param, count);
else
args->addr = param[0];
* VI:
* The ANISO_OVERRIDE sampler field enables this fix in TA.
*/
-static LLVMValueRef sici_fix_sampler_aniso(struct nir_to_llvm_context *ctx,
+static LLVMValueRef sici_fix_sampler_aniso(struct ac_nir_context *ctx,
LLVMValueRef res, LLVMValueRef samp)
{
- LLVMBuilderRef builder = ctx->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef img7, samp0;
- if (ctx->options->chip_class >= VI)
+ if (ctx->ac.chip_class >= VI)
return samp;
img7 = LLVMBuildExtractElement(builder, res,
- LLVMConstInt(ctx->i32, 7, 0), "");
+ LLVMConstInt(ctx->ac.i32, 7, 0), "");
samp0 = LLVMBuildExtractElement(builder, samp,
- LLVMConstInt(ctx->i32, 0, 0), "");
+ LLVMConstInt(ctx->ac.i32, 0, 0), "");
samp0 = LLVMBuildAnd(builder, samp0, img7, "");
return LLVMBuildInsertElement(builder, samp, samp0,
- LLVMConstInt(ctx->i32, 0, 0), "");
+ LLVMConstInt(ctx->ac.i32, 0, 0), "");
}
-static void tex_fetch_ptrs(struct nir_to_llvm_context *ctx,
+static void tex_fetch_ptrs(struct ac_nir_context *ctx,
nir_tex_instr *instr,
LLVMValueRef *res_ptr, LLVMValueRef *samp_ptr,
LLVMValueRef *fmask_ptr)
{
if (instr->sampler_dim == GLSL_SAMPLER_DIM_BUF)
- *res_ptr = get_sampler_desc(ctx, instr->texture, DESC_BUFFER);
+ *res_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_BUFFER, false, false);
else
- *res_ptr = get_sampler_desc(ctx, instr->texture, DESC_IMAGE);
+ *res_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_IMAGE, false, false);
if (samp_ptr) {
if (instr->sampler)
- *samp_ptr = get_sampler_desc(ctx, instr->sampler, DESC_SAMPLER);
+ *samp_ptr = get_sampler_desc(ctx, instr->sampler, AC_DESC_SAMPLER, false, false);
else
- *samp_ptr = get_sampler_desc(ctx, instr->texture, DESC_SAMPLER);
+ *samp_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_SAMPLER, false, false);
if (instr->sampler_dim < GLSL_SAMPLER_DIM_RECT)
*samp_ptr = sici_fix_sampler_aniso(ctx, *res_ptr, *samp_ptr);
}
if (fmask_ptr && !instr->sampler && (instr->op == nir_texop_txf_ms ||
instr->op == nir_texop_samples_identical))
- *fmask_ptr = get_sampler_desc(ctx, instr->texture, DESC_FMASK);
+ *fmask_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_FMASK, false, false);
}
-static LLVMValueRef apply_round_slice(struct nir_to_llvm_context *ctx,
+static LLVMValueRef apply_round_slice(struct ac_llvm_context *ctx,
LLVMValueRef coord)
{
- coord = to_float(&ctx->ac, coord);
- coord = ac_build_intrinsic(&ctx->ac, "llvm.rint.f32", ctx->f32, &coord, 1, 0);
- coord = to_integer(&ctx->ac, coord);
+ coord = ac_to_float(ctx, coord);
+ coord = ac_build_intrinsic(ctx, "llvm.rint.f32", ctx->f32, &coord, 1, 0);
+ coord = ac_to_integer(ctx, coord);
return coord;
}
-static void visit_tex(struct nir_to_llvm_context *ctx, nir_tex_instr *instr)
+static void visit_tex(struct ac_nir_context *ctx, nir_tex_instr *instr)
{
LLVMValueRef result = NULL;
struct ac_image_args args = { 0 };
unsigned chan, count = 0;
unsigned const_src = 0, num_deriv_comp = 0;
bool lod_is_zero = false;
+
tex_fetch_ptrs(ctx, instr, &res_ptr, &samp_ptr, &fmask_ptr);
for (unsigned i = 0; i < instr->num_srcs; i++) {
if (instr->op == nir_texop_texture_samples) {
LLVMValueRef res, samples, is_msaa;
- res = LLVMBuildBitCast(ctx->builder, res_ptr, ctx->v8i32, "");
- samples = LLVMBuildExtractElement(ctx->builder, res,
- LLVMConstInt(ctx->i32, 3, false), "");
- is_msaa = LLVMBuildLShr(ctx->builder, samples,
- LLVMConstInt(ctx->i32, 28, false), "");
- is_msaa = LLVMBuildAnd(ctx->builder, is_msaa,
- LLVMConstInt(ctx->i32, 0xe, false), "");
- is_msaa = LLVMBuildICmp(ctx->builder, LLVMIntEQ, is_msaa,
- LLVMConstInt(ctx->i32, 0xe, false), "");
-
- samples = LLVMBuildLShr(ctx->builder, samples,
- LLVMConstInt(ctx->i32, 16, false), "");
- samples = LLVMBuildAnd(ctx->builder, samples,
- LLVMConstInt(ctx->i32, 0xf, false), "");
- samples = LLVMBuildShl(ctx->builder, ctx->i32one,
+ res = LLVMBuildBitCast(ctx->ac.builder, res_ptr, ctx->ac.v8i32, "");
+ samples = LLVMBuildExtractElement(ctx->ac.builder, res,
+ LLVMConstInt(ctx->ac.i32, 3, false), "");
+ is_msaa = LLVMBuildLShr(ctx->ac.builder, samples,
+ LLVMConstInt(ctx->ac.i32, 28, false), "");
+ is_msaa = LLVMBuildAnd(ctx->ac.builder, is_msaa,
+ LLVMConstInt(ctx->ac.i32, 0xe, false), "");
+ is_msaa = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, is_msaa,
+ LLVMConstInt(ctx->ac.i32, 0xe, false), "");
+
+ samples = LLVMBuildLShr(ctx->ac.builder, samples,
+ LLVMConstInt(ctx->ac.i32, 16, false), "");
+ samples = LLVMBuildAnd(ctx->ac.builder, samples,
+ LLVMConstInt(ctx->ac.i32, 0xf, false), "");
+ samples = LLVMBuildShl(ctx->ac.builder, ctx->ac.i32_1,
samples, "");
- samples = LLVMBuildSelect(ctx->builder, is_msaa, samples,
- ctx->i32one, "");
+ samples = LLVMBuildSelect(ctx->ac.builder, is_msaa, samples,
+ ctx->ac.i32_1, "");
result = samples;
goto write_result;
}
if (coord)
for (chan = 0; chan < instr->coord_components; chan++)
- coords[chan] = llvm_extract_elem(ctx, coord, chan);
+ coords[chan] = llvm_extract_elem(&ctx->ac, coord, chan);
if (offsets && instr->op != nir_texop_txf) {
LLVMValueRef offset[3], pack;
for (chan = 0; chan < 3; ++chan)
- offset[chan] = ctx->i32zero;
+ offset[chan] = ctx->ac.i32_0;
args.offset = true;
for (chan = 0; chan < get_llvm_num_components(offsets); chan++) {
- offset[chan] = llvm_extract_elem(ctx, offsets, chan);
- offset[chan] = LLVMBuildAnd(ctx->builder, offset[chan],
- LLVMConstInt(ctx->i32, 0x3f, false), "");
+ offset[chan] = llvm_extract_elem(&ctx->ac, offsets, chan);
+ offset[chan] = LLVMBuildAnd(ctx->ac.builder, offset[chan],
+ LLVMConstInt(ctx->ac.i32, 0x3f, false), "");
if (chan)
- offset[chan] = LLVMBuildShl(ctx->builder, offset[chan],
- LLVMConstInt(ctx->i32, chan * 8, false), "");
+ offset[chan] = LLVMBuildShl(ctx->ac.builder, offset[chan],
+ LLVMConstInt(ctx->ac.i32, chan * 8, false), "");
}
- pack = LLVMBuildOr(ctx->builder, offset[0], offset[1], "");
- pack = LLVMBuildOr(ctx->builder, pack, offset[2], "");
+ pack = LLVMBuildOr(ctx->ac.builder, offset[0], offset[1], "");
+ pack = LLVMBuildOr(ctx->ac.builder, pack, offset[2], "");
address[count++] = pack;
}
/* Pack depth comparison value */
if (instr->is_shadow && comparator) {
- address[count++] = llvm_extract_elem(ctx, 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
+ * Z24 anymore. Do it manually here.
+ *
+ * It's unnecessary if the original texture format was
+ * Z32_FLOAT, but we don't know that here.
+ */
+ 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, ddx, i));
- derivs[num_deriv_comp + i] = to_float(&ctx->ac, llvm_extract_elem(ctx, 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, 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->f32);
+ 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[0];
if (instr->coord_components > 1) {
if (instr->sampler_dim == GLSL_SAMPLER_DIM_1D && instr->is_array && instr->op != nir_texop_txf) {
- coords[1] = apply_round_slice(ctx, coords[1]);
+ coords[1] = apply_round_slice(&ctx->ac, coords[1]);
}
address[count++] = coords[1];
}
if (instr->sampler_dim != GLSL_SAMPLER_DIM_3D &&
instr->sampler_dim != GLSL_SAMPLER_DIM_CUBE &&
instr->op != nir_texop_txf) {
- coords[2] = apply_round_slice(ctx, coords[2]);
+ coords[2] = apply_round_slice(&ctx->ac, coords[2]);
}
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 */
if (lod)
address[count++] = lod;
else
- address[count++] = ctx->i32zero;
+ address[count++] = ctx->ac.i32_0;
}
for (chan = 0; chan < count; chan++) {
- address[chan] = LLVMBuildBitCast(ctx->builder,
- address[chan], ctx->i32, "");
+ address[chan] = LLVMBuildBitCast(ctx->ac.builder,
+ address[chan], ctx->ac.i32, "");
}
if (instr->op == nir_texop_samples_identical) {
memcpy(txf_address, address, sizeof(txf_address));
if (!instr->is_array)
- txf_address[2] = ctx->i32zero;
- txf_address[3] = ctx->i32zero;
+ txf_address[2] = ctx->ac.i32_0;
+ txf_address[3] = ctx->ac.i32_0;
- set_tex_fetch_args(ctx, &txf_args, instr, nir_texop_txf,
+ set_tex_fetch_args(&ctx->ac, &txf_args, instr, nir_texop_txf,
fmask_ptr, NULL,
txf_address, txf_count, 0xf);
result = build_tex_intrinsic(ctx, instr, false, &txf_args);
- result = LLVMBuildExtractElement(ctx->builder, result, ctx->i32zero, "");
- result = emit_int_cmp(&ctx->ac, LLVMIntEQ, result, ctx->i32zero);
+ result = LLVMBuildExtractElement(ctx->ac.builder, result, ctx->ac.i32_0, "");
+ result = emit_int_cmp(&ctx->ac, LLVMIntEQ, result, ctx->ac.i32_0);
goto write_result;
}
if (instr->sampler_dim == GLSL_SAMPLER_DIM_MS &&
instr->op != nir_texop_txs) {
unsigned sample_chan = instr->is_array ? 3 : 2;
- address[sample_chan] = adjust_sample_index_using_fmask(ctx,
+ address[sample_chan] = adjust_sample_index_using_fmask(&ctx->ac,
address[0],
address[1],
instr->is_array ? address[2] : NULL,
assert(const_offset);
num_offsets = MIN2(num_offsets, instr->coord_components);
if (num_offsets > 2)
- address[2] = LLVMBuildAdd(ctx->builder,
- address[2], LLVMConstInt(ctx->i32, const_offset->i32[2], false), "");
+ address[2] = LLVMBuildAdd(ctx->ac.builder,
+ address[2], LLVMConstInt(ctx->ac.i32, const_offset->i32[2], false), "");
if (num_offsets > 1)
- address[1] = LLVMBuildAdd(ctx->builder,
- address[1], LLVMConstInt(ctx->i32, const_offset->i32[1], false), "");
- address[0] = LLVMBuildAdd(ctx->builder,
- address[0], LLVMConstInt(ctx->i32, const_offset->i32[0], false), "");
+ address[1] = LLVMBuildAdd(ctx->ac.builder,
+ address[1], LLVMConstInt(ctx->ac.i32, const_offset->i32[1], false), "");
+ address[0] = LLVMBuildAdd(ctx->ac.builder,
+ address[0], LLVMConstInt(ctx->ac.i32, const_offset->i32[0], false), "");
}
else
dmask = 1 << instr->component;
}
- set_tex_fetch_args(ctx, &args, instr, instr->op,
+ set_tex_fetch_args(&ctx->ac, &args, instr, instr->op,
res_ptr, samp_ptr, address, count, dmask);
result = build_tex_intrinsic(ctx, instr, lod_is_zero, &args);
if (instr->op == nir_texop_query_levels)
- result = LLVMBuildExtractElement(ctx->builder, result, LLVMConstInt(ctx->i32, 3, false), "");
- else if (instr->is_shadow && instr->op != nir_texop_txs && instr->op != nir_texop_lod && instr->op != nir_texop_tg4)
- result = LLVMBuildExtractElement(ctx->builder, result, ctx->i32zero, "");
+ result = LLVMBuildExtractElement(ctx->ac.builder, result, LLVMConstInt(ctx->ac.i32, 3, false), "");
+ else if (instr->is_shadow && instr->is_new_style_shadow &&
+ instr->op != nir_texop_txs && instr->op != nir_texop_lod &&
+ instr->op != nir_texop_tg4)
+ result = LLVMBuildExtractElement(ctx->ac.builder, result, ctx->ac.i32_0, "");
else if (instr->op == nir_texop_txs &&
instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE &&
instr->is_array) {
- LLVMValueRef two = LLVMConstInt(ctx->i32, 2, false);
- LLVMValueRef six = LLVMConstInt(ctx->i32, 6, false);
- LLVMValueRef z = LLVMBuildExtractElement(ctx->builder, result, two, "");
- z = LLVMBuildSDiv(ctx->builder, z, six, "");
- result = LLVMBuildInsertElement(ctx->builder, result, z, two, "");
+ LLVMValueRef two = LLVMConstInt(ctx->ac.i32, 2, false);
+ LLVMValueRef six = LLVMConstInt(ctx->ac.i32, 6, false);
+ 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, result, instr->dest.ssa.num_components);
+ 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);
}
}
-static void visit_phi(struct nir_to_llvm_context *ctx, nir_phi_instr *instr)
+static void visit_phi(struct ac_nir_context *ctx, nir_phi_instr *instr)
{
LLVMTypeRef type = get_def_type(ctx, &instr->dest.ssa);
- LLVMValueRef result = LLVMBuildPhi(ctx->builder, type, "");
+ LLVMValueRef result = LLVMBuildPhi(ctx->ac.builder, type, "");
_mesa_hash_table_insert(ctx->defs, &instr->dest.ssa, result);
_mesa_hash_table_insert(ctx->phis, instr, result);
}
-static void visit_post_phi(struct nir_to_llvm_context *ctx,
+static void visit_post_phi(struct ac_nir_context *ctx,
nir_phi_instr *instr,
LLVMValueRef llvm_phi)
{
}
}
-static void phi_post_pass(struct nir_to_llvm_context *ctx)
+static void phi_post_pass(struct ac_nir_context *ctx)
{
struct hash_entry *entry;
hash_table_foreach(ctx->phis, entry) {
}
-static void visit_ssa_undef(struct nir_to_llvm_context *ctx,
+static void visit_ssa_undef(struct ac_nir_context *ctx,
const nir_ssa_undef_instr *instr)
{
unsigned num_components = instr->def.num_components;
LLVMValueRef undef;
if (num_components == 1)
- undef = LLVMGetUndef(ctx->i32);
+ undef = LLVMGetUndef(ctx->ac.i32);
else {
- undef = LLVMGetUndef(LLVMVectorType(ctx->i32, num_components));
+ undef = LLVMGetUndef(LLVMVectorType(ctx->ac.i32, num_components));
}
_mesa_hash_table_insert(ctx->defs, &instr->def, undef);
}
-static void visit_jump(struct nir_to_llvm_context *ctx,
+static void visit_jump(struct ac_nir_context *ctx,
const nir_jump_instr *instr)
{
switch (instr->type) {
case nir_jump_break:
- LLVMBuildBr(ctx->builder, ctx->break_block);
- LLVMClearInsertionPosition(ctx->builder);
+ LLVMBuildBr(ctx->ac.builder, ctx->break_block);
+ LLVMClearInsertionPosition(ctx->ac.builder);
break;
case nir_jump_continue:
- LLVMBuildBr(ctx->builder, ctx->continue_block);
- LLVMClearInsertionPosition(ctx->builder);
+ LLVMBuildBr(ctx->ac.builder, ctx->continue_block);
+ LLVMClearInsertionPosition(ctx->ac.builder);
break;
default:
fprintf(stderr, "Unknown NIR jump instr: ");
}
}
-static void visit_cf_list(struct nir_to_llvm_context *ctx,
+static void visit_cf_list(struct ac_nir_context *ctx,
struct exec_list *list);
-static void visit_block(struct nir_to_llvm_context *ctx, nir_block *block)
+static void visit_block(struct ac_nir_context *ctx, nir_block *block)
{
- LLVMBasicBlockRef llvm_block = LLVMGetInsertBlock(ctx->builder);
+ LLVMBasicBlockRef llvm_block = LLVMGetInsertBlock(ctx->ac.builder);
nir_foreach_instr(instr, block)
{
switch (instr->type) {
_mesa_hash_table_insert(ctx->defs, block, llvm_block);
}
-static void visit_if(struct nir_to_llvm_context *ctx, nir_if *if_stmt)
+static void visit_if(struct ac_nir_context *ctx, nir_if *if_stmt)
{
LLVMValueRef value = get_src(ctx, if_stmt->condition);
+ LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx->ac.builder));
LLVMBasicBlockRef merge_block =
- LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, "");
+ LLVMAppendBasicBlockInContext(ctx->ac.context, fn, "");
LLVMBasicBlockRef if_block =
- LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, "");
+ LLVMAppendBasicBlockInContext(ctx->ac.context, fn, "");
LLVMBasicBlockRef else_block = merge_block;
if (!exec_list_is_empty(&if_stmt->else_list))
else_block = LLVMAppendBasicBlockInContext(
- ctx->context, ctx->main_function, "");
+ ctx->ac.context, fn, "");
- LLVMValueRef cond = LLVMBuildICmp(ctx->builder, LLVMIntNE, value,
- LLVMConstInt(ctx->i32, 0, false), "");
- LLVMBuildCondBr(ctx->builder, cond, if_block, else_block);
+ LLVMValueRef cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE, value,
+ LLVMConstInt(ctx->ac.i32, 0, false), "");
+ LLVMBuildCondBr(ctx->ac.builder, cond, if_block, else_block);
- LLVMPositionBuilderAtEnd(ctx->builder, if_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, if_block);
visit_cf_list(ctx, &if_stmt->then_list);
- if (LLVMGetInsertBlock(ctx->builder))
- LLVMBuildBr(ctx->builder, merge_block);
+ if (LLVMGetInsertBlock(ctx->ac.builder))
+ LLVMBuildBr(ctx->ac.builder, merge_block);
if (!exec_list_is_empty(&if_stmt->else_list)) {
- LLVMPositionBuilderAtEnd(ctx->builder, else_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, else_block);
visit_cf_list(ctx, &if_stmt->else_list);
- if (LLVMGetInsertBlock(ctx->builder))
- LLVMBuildBr(ctx->builder, merge_block);
+ if (LLVMGetInsertBlock(ctx->ac.builder))
+ LLVMBuildBr(ctx->ac.builder, merge_block);
}
- LLVMPositionBuilderAtEnd(ctx->builder, merge_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, merge_block);
}
-static void visit_loop(struct nir_to_llvm_context *ctx, nir_loop *loop)
+static void visit_loop(struct ac_nir_context *ctx, nir_loop *loop)
{
+ LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx->ac.builder));
LLVMBasicBlockRef continue_parent = ctx->continue_block;
LLVMBasicBlockRef break_parent = ctx->break_block;
ctx->continue_block =
- LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, "");
+ LLVMAppendBasicBlockInContext(ctx->ac.context, fn, "");
ctx->break_block =
- LLVMAppendBasicBlockInContext(ctx->context, ctx->main_function, "");
+ LLVMAppendBasicBlockInContext(ctx->ac.context, fn, "");
- LLVMBuildBr(ctx->builder, ctx->continue_block);
- LLVMPositionBuilderAtEnd(ctx->builder, ctx->continue_block);
+ LLVMBuildBr(ctx->ac.builder, ctx->continue_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, ctx->continue_block);
visit_cf_list(ctx, &loop->body);
- if (LLVMGetInsertBlock(ctx->builder))
- LLVMBuildBr(ctx->builder, ctx->continue_block);
- LLVMPositionBuilderAtEnd(ctx->builder, ctx->break_block);
+ if (LLVMGetInsertBlock(ctx->ac.builder))
+ LLVMBuildBr(ctx->ac.builder, ctx->continue_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, ctx->break_block);
ctx->continue_block = continue_parent;
ctx->break_block = break_parent;
}
-static void visit_cf_list(struct nir_to_llvm_context *ctx,
+static void visit_cf_list(struct ac_nir_context *ctx,
struct exec_list *list)
{
foreach_list_typed(nir_cf_node, node, node, list)
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, ""));
}
}
unsigned interp_type;
if (variable->data.sample) {
interp_type = INTERP_SAMPLE;
- ctx->shader_info->fs.force_persample = true;
+ ctx->shader_info->info.ps.force_persample = true;
} else if (variable->data.centroid)
interp_type = INTERP_CENTROID;
else
}
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);
++index;
} else if (i == VARYING_SLOT_POS) {
for(int i = 0; i < 3; ++i)
- inputs[i] = ctx->frag_pos[i];
+ inputs[i] = ctx->abi.frag_pos[i];
- inputs[3] = ac_build_fdiv(&ctx->ac, ctx->f32one, ctx->frag_pos[3]);
+ inputs[3] = ac_build_fdiv(&ctx->ac, ctx->f32one,
+ ctx->abi.frag_pos[3]);
}
}
ctx->shader_info->fs.num_interp = index;
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
-ac_build_alloca(struct nir_to_llvm_context *ctx,
+ac_build_alloca(struct ac_llvm_context *ac,
LLVMTypeRef type,
const char *name)
{
- LLVMBuilderRef builder = ctx->builder;
+ LLVMBuilderRef builder = ac->builder;
LLVMBasicBlockRef current_block = LLVMGetInsertBlock(builder);
LLVMValueRef function = LLVMGetBasicBlockParent(current_block);
LLVMBasicBlockRef first_block = LLVMGetEntryBasicBlock(function);
LLVMValueRef first_instr = LLVMGetFirstInstruction(first_block);
- LLVMBuilderRef first_builder = LLVMCreateBuilderInContext(ctx->context);
+ LLVMBuilderRef first_builder = LLVMCreateBuilderInContext(ac->context);
LLVMValueRef res;
if (first_instr) {
return res;
}
-static LLVMValueRef si_build_alloca_undef(struct nir_to_llvm_context *ctx,
+static LLVMValueRef si_build_alloca_undef(struct ac_llvm_context *ac,
LLVMTypeRef type,
const char *name)
{
- LLVMValueRef ptr = ac_build_alloca(ctx, type, name);
- LLVMBuildStore(ctx->builder, LLVMGetUndef(type), ptr);
+ LLVMValueRef ptr = ac_build_alloca(ac, type, name);
+ LLVMBuildStore(ac->builder, LLVMGetUndef(type), ptr);
return ptr;
}
static void
-handle_shader_output_decl(struct nir_to_llvm_context *ctx,
- struct nir_variable *variable)
+scan_shader_output_decl(struct nir_to_llvm_context *ctx,
+ 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);
uint64_t mask_attribs;
+
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)
}
}
+ ctx->output_mask |= mask_attribs;
+}
+
+static void
+handle_shader_output_decl(struct ac_nir_context *ctx,
+ struct nir_shader *nir,
+ struct nir_variable *variable)
+{
+ unsigned output_loc = variable->data.driver_location / 4;
+ unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
+
+ /* tess ctrl has it's own load/store paths for outputs */
+ if (ctx->stage == MESA_SHADER_TESS_CTRL)
+ return;
+
+ if (ctx->stage == MESA_SHADER_VERTEX ||
+ ctx->stage == MESA_SHADER_TESS_EVAL ||
+ ctx->stage == MESA_SHADER_GEOMETRY) {
+ int idx = variable->data.location + variable->data.index;
+ if (idx == VARYING_SLOT_CLIP_DIST0) {
+ int length = nir->info.clip_distance_array_size +
+ nir->info.cull_distance_array_size;
+
+ if (length > 4)
+ attrib_count = 2;
+ else
+ attrib_count = 1;
+ }
+ }
+
for (unsigned i = 0; i < attrib_count; ++i) {
for (unsigned chan = 0; chan < 4; chan++) {
- ctx->outputs[radeon_llvm_reg_index_soa(idx + i, chan)] =
- si_build_alloca_undef(ctx, ctx->f32, "");
+ ctx->outputs[radeon_llvm_reg_index_soa(output_loc + i, chan)] =
+ si_build_alloca_undef(&ctx->ac, ctx->ac.f32, "");
}
}
- ctx->output_mask |= mask_attribs;
}
static LLVMTypeRef
}
static void
-setup_locals(struct nir_to_llvm_context *ctx,
+setup_locals(struct ac_nir_context *ctx,
struct nir_function *func)
{
int i, j;
for (i = 0; i < ctx->num_locals; i++) {
for (j = 0; j < 4; j++) {
ctx->locals[i * 4 + j] =
- si_build_alloca_undef(ctx, ctx->f32, "temp");
+ si_build_alloca_undef(&ctx->ac, ctx->ac.f32, "temp");
}
}
}
static void
-setup_shared(struct nir_to_llvm_context *ctx,
+setup_shared(struct ac_nir_context *ctx,
struct nir_shader *nir)
{
nir_foreach_variable(variable, &nir->shared) {
LLVMValueRef shared =
LLVMAddGlobalInAddressSpace(
- ctx->module, glsl_to_llvm_type(ctx, variable->type),
+ ctx->ac.module, glsl_to_llvm_type(ctx->nctx, variable->type),
variable->name ? variable->name : "",
LOCAL_ADDR_SPACE);
_mesa_hash_table_insert(ctx->vars, variable, shared);
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));
}
unsigned index = target - V_008DFC_SQ_EXP_MRT;
unsigned col_format = (ctx->options->key.fs.col_format >> (4 * index)) & 0xf;
bool is_int8 = (ctx->options->key.fs.is_int8 >> index) & 1;
+ bool is_int10 = (ctx->options->key.fs.is_int10 >> index) & 1;
switch(col_format) {
case V_028714_SPI_SHADER_ZERO:
break;
case V_028714_SPI_SHADER_UINT16_ABGR: {
- LLVMValueRef max = LLVMConstInt(ctx->i32, is_int8 ? 255 : 65535, 0);
+ LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
+ is_int8 ? 255 : is_int10 ? 1023 : 65535, 0);
+ 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] = emit_minmax_int(&ctx->ac, LLVMIntULT, val[chan], max);
+ 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);
}
args->compr = 1;
}
case V_028714_SPI_SHADER_SINT16_ABGR: {
- LLVMValueRef max = LLVMConstInt(ctx->i32, is_int8 ? 127 : 32767, 0);
- LLVMValueRef min = LLVMConstInt(ctx->i32, is_int8 ? -128 : -32768, 0);
+ LLVMValueRef max_rgb = LLVMConstInt(ctx->i32,
+ is_int8 ? 127 : is_int10 ? 511 : 32767, 0);
+ LLVMValueRef min_rgb = LLVMConstInt(ctx->i32,
+ is_int8 ? -128 : is_int10 ? -512 : -32768, 0);
+ LLVMValueRef max_alpha = !is_int10 ? max_rgb : ctx->i32one;
+ LLVMValueRef min_alpha = !is_int10 ? min_rgb : LLVMConstInt(ctx->i32, -2, 0);
/* Clamp. */
for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = to_integer(&ctx->ac, values[chan]);
- val[chan] = emit_minmax_int(&ctx->ac, LLVMIntSLT, val[chan], max);
- val[chan] = emit_minmax_int(&ctx->ac, LLVMIntSGT, val[chan], min);
+ 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);
}
args->compr = 1;
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,
- ctx->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
+ 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++)
slots[i] = LLVMGetUndef(ctx->f32);
if (ctx->output_mask & (1ull << VARYING_SLOT_POS)) {
for (unsigned j = 0; j < 4; j++)
pos_values[j] = LLVMBuildLoad(ctx->builder,
- ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_POS, j)], "");
+ ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_POS, j)], "");
}
si_llvm_init_export_args(ctx, pos_values, V_008DFC_SQ_EXP_POS, &pos_args[0]);
if (ctx->output_mask & (1ull << VARYING_SLOT_PSIZ)) {
outinfo->writes_pointsize = true;
psize_value = LLVMBuildLoad(ctx->builder,
- ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_PSIZ, 0)], "");
+ ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_PSIZ, 0)], "");
}
if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) {
outinfo->writes_layer = true;
layer_value = LLVMBuildLoad(ctx->builder,
- ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)], "");
+ ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)], "");
}
if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) {
outinfo->writes_viewport_index = true;
viewport_index_value = LLVMBuildLoad(ctx->builder,
- ctx->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_VIEWPORT, 0)], "");
+ 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->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;
{
int j;
uint64_t max_output_written = 0;
+ LLVMValueRef lds_base = NULL;
+
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
- LLVMValueRef *out_ptr = &ctx->outputs[i * 4];
int param_index;
int length = 4;
param_index = shader_io_get_unique_index(i);
max_output_written = MAX2(param_index + (length > 4), max_output_written);
+ }
+
+ outinfo->esgs_itemsize = (max_output_written + 1) * 16;
+
+ if (ctx->ac.chip_class >= GFX9) {
+ unsigned itemsize_dw = outinfo->esgs_itemsize / 4;
+ LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
+ LLVMValueRef wave_idx = ac_build_bfe(&ctx->ac, ctx->merged_wave_info,
+ LLVMConstInt(ctx->ac.i32, 24, false),
+ LLVMConstInt(ctx->ac.i32, 4, false), false);
+ vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
+ LLVMBuildMul(ctx->ac.builder, wave_idx,
+ LLVMConstInt(ctx->i32, 64, false), ""), "");
+ lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
+ LLVMConstInt(ctx->i32, itemsize_dw, 0), "");
+ }
+
+ for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
+ LLVMValueRef dw_addr;
+ LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4];
+ int param_index;
+ int length = 4;
+ if (!(ctx->output_mask & (1ull << i)))
+ continue;
+
+ if (i == VARYING_SLOT_CLIP_DIST0)
+ length = ctx->num_output_clips + ctx->num_output_culls;
+
+ param_index = shader_io_get_unique_index(i);
+
+ if (lds_base) {
+ dw_addr = LLVMBuildAdd(ctx->builder, lds_base,
+ LLVMConstInt(ctx->i32, param_index * 4, false),
+ "");
+ }
for (j = 0; j < length; j++) {
LLVMValueRef out_val = LLVMBuildLoad(ctx->builder, out_ptr[j], "");
out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->i32, "");
- ac_build_buffer_store_dword(&ctx->ac,
- ctx->esgs_ring,
- out_val, 1,
- NULL, ctx->es2gs_offset,
- (4 * param_index + j) * 4,
- 1, 1, true, true);
+ if (ctx->ac.chip_class >= GFX9) {
+ lds_store(ctx, dw_addr,
+ LLVMBuildLoad(ctx->builder, out_ptr[j], ""));
+ dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->i32one, "");
+ } else {
+ ac_build_buffer_store_dword(&ctx->ac,
+ ctx->esgs_ring,
+ out_val, 1,
+ NULL, ctx->es2gs_offset,
+ (4 * param_index + j) * 4,
+ 1, 1, true, true);
+ }
}
}
- outinfo->esgs_itemsize = (max_output_written + 1) * 16;
}
static void
handle_ls_outputs_post(struct nir_to_llvm_context *ctx)
{
LLVMValueRef vertex_id = ctx->rel_auto_id;
- LLVMValueRef vertex_dw_stride = unpack_param(ctx, ctx->ls_out_layout, 13, 8);
+ LLVMValueRef vertex_dw_stride = unpack_param(&ctx->ac, ctx->ls_out_layout, 13, 8);
LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->builder, vertex_id,
vertex_dw_stride, "");
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
- LLVMValueRef *out_ptr = &ctx->outputs[i * 4];
+ LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4];
int length = 4;
if (!(ctx->output_mask & (1ull << i)))
{
unsigned stride, outer_comps, inner_comps;
struct ac_build_if_state if_ctx, inner_if_ctx;
- LLVMValueRef invocation_id = unpack_param(ctx, ctx->tcs_rel_ids, 8, 5);
- LLVMValueRef rel_patch_id = unpack_param(ctx, ctx->tcs_rel_ids, 0, 8);
+ LLVMValueRef invocation_id = unpack_param(&ctx->ac, ctx->tcs_rel_ids, 8, 5);
+ LLVMValueRef rel_patch_id = unpack_param(&ctx->ac, ctx->tcs_rel_ids, 0, 8);
unsigned tess_inner_index, tess_outer_index;
LLVMValueRef lds_base, lds_inner, lds_outer, byteoffset, buffer;
LLVMValueRef out[6], vec0, vec1, tf_base, inner[4], outer[4];
tf_base = ctx->tess_factor_offset;
byteoffset = LLVMBuildMul(ctx->builder, rel_patch_id,
LLVMConstInt(ctx->i32, 4 * stride, false), "");
+ unsigned tf_offset = 0;
- ac_nir_build_if(&inner_if_ctx, ctx,
- LLVMBuildICmp(ctx->builder, LLVMIntEQ,
- rel_patch_id, ctx->i32zero, ""));
+ if (ctx->options->chip_class <= VI) {
+ ac_nir_build_if(&inner_if_ctx, ctx,
+ LLVMBuildICmp(ctx->builder, LLVMIntEQ,
+ rel_patch_id, ctx->i32zero, ""));
- /* Store the dynamic HS control word. */
- ac_build_buffer_store_dword(&ctx->ac, buffer,
- LLVMConstInt(ctx->i32, 0x80000000, false),
- 1, ctx->i32zero, tf_base,
- 0, 1, 0, true, false);
- ac_nir_build_endif(&inner_if_ctx);
+ /* Store the dynamic HS control word. */
+ ac_build_buffer_store_dword(&ctx->ac, buffer,
+ LLVMConstInt(ctx->i32, 0x80000000, false),
+ 1, ctx->i32zero, tf_base,
+ 0, 1, 0, true, false);
+ tf_offset += 4;
+
+ ac_nir_build_endif(&inner_if_ctx);
+ }
/* Store the tessellation factors. */
ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
MIN2(stride, 4), byteoffset, tf_base,
- 4, 1, 0, true, false);
+ tf_offset, 1, 0, true, false);
if (vec1)
ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
stride - 4, byteoffset, tf_base,
- 20, 1, 0, true, false);
+ 16 + tf_offset, 1, 0, true, false);
//TODO store to offchip for TES to read - only if TES reads them
if (1) {
if (i == FRAG_RESULT_DEPTH) {
ctx->shader_info->fs.writes_z = true;
- depth = to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
- ctx->outputs[radeon_llvm_reg_index_soa(i, 0)], ""));
+ 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,
- ctx->outputs[radeon_llvm_reg_index_soa(i, 0)], ""));
+ 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,
- ctx->outputs[radeon_llvm_reg_index_soa(i, 0)], ""));
+ 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,
- ctx->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 (!ctx->shader_info->fs.writes_z && !ctx->shader_info->fs.writes_stencil && !ctx->shader_info->fs.writes_sample_mask)
last = ctx->output_mask <= ((1ull << (i + 1)) - 1);
}
static void
-handle_shader_outputs_post(struct nir_to_llvm_context *ctx)
+handle_shader_outputs_post(struct ac_shader_abi *abi, unsigned max_outputs,
+ LLVMValueRef *addrs)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+
switch (ctx->stage) {
case MESA_SHADER_VERTEX:
if (ctx->options->key.vs.as_ls)
{
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));
}
}
ac_nir_get_max_workgroup_size(enum chip_class chip_class,
const struct nir_shader *nir)
{
- switch (nir->stage) {
+ switch (nir->info.stage) {
case MESA_SHADER_TESS_CTRL:
return chip_class >= CIK ? 128 : 64;
case MESA_SHADER_GEOMETRY:
return max_workgroup_size;
}
+/* Fixup the HW not emitting the TCS regs if there are no HS threads. */
+static void ac_nir_fixup_ls_hs_input_vgprs(struct nir_to_llvm_context *ctx)
+{
+ LLVMValueRef count = ac_build_bfe(&ctx->ac, ctx->merged_wave_info,
+ LLVMConstInt(ctx->ac.i32, 8, false),
+ LLVMConstInt(ctx->ac.i32, 8, false), false);
+ LLVMValueRef hs_empty = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, count,
+ LLVMConstInt(ctx->ac.i32, 0, false), "");
+ ctx->abi.instance_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->rel_auto_id, ctx->abi.instance_id, "");
+ ctx->vs_prim_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.vertex_id, ctx->vs_prim_id, "");
+ ctx->rel_auto_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->tcs_rel_ids, ctx->rel_auto_id, "");
+ ctx->abi.vertex_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->tcs_patch_id, ctx->abi.vertex_id, "");
+}
+
+static void prepare_gs_input_vgprs(struct nir_to_llvm_context *ctx)
+{
+ for(int i = 5; i >= 0; --i) {
+ ctx->gs_vtx_offset[i] = ac_build_bfe(&ctx->ac, ctx->gs_vtx_offset[i & ~1],
+ LLVMConstInt(ctx->ac.i32, (i & 1) * 16, false),
+ LLVMConstInt(ctx->ac.i32, 16, false), false);
+ }
+
+ ctx->gs_wave_id = ac_build_bfe(&ctx->ac, ctx->merged_wave_info,
+ LLVMConstInt(ctx->ac.i32, 16, false),
+ LLVMConstInt(ctx->ac.i32, 8, false), false);
+}
+
+void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
+ struct nir_shader *nir, struct nir_to_llvm_context *nctx)
+{
+ struct ac_nir_context ctx = {};
+ struct nir_function *func;
+
+ ctx.ac = *ac;
+ ctx.abi = abi;
+
+ ctx.nctx = nctx;
+ if (nctx)
+ nctx->nir = &ctx;
+
+ ctx.stage = nir->info.stage;
+
+ ctx.main_function = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx.ac.builder));
+
+ nir_foreach_variable(variable, &nir->outputs)
+ handle_shader_output_decl(&ctx, nir, variable);
+
+ ctx.defs = _mesa_hash_table_create(NULL, _mesa_hash_pointer,
+ _mesa_key_pointer_equal);
+ ctx.phis = _mesa_hash_table_create(NULL, _mesa_hash_pointer,
+ _mesa_key_pointer_equal);
+ ctx.vars = _mesa_hash_table_create(NULL, _mesa_hash_pointer,
+ _mesa_key_pointer_equal);
+
+ func = (struct nir_function *)exec_list_get_head(&nir->functions);
+
+ setup_locals(&ctx, func);
+
+ if (nir->info.stage == MESA_SHADER_COMPUTE)
+ setup_shared(&ctx, nir);
+
+ visit_cf_list(&ctx, &func->impl->body);
+ phi_post_pass(&ctx);
+
+ ctx.abi->emit_outputs(ctx.abi, RADEON_LLVM_MAX_OUTPUTS,
+ ctx.outputs);
+
+ free(ctx.locals);
+ ralloc_free(ctx.defs);
+ ralloc_free(ctx.phis);
+ ralloc_free(ctx.vars);
+
+ if (nctx)
+ nctx->nir = NULL;
+}
+
static
LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
- struct nir_shader *nir,
+ struct nir_shader *const *shaders,
+ int shader_count,
struct ac_shader_variant_info *shader_info,
const struct ac_nir_compiler_options *options)
{
struct nir_to_llvm_context ctx = {0};
- struct nir_function *func;
unsigned i;
ctx.options = options;
ctx.shader_info = shader_info;
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;
-
- ctx.has_ds_bpermute = ctx.options->chip_class >= VI;
-
- memset(shader_info, 0, sizeof(*shader_info));
-
- ac_nir_shader_info_pass(nir, options, &shader_info->info);
-
LLVMSetTarget(ctx.module, options->supports_spill ? "amdgcn-mesa-mesa3d" : "amdgcn--");
LLVMTargetDataRef data_layout = LLVMCreateTargetDataLayout(tm);
LLVMDisposeMessage(data_layout_str);
setup_types(&ctx);
-
ctx.builder = LLVMCreateBuilderInContext(ctx.context);
ctx.ac.builder = ctx.builder;
- ctx.stage = nir->stage;
- ctx.max_workgroup_size = ac_nir_get_max_workgroup_size(ctx.options->chip_class, nir);
+
+ memset(shader_info, 0, sizeof(*shader_info));
+
+ for(int i = 0; i < shader_count; ++i)
+ ac_nir_shader_info_pass(shaders[i], options, &shader_info->info);
for (i = 0; i < AC_UD_MAX_SETS; i++)
shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
for (i = 0; i < AC_UD_MAX_UD; i++)
shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
- create_function(&ctx);
+ ctx.max_workgroup_size = ac_nir_get_max_workgroup_size(ctx.options->chip_class, shaders[0]);
- if (nir->stage == MESA_SHADER_GEOMETRY) {
- ctx.gs_next_vertex = ac_build_alloca(&ctx, ctx.i32, "gs_next_vertex");
+ create_function(&ctx, shaders[shader_count - 1]->info.stage, shader_count >= 2,
+ shader_count >= 2 ? shaders[shader_count - 2]->info.stage : MESA_SHADER_VERTEX);
- ctx.gs_max_out_vertices = nir->info.gs.vertices_out;
- } else if (nir->stage == MESA_SHADER_TESS_EVAL) {
- ctx.tes_primitive_mode = nir->info.tess.primitive_mode;
- }
+ 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;
- ac_setup_rings(&ctx);
+ if (shader_count >= 2)
+ ac_init_exec_full_mask(&ctx.ac);
- nir_foreach_variable(variable, &nir->inputs)
- handle_shader_input_decl(&ctx, variable);
+ if (ctx.ac.chip_class == GFX9 &&
+ shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
+ ac_nir_fixup_ls_hs_input_vgprs(&ctx);
- if (nir->stage == MESA_SHADER_FRAGMENT)
- handle_fs_inputs_pre(&ctx, nir);
+ for(int i = 0; i < shader_count; ++i) {
+ ctx.stage = shaders[i]->info.stage;
+ ctx.output_mask = 0;
+ ctx.tess_outputs_written = 0;
+ ctx.num_output_clips = shaders[i]->info.clip_distance_array_size;
+ ctx.num_output_culls = shaders[i]->info.cull_distance_array_size;
- ctx.num_output_clips = nir->info.clip_distance_array_size;
- ctx.num_output_culls = nir->info.cull_distance_array_size;
+ if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
+ ctx.gs_next_vertex = ac_build_alloca(&ctx.ac, ctx.i32, "gs_next_vertex");
- nir_foreach_variable(variable, &nir->outputs)
- handle_shader_output_decl(&ctx, variable);
+ ctx.gs_max_out_vertices = shaders[i]->info.gs.vertices_out;
+ } else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
+ ctx.tes_primitive_mode = shaders[i]->info.tess.primitive_mode;
+ } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
+ if (shader_info->info.vs.needs_instance_id) {
+ ctx.shader_info->vs.vgpr_comp_cnt =
+ MAX2(3, ctx.shader_info->vs.vgpr_comp_cnt);
+ }
+ } else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
+ shader_info->fs.can_discard = shaders[i]->info.fs.uses_discard;
+ }
- ctx.defs = _mesa_hash_table_create(NULL, _mesa_hash_pointer,
- _mesa_key_pointer_equal);
- ctx.phis = _mesa_hash_table_create(NULL, _mesa_hash_pointer,
- _mesa_key_pointer_equal);
- ctx.vars = _mesa_hash_table_create(NULL, _mesa_hash_pointer,
- _mesa_key_pointer_equal);
+ if (i)
+ emit_barrier(&ctx);
- func = (struct nir_function *)exec_list_get_head(&nir->functions);
+ ac_setup_rings(&ctx);
- setup_locals(&ctx, func);
+ LLVMBasicBlockRef merge_block;
+ if (shader_count >= 2) {
+ LLVMValueRef fn = LLVMGetBasicBlockParent(LLVMGetInsertBlock(ctx.ac.builder));
+ LLVMBasicBlockRef then_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
+ merge_block = LLVMAppendBasicBlockInContext(ctx.ac.context, fn, "");
- if (nir->stage == MESA_SHADER_COMPUTE)
- setup_shared(&ctx, nir);
+ LLVMValueRef count = ac_build_bfe(&ctx.ac, ctx.merged_wave_info,
+ LLVMConstInt(ctx.ac.i32, 8 * i, false),
+ LLVMConstInt(ctx.ac.i32, 8, false), false);
+ LLVMValueRef thread_id = ac_get_thread_id(&ctx.ac);
+ LLVMValueRef cond = LLVMBuildICmp(ctx.ac.builder, LLVMIntULT,
+ thread_id, count, "");
+ LLVMBuildCondBr(ctx.ac.builder, cond, then_block, merge_block);
- visit_cf_list(&ctx, &func->impl->body);
- phi_post_pass(&ctx);
+ LLVMPositionBuilderAtEnd(ctx.ac.builder, then_block);
+ }
- handle_shader_outputs_post(&ctx);
- LLVMBuildRetVoid(ctx.builder);
+ if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT)
+ handle_fs_inputs(&ctx, shaders[i]);
+ else if(shaders[i]->info.stage == MESA_SHADER_VERTEX)
+ handle_vs_inputs(&ctx, shaders[i]);
+ else if(shader_count >= 2 && shaders[i]->info.stage == MESA_SHADER_GEOMETRY)
+ prepare_gs_input_vgprs(&ctx);
- ac_llvm_finalize_module(&ctx);
+ nir_foreach_variable(variable, &shaders[i]->outputs)
+ scan_shader_output_decl(&ctx, variable, shaders[i], shaders[i]->info.stage);
- ac_nir_eliminate_const_vs_outputs(&ctx);
- free(ctx.locals);
- ralloc_free(ctx.defs);
- ralloc_free(ctx.phis);
- ralloc_free(ctx.vars);
+ ac_nir_translate(&ctx.ac, &ctx.abi, shaders[i], &ctx);
- if (nir->stage == MESA_SHADER_GEOMETRY) {
- unsigned addclip = ctx.num_output_clips + ctx.num_output_culls > 4;
- shader_info->gs.gsvs_vertex_size = (util_bitcount64(ctx.output_mask) + addclip) * 16;
- shader_info->gs.max_gsvs_emit_size = shader_info->gs.gsvs_vertex_size *
- nir->info.gs.vertices_out;
- } else if (nir->stage == MESA_SHADER_TESS_CTRL) {
- shader_info->tcs.outputs_written = ctx.tess_outputs_written;
- shader_info->tcs.patch_outputs_written = ctx.tess_patch_outputs_written;
- } else if (nir->stage == MESA_SHADER_VERTEX && ctx.options->key.vs.as_ls) {
- shader_info->vs.outputs_written = ctx.tess_outputs_written;
+ if (shader_count >= 2) {
+ LLVMBuildBr(ctx.ac.builder, merge_block);
+ LLVMPositionBuilderAtEnd(ctx.ac.builder, merge_block);
+ }
+
+ if (shaders[i]->info.stage == MESA_SHADER_GEOMETRY) {
+ unsigned addclip = shaders[i]->info.clip_distance_array_size +
+ shaders[i]->info.cull_distance_array_size > 4;
+ shader_info->gs.gsvs_vertex_size = (util_bitcount64(ctx.output_mask) + addclip) * 16;
+ shader_info->gs.max_gsvs_emit_size = shader_info->gs.gsvs_vertex_size *
+ shaders[i]->info.gs.vertices_out;
+ } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
+ shader_info->tcs.outputs_written = ctx.tess_outputs_written;
+ shader_info->tcs.patch_outputs_written = ctx.tess_patch_outputs_written;
+ } else if (shaders[i]->info.stage == MESA_SHADER_VERTEX && ctx.options->key.vs.as_ls) {
+ shader_info->vs.outputs_written = ctx.tess_outputs_written;
+ }
}
+ LLVMBuildRetVoid(ctx.builder);
+
+ ac_llvm_finalize_module(&ctx);
+
+ if (shader_count == 1)
+ ac_nir_eliminate_const_vs_outputs(&ctx);
+
return ctx.module;
}
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->info.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, nir_count, 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]->info.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->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.num_output_clips = geom_shader->info.clip_distance_array_size;
ctx.num_output_culls = geom_shader->info.cull_distance_array_size;
- nir_foreach_variable(variable, &geom_shader->outputs)
- handle_shader_output_decl(&ctx, variable);
+ struct ac_nir_context nir_ctx = {};
+ nir_ctx.ac = ctx.ac;
+ nir_ctx.abi = &ctx.abi;
+
+ nir_ctx.nctx = &ctx;
+ ctx.nir = &nir_ctx;
+
+ nir_foreach_variable(variable, &geom_shader->outputs) {
+ scan_shader_output_decl(&ctx, variable, geom_shader, MESA_SHADER_VERTEX);
+ handle_shader_output_decl(&nir_ctx, geom_shader, variable);
+ }
ac_gs_copy_shader_emit(&ctx);
+ ctx.nir = NULL;
+
LLVMBuildRetVoid(ctx.builder);
ac_llvm_finalize_module(&ctx);
projects / mesa.git / blobdiff