#include <llvm-c/Transforms/Scalar.h>
#include "ac_shader_abi.h"
#include "ac_shader_info.h"
+#include "ac_shader_util.h"
#include "ac_exp_param.h"
enum radeon_llvm_calling_convention {
RADEON_LLVM_AMDGPU_HS = 93,
};
-#define CONST_ADDR_SPACE 2
-#define LOCAL_ADDR_SPACE 3
-
#define RADEON_LLVM_MAX_INPUTS (VARYING_SLOT_VAR31 + 1)
#define RADEON_LLVM_MAX_OUTPUTS (VARYING_SLOT_VAR31 + 1)
LLVMValueRef ring_offsets;
LLVMValueRef push_constants;
LLVMValueRef view_index;
- LLVMValueRef num_work_groups;
- LLVMValueRef workgroup_ids;
- LLVMValueRef local_invocation_ids;
LLVMValueRef tg_size;
LLVMValueRef vertex_buffers;
LLVMValueRef oc_lds;
LLVMValueRef merged_wave_info;
LLVMValueRef tess_factor_offset;
- LLVMValueRef tcs_patch_id;
- LLVMValueRef tcs_rel_ids;
LLVMValueRef tes_rel_patch_id;
- LLVMValueRef tes_patch_id;
LLVMValueRef tes_u;
LLVMValueRef tes_v;
LLVMValueRef gs2vs_offset;
LLVMValueRef gs_wave_id;
LLVMValueRef gs_vtx_offset[6];
- LLVMValueRef gs_prim_id, gs_invocation_id;
LLVMValueRef esgs_ring;
LLVMValueRef gsvs_ring;
LLVMValueRef hs_ring_tess_offchip;
LLVMValueRef hs_ring_tess_factor;
- LLVMValueRef prim_mask;
LLVMValueRef sample_pos_offset;
LLVMValueRef persp_sample, persp_center, persp_centroid;
LLVMValueRef linear_sample, linear_center, linear_centroid;
- LLVMTypeRef i1;
- LLVMTypeRef i8;
- LLVMTypeRef i16;
- LLVMTypeRef i32;
- LLVMTypeRef i64;
- LLVMTypeRef v2i32;
- LLVMTypeRef v3i32;
- LLVMTypeRef v4i32;
- LLVMTypeRef v8i32;
- LLVMTypeRef f64;
- LLVMTypeRef f32;
- LLVMTypeRef f16;
- LLVMTypeRef v2f32;
- LLVMTypeRef v4f32;
- LLVMTypeRef voidt;
-
- LLVMValueRef i1true;
- LLVMValueRef i1false;
- LLVMValueRef i32zero;
- LLVMValueRef i32one;
- LLVMValueRef f32zero;
- LLVMValueRef f32one;
- LLVMValueRef v4f32empty;
-
- unsigned uniform_md_kind;
- LLVMValueRef empty_md;
gl_shader_stage stage;
- LLVMValueRef lds;
LLVMValueRef inputs[RADEON_LLVM_MAX_INPUTS * 4];
uint64_t input_mask;
unsigned tes_primitive_mode;
uint64_t tess_outputs_written;
uint64_t tess_patch_outputs_written;
+
+ uint32_t tcs_patch_outputs_read;
+ uint64_t tcs_outputs_read;
};
static inline struct nir_to_llvm_context *
static LLVMValueRef get_sampler_desc(struct ac_nir_context *ctx,
const nir_deref_var *deref,
enum ac_descriptor_type desc_type,
+ const nir_tex_instr *instr,
bool image, bool write);
static unsigned radeon_llvm_reg_index_soa(unsigned index, unsigned chan)
LLVMValueRef *assign[MAX_ARGS];
unsigned array_params_mask;
uint8_t count;
- uint8_t user_sgpr_count;
uint8_t sgpr_count;
- uint8_t num_user_sgprs_used;
uint8_t num_sgprs_used;
uint8_t num_vgprs_used;
};
-static inline void
-add_argument(struct arg_info *info,
- LLVMTypeRef type, LLVMValueRef *param_ptr)
+enum ac_arg_regfile {
+ ARG_SGPR,
+ ARG_VGPR,
+};
+
+static void
+add_arg(struct arg_info *info, enum ac_arg_regfile regfile, LLVMTypeRef type,
+ LLVMValueRef *param_ptr)
{
assert(info->count < MAX_ARGS);
+
info->assign[info->count] = param_ptr;
info->types[info->count] = type;
info->count++;
-}
-
-static inline void
-add_sgpr_argument(struct arg_info *info,
- LLVMTypeRef type, LLVMValueRef *param_ptr)
-{
- add_argument(info, type, param_ptr);
- info->num_sgprs_used += ac_get_type_size(type) / 4;
- info->sgpr_count++;
-}
-static inline void
-add_user_sgpr_argument(struct arg_info *info,
- LLVMTypeRef type,
- LLVMValueRef *param_ptr)
-{
- add_sgpr_argument(info, type, param_ptr);
- info->num_user_sgprs_used += ac_get_type_size(type) / 4;
- info->user_sgpr_count++;
-}
-
-static inline void
-add_vgpr_argument(struct arg_info *info,
- LLVMTypeRef type,
- LLVMValueRef *param_ptr)
-{
- add_argument(info, type, param_ptr);
- info->num_vgprs_used += ac_get_type_size(type) / 4;
+ if (regfile == ARG_SGPR) {
+ info->num_sgprs_used += ac_get_type_size(type) / 4;
+ info->sgpr_count++;
+ } else {
+ assert(regfile == ARG_VGPR);
+ info->num_vgprs_used += ac_get_type_size(type) / 4;
+ }
}
static inline void
-add_user_sgpr_array_argument(struct arg_info *info,
- LLVMTypeRef type,
- LLVMValueRef *param_ptr)
+add_array_arg(struct arg_info *info, LLVMTypeRef type, LLVMValueRef *param_ptr)
{
info->array_params_mask |= (1 << info->count);
- add_user_sgpr_argument(info, type, param_ptr);
+ add_arg(info, ARG_SGPR, type, param_ptr);
}
static void assign_arguments(LLVMValueRef main_function,
LLVMSetFunctionCallConv(main_function, RADEON_LLVM_AMDGPU_CS);
for (unsigned i = 0; i < args->sgpr_count; ++i) {
+ ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
+
if (args->array_params_mask & (1 << i)) {
LLVMValueRef P = LLVMGetParam(main_function, i);
- ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_BYVAL);
+ ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_NOALIAS);
ac_add_attr_dereferenceable(P, UINT64_MAX);
}
- else {
- ac_add_function_attr(ctx, main_function, i + 1, AC_FUNC_ATTR_INREG);
- }
}
if (max_workgroup_size) {
LLVMAddTargetDependentFunctionAttr(main_function,
"unsafe-fp-math",
"true");
+ LLVMAddTargetDependentFunctionAttr(main_function,
+ "no-signed-zeros-fp-math",
+ "true");
}
return main_function;
}
-static LLVMTypeRef const_array(LLVMTypeRef elem_type, int num_elements)
-{
- return LLVMPointerType(LLVMArrayType(elem_type, num_elements),
- CONST_ADDR_SPACE);
-}
-
static int get_elem_bits(struct ac_llvm_context *ctx, LLVMTypeRef type)
{
if (LLVMGetTypeKind(type) == LLVMVectorTypeKind)
{
switch (ctx->stage) {
case MESA_SHADER_TESS_CTRL:
- return unpack_param(&ctx->ac, ctx->tcs_rel_ids, 0, 8);
+ return unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 0, 8);
case MESA_SHADER_TESS_EVAL:
return ctx->tes_rel_patch_id;
break;
{
return LLVMBuildMul(ctx->builder,
unpack_param(&ctx->ac, ctx->tcs_out_offsets, 0, 16),
- LLVMConstInt(ctx->i32, 4, false), "");
+ LLVMConstInt(ctx->ac.i32, 4, false), "");
}
static LLVMValueRef
{
return LLVMBuildMul(ctx->builder,
unpack_param(&ctx->ac, ctx->tcs_out_offsets, 16, 16),
- LLVMConstInt(ctx->i32, 4, false), "");
+ LLVMConstInt(ctx->ac.i32, 4, false), "");
}
static LLVMValueRef
"");
}
-static void set_userdata_location(struct ac_userdata_info *ud_info, uint8_t *sgpr_idx, uint8_t num_sgprs)
+static void
+set_loc(struct ac_userdata_info *ud_info, uint8_t *sgpr_idx, uint8_t num_sgprs,
+ uint32_t indirect_offset)
{
ud_info->sgpr_idx = *sgpr_idx;
ud_info->num_sgprs = num_sgprs;
- ud_info->indirect = false;
- ud_info->indirect_offset = 0;
+ ud_info->indirect = indirect_offset > 0;
+ ud_info->indirect_offset = indirect_offset;
*sgpr_idx += num_sgprs;
}
-static void set_userdata_location_shader(struct nir_to_llvm_context *ctx,
- int idx, uint8_t *sgpr_idx, uint8_t num_sgprs)
+static void
+set_loc_shader(struct nir_to_llvm_context *ctx, int idx, uint8_t *sgpr_idx,
+ uint8_t num_sgprs)
{
- set_userdata_location(&ctx->shader_info->user_sgprs_locs.shader_data[idx], sgpr_idx, num_sgprs);
-}
-
+ struct ac_userdata_info *ud_info =
+ &ctx->shader_info->user_sgprs_locs.shader_data[idx];
+ assert(ud_info);
-static void set_userdata_location_indirect(struct ac_userdata_info *ud_info, uint8_t sgpr_idx, uint8_t num_sgprs,
- uint32_t indirect_offset)
-{
- ud_info->sgpr_idx = sgpr_idx;
- ud_info->num_sgprs = num_sgprs;
- ud_info->indirect = true;
- ud_info->indirect_offset = indirect_offset;
+ set_loc(ud_info, sgpr_idx, num_sgprs, 0);
}
-static void declare_tess_lds(struct nir_to_llvm_context *ctx)
+static void
+set_loc_desc(struct nir_to_llvm_context *ctx, int idx, uint8_t *sgpr_idx,
+ uint32_t indirect_offset)
{
- unsigned lds_size = ctx->options->chip_class >= CIK ? 65536 : 32768;
- ctx->lds = LLVMBuildIntToPtr(ctx->builder, ctx->i32zero,
- LLVMPointerType(LLVMArrayType(ctx->i32, lds_size / 4), LOCAL_ADDR_SPACE),
- "tess_lds");
+ struct ac_userdata_info *ud_info =
+ &ctx->shader_info->user_sgprs_locs.descriptor_sets[idx];
+ assert(ud_info);
+
+ set_loc(ud_info, sgpr_idx, 2, indirect_offset);
}
struct user_sgpr_info {
bool indirect_all_descriptor_sets;
};
+static bool needs_view_index_sgpr(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage)
+{
+ switch (stage) {
+ case MESA_SHADER_VERTEX:
+ 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))
+ return true;
+ break;
+ case MESA_SHADER_TESS_EVAL:
+ if (ctx->shader_info->info.needs_multiview_view_index || (!ctx->options->key.tes.as_es && ctx->options->key.has_multiview_view_index))
+ return true;
+ break;
+ case MESA_SHADER_GEOMETRY:
+ case MESA_SHADER_TESS_CTRL:
+ if (ctx->shader_info->info.needs_multiview_view_index)
+ return true;
+ break;
+ default:
+ break;
+ }
+ return false;
+}
+
static void allocate_user_sgprs(struct nir_to_llvm_context *ctx,
+ gl_shader_stage stage,
+ bool needs_view_index,
struct user_sgpr_info *user_sgpr_info)
{
memset(user_sgpr_info, 0, sizeof(struct user_sgpr_info));
/* until we sort out scratch/global buffers always assign ring offsets for gs/vs/es */
- if (ctx->stage == MESA_SHADER_GEOMETRY ||
- ctx->stage == MESA_SHADER_VERTEX ||
- ctx->stage == MESA_SHADER_TESS_CTRL ||
- ctx->stage == MESA_SHADER_TESS_EVAL ||
+ if (stage == MESA_SHADER_GEOMETRY ||
+ stage == MESA_SHADER_VERTEX ||
+ stage == MESA_SHADER_TESS_CTRL ||
+ stage == MESA_SHADER_TESS_EVAL ||
ctx->is_gs_copy_shader)
user_sgpr_info->need_ring_offsets = true;
- if (ctx->stage == MESA_SHADER_FRAGMENT &&
+ if (stage == MESA_SHADER_FRAGMENT &&
ctx->shader_info->info.ps.needs_sample_positions)
user_sgpr_info->need_ring_offsets = true;
user_sgpr_info->sgpr_count += 2;
}
- switch (ctx->stage) {
+ /* FIXME: fix the number of user sgprs for merged shaders on GFX9 */
+ switch (stage) {
case MESA_SHADER_COMPUTE:
- user_sgpr_info->sgpr_count += ctx->shader_info->info.cs.grid_components_used;
+ if (ctx->shader_info->info.cs.uses_grid_size)
+ user_sgpr_info->sgpr_count += 3;
break;
case MESA_SHADER_FRAGMENT:
user_sgpr_info->sgpr_count += ctx->shader_info->info.ps.needs_sample_positions;
break;
}
- if (ctx->shader_info->info.needs_push_constants)
+ if (needs_view_index)
+ user_sgpr_info->sgpr_count++;
+
+ if (ctx->shader_info->info.loads_push_constants)
user_sgpr_info->sgpr_count += 2;
- uint32_t remaining_sgprs = 16 - user_sgpr_info->sgpr_count;
+ uint32_t available_sgprs = ctx->options->chip_class >= GFX9 ? 32 : 16;
+ uint32_t remaining_sgprs = available_sgprs - user_sgpr_info->sgpr_count;
+
if (remaining_sgprs / 2 < util_bitcount(ctx->shader_info->info.desc_set_used_mask)) {
user_sgpr_info->sgpr_count += 2;
user_sgpr_info->indirect_all_descriptor_sets = true;
}
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;
+declare_global_input_sgprs(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)
+{
+ LLVMTypeRef type = ac_array_in_const_addr_space(ctx->ac.i8);
+ 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) {
- add_user_sgpr_array_argument(args, const_array(ctx->i8, 1024 * 1024), &ctx->descriptor_sets[i]);
+ add_array_arg(args, type,
+ &ctx->descriptor_sets[i]);
}
}
- } else
- add_user_sgpr_array_argument(args, const_array(const_array(ctx->i8, 1024 * 1024), 32), desc_sets);
+ } else {
+ add_array_arg(args, ac_array_in_const_addr_space(type), desc_sets);
+ }
- if (ctx->shader_info->info.needs_push_constants) {
+ if (ctx->shader_info->info.loads_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_array_arg(args, type, &ctx->push_constants);
+ }
+}
+
+static void
+declare_vs_specific_input_sgprs(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_arg(args, ARG_SGPR, ac_array_in_const_addr_space(ctx->ac.v4i32),
+ &ctx->vertex_buffers);
+ }
+ add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.base_vertex);
+ add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.start_instance);
+ if (ctx->shader_info->info.vs.needs_draw_id) {
+ add_arg(args, ARG_SGPR, ctx->ac.i32, &ctx->abi.draw_id);
+ }
+ }
+}
+
+static void
+declare_vs_input_vgprs(struct nir_to_llvm_context *ctx, struct arg_info *args)
+{
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.vertex_id);
+ if (!ctx->is_gs_copy_shader) {
+ if (ctx->options->key.vs.as_ls) {
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->rel_auto_id);
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
+ } else {
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.instance_id);
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->vs_prim_id);
+ }
+ add_arg(args, ARG_VGPR, ctx->ac.i32, NULL); /* unused */
}
}
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;
+declare_tes_input_vgprs(struct nir_to_llvm_context *ctx, struct arg_info *args)
+{
+ add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_u);
+ add_arg(args, ARG_VGPR, ctx->ac.f32, &ctx->tes_v);
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->tes_rel_patch_id);
+ add_arg(args, ARG_VGPR, ctx->ac.i32, &ctx->abi.tes_patch_id);
+}
+
+static void
+set_global_input_locs(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);
+ set_loc_desc(ctx, i, user_sgpr_idx, 0);
} 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);
+ set_loc_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));
+ set_loc_desc(ctx, i, user_sgpr_idx, i * 8);
+ ctx->descriptor_sets[i] =
+ ac_build_load_to_sgpr(&ctx->ac,
+ desc_sets,
+ LLVMConstInt(ctx->ac.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);
+ if (ctx->shader_info->info.loads_push_constants) {
+ set_loc_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))) {
+set_vs_specific_input_locs(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);
+ set_loc_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);
+ set_loc_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,
struct user_sgpr_info user_sgpr_info;
struct arg_info args = {};
LLVMValueRef desc_sets;
-
- allocate_user_sgprs(ctx, &user_sgpr_info);
+ bool needs_view_index = needs_view_index_sgpr(ctx, stage);
+ allocate_user_sgprs(ctx, stage, needs_view_index, &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 */
+ add_arg(&args, ARG_SGPR, ac_array_in_const_addr_space(ctx->ac.v4i32),
+ &ctx->ring_offsets);
}
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_sgpr_argument(&args, ctx->i32, &ctx->tg_size);
- add_vgpr_argument(&args, LLVMVectorType(ctx->i32, 3), &ctx->local_invocation_ids);
+ declare_global_input_sgprs(ctx, stage, has_previous_stage,
+ previous_stage, &user_sgpr_info,
+ &args, &desc_sets);
+
+ if (ctx->shader_info->info.cs.uses_grid_size) {
+ add_arg(&args, ARG_SGPR, ctx->ac.v3i32,
+ &ctx->abi.num_work_groups);
+ }
+
+ for (int i = 0; i < 3; i++) {
+ ctx->abi.workgroup_ids[i] = NULL;
+ if (ctx->shader_info->info.cs.uses_block_id[i]) {
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->abi.workgroup_ids[i]);
+ }
+ }
+
+ if (ctx->shader_info->info.cs.uses_local_invocation_idx)
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->tg_size);
+ add_arg(&args, ARG_VGPR, ctx->ac.v3i32,
+ &ctx->abi.local_invocation_ids);
break;
case MESA_SHADER_VERTEX:
- 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);
+ declare_global_input_sgprs(ctx, stage, has_previous_stage,
+ previous_stage, &user_sgpr_info,
+ &args, &desc_sets);
+ declare_vs_specific_input_sgprs(ctx, stage, has_previous_stage,
+ previous_stage, &args);
+
+ if (needs_view_index)
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->view_index);
if (ctx->options->key.vs.as_es)
- add_sgpr_argument(&args, ctx->i32, &ctx->es2gs_offset); // es2gs offset
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->es2gs_offset);
else if (ctx->options->key.vs.as_ls)
- add_user_sgpr_argument(&args, ctx->i32, &ctx->ls_out_layout); // ls out layout
- add_vgpr_argument(&args, ctx->i32, &ctx->abi.vertex_id); // vertex id
- if (!ctx->is_gs_copy_shader) {
- 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
- }
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->ls_out_layout);
+
+ declare_vs_input_vgprs(ctx, &args);
break;
case MESA_SHADER_TESS_CTRL:
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
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->merged_wave_info);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tess_factor_offset);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
+
+ declare_global_input_sgprs(ctx, stage,
+ has_previous_stage,
+ previous_stage,
+ &user_sgpr_info, &args,
+ &desc_sets);
+ declare_vs_specific_input_sgprs(ctx, stage,
+ has_previous_stage,
+ previous_stage, &args);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->ls_out_layout);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_offchip_layout);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_out_offsets);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_out_layout);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_in_layout);
+ if (needs_view_index)
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->view_index);
+
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.tcs_patch_id);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.tcs_rel_ids);
+
+ declare_vs_input_vgprs(ctx, &args);
} 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;
+ declare_global_input_sgprs(ctx, stage,
+ has_previous_stage,
+ previous_stage,
+ &user_sgpr_info, &args,
+ &desc_sets);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_offchip_layout);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_out_offsets);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_out_layout);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_in_layout);
+ if (needs_view_index)
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->view_index);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tess_factor_offset);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.tcs_patch_id);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.tcs_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);
+ declare_global_input_sgprs(ctx, stage, has_previous_stage,
+ previous_stage, &user_sgpr_info,
+ &args, &desc_sets);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->tcs_offchip_layout);
+ if (needs_view_index)
+ add_arg(&args, ARG_SGPR, ctx->ac.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_sgpr_argument(&args, ctx->i32, &ctx->es2gs_offset); // es2gs offset
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->es2gs_offset);
} else {
- add_sgpr_argument(&args, ctx->i32, NULL); //
- add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // OC LDS
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
}
- 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
+ declare_tes_input_vgprs(ctx, &args);
break;
case MESA_SHADER_GEOMETRY:
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
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->gs2vs_offset);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->merged_wave_info);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->oc_lds);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // scratch offset
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, NULL); // unknown
+
+ declare_global_input_sgprs(ctx, stage,
+ has_previous_stage,
+ previous_stage,
+ &user_sgpr_info, &args,
+ &desc_sets);
+
+ if (previous_stage == MESA_SHADER_TESS_EVAL) {
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->tcs_offchip_layout);
+ } else {
+ declare_vs_specific_input_sgprs(ctx, stage,
+ has_previous_stage,
+ previous_stage,
+ &args);
+ }
- 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]);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->gsvs_ring_stride);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->gsvs_num_entries);
+ if (needs_view_index)
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->view_index);
+
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[0]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[2]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.gs_prim_id);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.gs_invocation_id);
+ add_arg(&args, ARG_VGPR, ctx->ac.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
+ declare_vs_input_vgprs(ctx, &args);
} 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
+ declare_tes_input_vgprs(ctx, &args);
}
} 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);
+ declare_global_input_sgprs(ctx, stage,
+ has_previous_stage,
+ previous_stage,
+ &user_sgpr_info, &args,
+ &desc_sets);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->gsvs_ring_stride);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->gsvs_num_entries);
+ if (needs_view_index)
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->view_index);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs2vs_offset);
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->gs_wave_id);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[0]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[1]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.gs_prim_id);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[2]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[3]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[4]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->gs_vtx_offset[5]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32,
+ &ctx->abi.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);
+ declare_global_input_sgprs(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->persp_sample); /* persp sample */
- add_vgpr_argument(&args, ctx->v2i32, &ctx->persp_center); /* persp center */
- add_vgpr_argument(&args, ctx->v2i32, &ctx->persp_centroid); /* persp centroid */
- add_vgpr_argument(&args, ctx->v3i32, NULL); /* persp pull model */
- add_vgpr_argument(&args, ctx->v2i32, &ctx->linear_sample); /* linear sample */
- 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->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 */
+ add_arg(&args, ARG_SGPR, ctx->ac.i32,
+ &ctx->sample_pos_offset);
+
+ add_arg(&args, ARG_SGPR, ctx->ac.i32, &ctx->abi.prim_mask);
+ add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->persp_sample);
+ add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->persp_center);
+ add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->persp_centroid);
+ add_arg(&args, ARG_VGPR, ctx->ac.v3i32, NULL); /* persp pull model */
+ add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->linear_sample);
+ add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->linear_center);
+ add_arg(&args, ARG_VGPR, ctx->ac.v2i32, &ctx->linear_centroid);
+ add_arg(&args, ARG_VGPR, ctx->ac.f32, NULL); /* line stipple tex */
+ add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[0]);
+ add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[1]);
+ add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[2]);
+ add_arg(&args, ARG_VGPR, ctx->ac.f32, &ctx->abi.frag_pos[3]);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.front_face);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.ancillary);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32, &ctx->abi.sample_coverage);
+ add_arg(&args, ARG_VGPR, ctx->ac.i32, NULL); /* fixed pt */
break;
default:
unreachable("Shader stage not implemented");
user_sgpr_idx = 0;
if (ctx->options->supports_spill || user_sgpr_info.need_ring_offsets) {
- set_userdata_location_shader(ctx, AC_UD_SCRATCH_RING_OFFSETS, &user_sgpr_idx, 2);
+ set_loc_shader(ctx, AC_UD_SCRATCH_RING_OFFSETS,
+ &user_sgpr_idx, 2);
if (ctx->options->supports_spill) {
ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
- LLVMPointerType(ctx->i8, CONST_ADDR_SPACE),
+ LLVMPointerType(ctx->ac.i8, AC_CONST_ADDR_SPACE),
NULL, 0, AC_FUNC_ATTR_READNONE);
ctx->ring_offsets = LLVMBuildBitCast(ctx->builder, ctx->ring_offsets,
- const_array(ctx->v4i32, 16), "");
+ ac_array_in_const_addr_space(ctx->ac.v4i32), "");
}
}
if (has_previous_stage)
user_sgpr_idx = 0;
- radv_define_common_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_info, desc_sets, &user_sgpr_idx);
+ set_global_input_locs(ctx, stage, has_previous_stage, previous_stage,
+ &user_sgpr_info, desc_sets, &user_sgpr_idx);
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);
+ if (ctx->shader_info->info.cs.uses_grid_size) {
+ set_loc_shader(ctx, AC_UD_CS_GRID_SIZE,
+ &user_sgpr_idx, 3);
}
break;
case MESA_SHADER_VERTEX:
- radv_define_vs_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_idx);
+ set_vs_specific_input_locs(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);
+ set_loc_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);
+ set_loc_shader(ctx, AC_UD_VS_LS_TCS_IN_LAYOUT,
+ &user_sgpr_idx, 1);
}
if (ctx->options->key.vs.as_ls)
- declare_tess_lds(ctx);
+ ac_declare_lds_as_pointer(&ctx->ac);
break;
case MESA_SHADER_TESS_CTRL:
- radv_define_vs_user_sgprs_phase2(ctx, stage, has_previous_stage, previous_stage, &user_sgpr_idx);
+ set_vs_specific_input_locs(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);
+ set_loc_shader(ctx, AC_UD_VS_LS_TCS_IN_LAYOUT,
+ &user_sgpr_idx, 1);
+ set_loc_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);
+ set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
+ ac_declare_lds_as_pointer(&ctx->ac);
break;
case MESA_SHADER_TESS_EVAL:
- set_userdata_location_shader(ctx, AC_UD_TES_OFFCHIP_LAYOUT, &user_sgpr_idx, 1);
+ set_loc_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);
+ set_loc_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);
+ set_vs_specific_input_locs(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_loc_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);
+ set_loc_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);
+ set_loc_shader(ctx, AC_UD_VIEW_INDEX, &user_sgpr_idx, 1);
if (has_previous_stage)
- declare_tess_lds(ctx);
+ ac_declare_lds_as_pointer(&ctx->ac);
break;
case MESA_SHADER_FRAGMENT:
if (ctx->shader_info->info.ps.needs_sample_positions) {
- set_userdata_location_shader(ctx, AC_UD_PS_SAMPLE_POS_OFFSET, &user_sgpr_idx, 1);
+ set_loc_shader(ctx, AC_UD_PS_SAMPLE_POS_OFFSET,
+ &user_sgpr_idx, 1);
}
break;
default:
ctx->shader_info->num_user_sgprs = user_sgpr_idx;
}
-static void setup_types(struct nir_to_llvm_context *ctx)
-{
- LLVMValueRef args[4];
-
- ctx->voidt = LLVMVoidTypeInContext(ctx->context);
- ctx->i1 = LLVMIntTypeInContext(ctx->context, 1);
- ctx->i8 = LLVMIntTypeInContext(ctx->context, 8);
- ctx->i16 = LLVMIntTypeInContext(ctx->context, 16);
- ctx->i32 = LLVMIntTypeInContext(ctx->context, 32);
- ctx->i64 = LLVMIntTypeInContext(ctx->context, 64);
- ctx->v2i32 = LLVMVectorType(ctx->i32, 2);
- ctx->v3i32 = LLVMVectorType(ctx->i32, 3);
- ctx->v4i32 = LLVMVectorType(ctx->i32, 4);
- ctx->v8i32 = LLVMVectorType(ctx->i32, 8);
- ctx->f32 = LLVMFloatTypeInContext(ctx->context);
- ctx->f16 = LLVMHalfTypeInContext(ctx->context);
- ctx->f64 = LLVMDoubleTypeInContext(ctx->context);
- ctx->v2f32 = LLVMVectorType(ctx->f32, 2);
- ctx->v4f32 = LLVMVectorType(ctx->f32, 4);
-
- ctx->i1false = LLVMConstInt(ctx->i1, 0, false);
- ctx->i1true = LLVMConstInt(ctx->i1, 1, false);
- ctx->i32zero = LLVMConstInt(ctx->i32, 0, false);
- ctx->i32one = LLVMConstInt(ctx->i32, 1, false);
- ctx->f32zero = LLVMConstReal(ctx->f32, 0.0);
- ctx->f32one = LLVMConstReal(ctx->f32, 1.0);
-
- args[0] = ctx->f32zero;
- args[1] = ctx->f32zero;
- args[2] = ctx->f32zero;
- args[3] = ctx->f32one;
- ctx->v4f32empty = LLVMConstVector(args, 4);
-
- ctx->uniform_md_kind =
- LLVMGetMDKindIDInContext(ctx->context, "amdgpu.uniform", 14);
- ctx->empty_md = LLVMMDNodeInContext(ctx->context, NULL, 0);
-
- args[0] = LLVMConstReal(ctx->f32, 2.5);
-}
-
-static int get_llvm_num_components(LLVMValueRef value)
-{
- LLVMTypeRef type = LLVMTypeOf(value);
- unsigned num_components = LLVMGetTypeKind(type) == LLVMVectorTypeKind
- ? LLVMGetVectorSize(type)
- : 1;
- return num_components;
-}
-
-static LLVMValueRef llvm_extract_elem(struct ac_llvm_context *ac,
- LLVMValueRef value,
- int index)
-{
- int count = get_llvm_num_components(value);
-
- assert(index < count);
- if (count == 1)
- return value;
-
- return LLVMBuildExtractElement(ac->builder, value,
- LLVMConstInt(ac->i32, index, false), "");
-}
-
static LLVMValueRef trim_vector(struct ac_llvm_context *ctx,
LLVMValueRef value, unsigned count)
{
- unsigned num_components = get_llvm_num_components(value);
+ unsigned num_components = ac_get_llvm_num_components(value);
if (count == num_components)
return value;
LLVMValueRef result = LLVMBuildICmp(ctx->builder, pred, src0, src1, "");
return LLVMBuildSelect(ctx->builder, result,
LLVMConstInt(ctx->i32, 0xFFFFFFFF, false),
- LLVMConstInt(ctx->i32, 0, false), "");
+ ctx->i32_0, "");
}
static LLVMValueRef emit_float_cmp(struct ac_llvm_context *ctx,
result = LLVMBuildFCmp(ctx->builder, pred, src0, src1, "");
return LLVMBuildSelect(ctx->builder, result,
LLVMConstInt(ctx->i32, 0xFFFFFFFF, false),
- LLVMConstInt(ctx->i32, 0, false), "");
+ ctx->i32_0, "");
}
static LLVMValueRef emit_intrin_1f_param(struct ac_llvm_context *ctx,
return LLVMBuildSelect(ctx->builder, v, src1, src2, "");
}
-static LLVMValueRef emit_find_lsb(struct ac_llvm_context *ctx,
- LLVMValueRef src0)
-{
- LLVMValueRef params[2] = {
- src0,
-
- /* The value of 1 means that ffs(x=0) = undef, so LLVM won't
- * add special code to check for x=0. The reason is that
- * the LLVM behavior for x=0 is different from what we
- * need here.
- *
- * The hardware already implements the correct behavior.
- */
- LLVMConstInt(ctx->i1, 1, false),
- };
-
- 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)
-{
- return ac_build_imsb(ctx, src0, ctx->i32);
-}
-
-static LLVMValueRef emit_ufind_msb(struct ac_llvm_context *ctx,
- LLVMValueRef src0)
-{
- return ac_build_umsb(ctx, src0, ctx->i32);
-}
-
static LLVMValueRef emit_minmax_int(struct ac_llvm_context *ctx,
LLVMIntPredicate pred,
LLVMValueRef src0, LLVMValueRef src1)
}
static LLVMValueRef emit_fsign(struct ac_llvm_context *ctx,
- LLVMValueRef src0)
+ LLVMValueRef src0,
+ unsigned bitsize)
{
- LLVMValueRef cmp, val;
+ LLVMValueRef cmp, val, zero, one;
+ LLVMTypeRef type;
+
+ if (bitsize == 32) {
+ type = ctx->f32;
+ zero = ctx->f32_0;
+ one = ctx->f32_1;
+ } else {
+ type = ctx->f64;
+ zero = ctx->f64_0;
+ one = ctx->f64_1;
+ }
- cmp = LLVMBuildFCmp(ctx->builder, LLVMRealOGT, src0, ctx->f32_0, "");
- val = LLVMBuildSelect(ctx->builder, cmp, ctx->f32_1, src0, "");
- cmp = LLVMBuildFCmp(ctx->builder, LLVMRealOGE, val, ctx->f32_0, "");
- val = LLVMBuildSelect(ctx->builder, cmp, val, LLVMConstReal(ctx->f32, -1.0), "");
+ cmp = LLVMBuildFCmp(ctx->builder, LLVMRealOGT, src0, zero, "");
+ val = LLVMBuildSelect(ctx->builder, cmp, one, src0, "");
+ cmp = LLVMBuildFCmp(ctx->builder, LLVMRealOGE, val, zero, "");
+ val = LLVMBuildSelect(ctx->builder, cmp, val, LLVMConstReal(type, -1.0), "");
return val;
}
static LLVMValueRef emit_isign(struct ac_llvm_context *ctx,
- LLVMValueRef src0)
+ LLVMValueRef src0, unsigned bitsize)
{
- LLVMValueRef cmp, val;
+ LLVMValueRef cmp, val, zero, one;
+ LLVMTypeRef type;
+
+ if (bitsize == 32) {
+ type = ctx->i32;
+ zero = ctx->i32_0;
+ one = ctx->i32_1;
+ } else {
+ type = ctx->i64;
+ zero = ctx->i64_0;
+ one = ctx->i64_1;
+ }
- cmp = LLVMBuildICmp(ctx->builder, LLVMIntSGT, src0, ctx->i32_0, "");
- val = LLVMBuildSelect(ctx->builder, cmp, ctx->i32_1, src0, "");
- cmp = LLVMBuildICmp(ctx->builder, LLVMIntSGE, val, ctx->i32_0, "");
- val = LLVMBuildSelect(ctx->builder, cmp, val, LLVMConstInt(ctx->i32, -1, true), "");
+ cmp = LLVMBuildICmp(ctx->builder, LLVMIntSGT, src0, zero, "");
+ val = LLVMBuildSelect(ctx->builder, cmp, one, src0, "");
+ cmp = LLVMBuildICmp(ctx->builder, LLVMIntSGE, val, zero, "");
+ val = LLVMBuildSelect(ctx->builder, cmp, val, LLVMConstInt(type, -1, true), "");
return val;
}
static LLVMValueRef emit_ffract(struct ac_llvm_context *ctx,
- LLVMValueRef src0)
+ LLVMValueRef src0, unsigned bitsize)
{
- const char *intr = "llvm.floor.f32";
+ LLVMTypeRef type;
+ char *intr;
+
+ if (bitsize == 32) {
+ intr = "llvm.floor.f32";
+ type = ctx->f32;
+ } else {
+ intr = "llvm.floor.f64";
+ type = ctx->f64;
+ }
+
LLVMValueRef fsrc0 = ac_to_float(ctx, src0);
LLVMValueRef params[] = {
fsrc0,
};
- LLVMValueRef floor = ac_build_intrinsic(ctx, intr,
- ctx->f32, params, 1,
+ LLVMValueRef floor = ac_build_intrinsic(ctx, intr, type, params, 1,
AC_FUNC_ATTR_READNONE);
return LLVMBuildFSub(ctx->builder, fsrc0, floor, "");
}
}
static LLVMValueRef emit_b2i(struct ac_llvm_context *ctx,
- LLVMValueRef src0)
+ LLVMValueRef src0,
+ unsigned bitsize)
{
- return LLVMBuildAnd(ctx->builder, src0, ctx->i32_1, "");
+ LLVMValueRef result = LLVMBuildAnd(ctx->builder, src0, ctx->i32_1, "");
+
+ if (bitsize == 32)
+ return result;
+
+ return LLVMBuildZExt(ctx->builder, result, ctx->i64, "");
}
static LLVMValueRef emit_i2b(struct ac_llvm_context *ctx,
LLVMValueRef cond = NULL;
src0 = ac_to_float(&ctx->ac, src0);
- result = LLVMBuildFPTrunc(ctx->builder, src0, ctx->f16, "");
+ result = LLVMBuildFPTrunc(ctx->builder, src0, ctx->ac.f16, "");
if (ctx->options->chip_class >= VI) {
LLVMValueRef args[2];
/* Check if the result is a denormal - and flush to 0 if so. */
args[0] = result;
- args[1] = LLVMConstInt(ctx->i32, N_SUBNORMAL | P_SUBNORMAL, false);
- cond = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.class.f16", ctx->i1, args, 2, AC_FUNC_ATTR_READNONE);
+ args[1] = LLVMConstInt(ctx->ac.i32, N_SUBNORMAL | P_SUBNORMAL, false);
+ cond = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.class.f16", ctx->ac.i1, args, 2, AC_FUNC_ATTR_READNONE);
}
/* need to convert back up to f32 */
- result = LLVMBuildFPExt(ctx->builder, result, ctx->f32, "");
+ result = LLVMBuildFPExt(ctx->builder, result, ctx->ac.f32, "");
if (ctx->options->chip_class >= VI)
- result = LLVMBuildSelect(ctx->builder, cond, ctx->f32zero, result, "");
+ result = LLVMBuildSelect(ctx->builder, cond, ctx->ac.f32_0, result, "");
else {
/* for SI/CIK */
/* 0x38800000 is smallest half float value (2^-14) in 32-bit float,
*/
LLVMValueRef temp, cond2;
temp = emit_intrin_1f_param(&ctx->ac, "llvm.fabs",
- ctx->f32, result);
+ ctx->ac.f32, result);
cond = LLVMBuildFCmp(ctx->builder, LLVMRealUGT,
- LLVMBuildBitCast(ctx->builder, LLVMConstInt(ctx->i32, 0x38800000, false), ctx->f32, ""),
+ LLVMBuildBitCast(ctx->builder, LLVMConstInt(ctx->ac.i32, 0x38800000, false), ctx->ac.f32, ""),
temp, "");
cond2 = LLVMBuildFCmp(ctx->builder, LLVMRealUNE,
- temp, ctx->f32zero, "");
+ temp, ctx->ac.f32_0, "");
cond = LLVMBuildAnd(ctx->builder, cond, cond2, "");
- result = LLVMBuildSelect(ctx->builder, cond, ctx->f32zero, result, "");
+ result = LLVMBuildSelect(ctx->builder, cond, ctx->ac.f32_0, result, "");
}
return result;
}
static LLVMValueRef emit_pack_half_2x16(struct ac_llvm_context *ctx,
LLVMValueRef src0)
{
- LLVMValueRef const16 = LLVMConstInt(ctx->i32, 16, false);
- int i;
LLVMValueRef comp[2];
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++) {
- comp[i] = LLVMBuildFPTrunc(ctx->builder, comp[i], ctx->f16, "");
- comp[i] = LLVMBuildBitCast(ctx->builder, comp[i], ctx->i16, "");
- comp[i] = LLVMBuildZExt(ctx->builder, comp[i], ctx->i32, "");
- }
-
- comp[1] = LLVMBuildShl(ctx->builder, comp[1], const16, "");
- comp[0] = LLVMBuildOr(ctx->builder, comp[0], comp[1], "");
- return comp[0];
+ return ac_build_cvt_pkrtz_f16(ctx, comp);
}
static LLVMValueRef emit_unpack_half_2x16(struct ac_llvm_context *ctx,
temps[i] = LLVMBuildFPExt(ctx->builder, val, ctx->f32, "");
}
- LLVMTypeRef v2f32 = LLVMVectorType(ctx->f32, 2);
- result = LLVMBuildInsertElement(ctx->builder, LLVMGetUndef(v2f32), temps[0],
+ result = LLVMBuildInsertElement(ctx->builder, LLVMGetUndef(ctx->v2f32), temps[0],
ctx->i32_0, "");
result = LLVMBuildInsertElement(ctx->builder, result, temps[1],
ctx->i32_1, "");
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] = 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] = ac_to_float(&ctx->ac, src[0]);
- result = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, src[0]);
+ result = ac_build_fdiv(&ctx->ac, instr->dest.dest.ssa.bit_size == 32 ? ctx->ac.f32_1 : ctx->ac.f64_1,
+ src[0]);
break;
case nir_op_iand:
result = LLVMBuildAnd(ctx->ac.builder, src[0], src[1], "");
result = emit_minmax_int(&ctx->ac, LLVMIntULT, src[0], src[1]);
break;
case nir_op_isign:
- result = emit_isign(&ctx->ac, src[0]);
+ result = emit_isign(&ctx->ac, src[0], instr->dest.dest.ssa.bit_size);
break;
case nir_op_fsign:
src[0] = ac_to_float(&ctx->ac, src[0]);
- result = emit_fsign(&ctx->ac, src[0]);
+ result = emit_fsign(&ctx->ac, src[0], instr->dest.dest.ssa.bit_size);
break;
case nir_op_ffloor:
result = emit_intrin_1f_param(&ctx->ac, "llvm.floor",
ac_to_float_type(&ctx->ac, def_type),src[0]);
break;
case nir_op_ffract:
- result = emit_ffract(&ctx->ac, src[0]);
+ result = emit_ffract(&ctx->ac, src[0], instr->dest.dest.ssa.bit_size);
break;
case nir_op_fsin:
result = emit_intrin_1f_param(&ctx->ac, "llvm.sin",
case nir_op_frsq:
result = emit_intrin_1f_param(&ctx->ac, "llvm.sqrt",
ac_to_float_type(&ctx->ac, def_type), src[0]);
- result = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, result);
+ result = ac_build_fdiv(&ctx->ac, instr->dest.dest.ssa.bit_size == 32 ? ctx->ac.f32_1 : ctx->ac.f64_1,
+ result);
break;
case nir_op_fpow:
result = emit_intrin_2f_param(&ctx->ac, "llvm.pow",
case nir_op_fmax:
result = emit_intrin_2f_param(&ctx->ac, "llvm.maxnum",
ac_to_float_type(&ctx->ac, def_type), src[0], src[1]);
- if (instr->dest.dest.ssa.bit_size == 32)
+ if (ctx->ac.chip_class < GFX9 &&
+ instr->dest.dest.ssa.bit_size == 32) {
+ /* Only pre-GFX9 chips do not flush denorms. */
result = emit_intrin_1f_param(&ctx->ac, "llvm.canonicalize",
ac_to_float_type(&ctx->ac, def_type),
result);
+ }
break;
case nir_op_fmin:
result = emit_intrin_2f_param(&ctx->ac, "llvm.minnum",
ac_to_float_type(&ctx->ac, def_type), src[0], src[1]);
- if (instr->dest.dest.ssa.bit_size == 32)
+ if (ctx->ac.chip_class < GFX9 &&
+ instr->dest.dest.ssa.bit_size == 32) {
+ /* Only pre-GFX9 chips do not flush denorms. */
result = emit_intrin_1f_param(&ctx->ac, "llvm.canonicalize",
ac_to_float_type(&ctx->ac, def_type),
result);
+ }
break;
case nir_op_ffma:
result = emit_intrin_3f_param(&ctx->ac, "llvm.fmuladd",
result = LLVMBuildUIToFP(ctx->ac.builder, src[0], ac_to_float_type(&ctx->ac, def_type), "");
break;
case nir_op_f2f64:
+ src[0] = ac_to_float(&ctx->ac, src[0]);
result = LLVMBuildFPExt(ctx->ac.builder, src[0], ac_to_float_type(&ctx->ac, def_type), "");
break;
case nir_op_f2f32:
break;
case nir_op_find_lsb:
src[0] = ac_to_integer(&ctx->ac, src[0]);
- result = emit_find_lsb(&ctx->ac, src[0]);
+ result = ac_find_lsb(&ctx->ac, ctx->ac.i32, 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]);
+ result = ac_build_umsb(&ctx->ac, src[0], ctx->ac.i32);
break;
case nir_op_ifind_msb:
src[0] = ac_to_integer(&ctx->ac, src[0]);
- result = emit_ifind_msb(&ctx->ac, src[0]);
+ result = ac_build_imsb(&ctx->ac, src[0], ctx->ac.i32);
break;
case nir_op_uadd_carry:
src[0] = ac_to_integer(&ctx->ac, src[0]);
result = emit_f2b(&ctx->ac, src[0]);
break;
case nir_op_b2i:
- result = emit_b2i(&ctx->ac, src[0]);
+ result = emit_b2i(&ctx->ac, src[0], instr->dest.dest.ssa.bit_size);
break;
case nir_op_i2b:
src[0] = ac_to_integer(&ctx->ac, src[0]);
case nir_op_unpack_64_2x32_split_x: {
assert(instr->src[0].src.ssa->num_components == 1);
LLVMValueRef tmp = LLVMBuildBitCast(ctx->ac.builder, src[0],
- LLVMVectorType(ctx->ac.i32, 2),
+ ctx->ac.v2i32,
"");
result = LLVMBuildExtractElement(ctx->ac.builder, tmp,
ctx->ac.i32_0, "");
case nir_op_unpack_64_2x32_split_y: {
assert(instr->src[0].src.ssa->num_components == 1);
LLVMValueRef tmp = LLVMBuildBitCast(ctx->ac.builder, src[0],
- LLVMVectorType(ctx->ac.i32, 2),
+ ctx->ac.v2i32,
"");
result = LLVMBuildExtractElement(ctx->ac.builder, tmp,
ctx->ac.i32_1, "");
}
case nir_op_pack_64_2x32_split: {
- LLVMValueRef tmp = LLVMGetUndef(LLVMVectorType(ctx->ac.i32, 2));
+ LLVMValueRef tmp = LLVMGetUndef(ctx->ac.v2i32);
tmp = LLVMBuildInsertElement(ctx->ac.builder, tmp,
src[0], ctx->ac.i32_0, "");
tmp = LLVMBuildInsertElement(ctx->ac.builder, tmp,
*/
LLVMValueRef stride =
LLVMBuildExtractElement(ctx->ac.builder, descriptor,
- LLVMConstInt(ctx->ac.i32, 1, false), "");
+ ctx->ac.i32_1, "");
stride = LLVMBuildLShr(ctx->ac.builder, stride,
LLVMConstInt(ctx->ac.i32, 16, false), "");
stride = LLVMBuildAnd(ctx->ac.builder, stride,
struct ac_image_args *args)
{
if (instr->sampler_dim == GLSL_SAMPLER_DIM_BUF) {
+ unsigned mask = nir_ssa_def_components_read(&instr->dest.ssa);
+
return ac_build_buffer_load_format(&ctx->ac,
args->resource,
args->addr,
- LLVMConstInt(ctx->ac.i32, 0, false),
- true);
+ ctx->ac.i32_0,
+ util_last_bit(mask),
+ false, true);
}
args->opcode = ac_image_sample;
case nir_texop_txf:
case nir_texop_txf_ms:
case nir_texop_samples_identical:
- args->opcode = instr->sampler_dim == GLSL_SAMPLER_DIM_MS ? ac_image_load : ac_image_load_mip;
+ args->opcode = lod_is_zero ||
+ instr->sampler_dim == GLSL_SAMPLER_DIM_MS ?
+ ac_image_load : ac_image_load_mip;
args->compare = false;
args->offset = false;
break;
layout->binding[binding].dynamic_offset_offset;
desc_ptr = ctx->push_constants;
base_offset = pipeline_layout->push_constant_size + 16 * idx;
- stride = LLVMConstInt(ctx->i32, 16, false);
+ stride = LLVMConstInt(ctx->ac.i32, 16, false);
} else
- stride = LLVMConstInt(ctx->i32, layout->binding[binding].size, false);
+ stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
- offset = LLVMConstInt(ctx->i32, base_offset, false);
+ offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
index = LLVMBuildMul(ctx->builder, index, stride, "");
offset = LLVMBuildAdd(ctx->builder, offset, index, "");
desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
- desc_ptr = cast_ptr(ctx, desc_ptr, ctx->v4i32);
- LLVMSetMetadata(desc_ptr, ctx->uniform_md_kind, ctx->empty_md);
+ desc_ptr = cast_ptr(ctx, desc_ptr, ctx->ac.v4i32);
+ LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
- return LLVMBuildLoad(ctx->builder, desc_ptr, "");
+ return desc_ptr;
+}
+
+static LLVMValueRef visit_vulkan_resource_reindex(struct nir_to_llvm_context *ctx,
+ nir_intrinsic_instr *instr)
+{
+ LLVMValueRef ptr = get_src(ctx->nir, instr->src[0]);
+ LLVMValueRef index = get_src(ctx->nir, instr->src[1]);
+
+ LLVMValueRef result = LLVMBuildGEP(ctx->builder, ptr, &index, 1, "");
+ LLVMSetMetadata(result, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
+ return result;
}
static LLVMValueRef visit_load_push_constant(struct nir_to_llvm_context *ctx,
{
LLVMValueRef ptr, addr;
- addr = LLVMConstInt(ctx->i32, nir_intrinsic_base(instr), 0);
+ addr = LLVMConstInt(ctx->ac.i32, nir_intrinsic_base(instr), 0);
addr = LLVMBuildAdd(ctx->builder, addr, get_src(ctx->nir, instr->src[0]), "");
ptr = ac_build_gep0(&ctx->ac, ctx->push_constants, addr);
static LLVMValueRef visit_get_buffer_size(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
- LLVMValueRef desc = get_src(ctx, instr->src[0]);
+ LLVMValueRef index = get_src(ctx, instr->src[0]);
+
+ return get_buffer_size(ctx, ctx->abi->load_ssbo(ctx->abi, index, false), false);
+}
- return get_buffer_size(ctx, desc, false);
+static uint32_t widen_mask(uint32_t mask, unsigned multiplier)
+{
+ uint32_t new_mask = 0;
+ for(unsigned i = 0; i < 32 && (1u << i) <= mask; ++i)
+ if (mask & (1u << i))
+ new_mask |= ((1u << multiplier) - 1u) << (i * multiplier);
+ return new_mask;
}
+
+static LLVMValueRef extract_vector_range(struct ac_llvm_context *ctx, LLVMValueRef src,
+ unsigned start, unsigned count)
+{
+ LLVMTypeRef type = LLVMTypeOf(src);
+
+ if (LLVMGetTypeKind(type) != LLVMVectorTypeKind) {
+ assert(start == 0);
+ assert(count == 1);
+ return src;
+ }
+
+ unsigned src_elements = LLVMGetVectorSize(type);
+ assert(start < src_elements);
+ assert(start + count <= src_elements);
+
+ if (start == 0 && count == src_elements)
+ return src;
+
+ if (count == 1)
+ return LLVMBuildExtractElement(ctx->builder, src, LLVMConstInt(ctx->i32, start, false), "");
+
+ assert(count <= 8);
+ LLVMValueRef indices[8];
+ for (unsigned i = 0; i < count; ++i)
+ indices[i] = LLVMConstInt(ctx->i32, start + i, false);
+
+ LLVMValueRef swizzle = LLVMConstVector(indices, count);
+ return LLVMBuildShuffleVector(ctx->builder, src, src, swizzle, "");
+}
+
static void visit_store_ssbo(struct ac_nir_context *ctx,
nir_intrinsic_instr *instr)
{
unsigned writemask = nir_intrinsic_write_mask(instr);
LLVMValueRef base_data, base_offset;
LLVMValueRef params[6];
- LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
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 */
+ params[2] = ctx->ac.i32_0; /* vindex */
+ params[4] = ctx->ac.i1false; /* glc */
+ params[5] = ctx->ac.i1false; /* slc */
if (components_32bit > 1)
data_type = LLVMVectorType(ctx->ac.f32, components_32bit);
+ writemask = widen_mask(writemask, elem_size_mult);
+
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,
int start, count;
LLVMValueRef data;
LLVMValueRef offset;
- LLVMValueRef tmp;
+
u_bit_scan_consecutive_range(&writemask, &start, &count);
/* Due to an LLVM limitation, split 3-element writes
count = 2;
}
- start *= elem_size_mult;
- count *= elem_size_mult;
-
if (count > 4) {
writemask |= ((1u << (count - 4)) - 1u) << (start + 4);
count = 4;
if (count == 4) {
store_name = "llvm.amdgcn.buffer.store.v4f32";
- data = base_data;
} else if (count == 2) {
- 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->ac.builder, base_data,
- LLVMConstInt(ctx->ac.i32, start, false), "");
- else
- data = base_data;
store_name = "llvm.amdgcn.buffer.store.f32";
}
+ data = extract_vector_range(&ctx->ac, base_data, start, count);
offset = base_offset;
if (start != 0) {
int arg_count = 0;
if (instr->intrinsic == nir_intrinsic_ssbo_atomic_comp_swap) {
- params[arg_count++] = llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[3]), 0);
+ params[arg_count++] = ac_llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[3]), 0);
}
- params[arg_count++] = llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[2]), 0);
+ params[arg_count++] = ac_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++] = ctx->ac.i32_0; /* vindex */
params[arg_count++] = get_src(ctx, instr->src[1]); /* voffset */
params[arg_count++] = LLVMConstInt(ctx->ac.i1, 0, false); /* slc */
else
unreachable("unhandled number of components");
- LLVMValueRef i1false = LLVMConstInt(ctx->ac.i1, 0, false);
LLVMValueRef params[] = {
ctx->abi->load_ssbo(ctx->abi,
get_src(ctx, instr->src[0]),
false),
- LLVMConstInt(ctx->ac.i32, 0, false),
+ ctx->ac.i32_0,
offset,
- i1false,
- i1false,
+ ctx->ac.i1false,
+ ctx->ac.i1false,
};
- results[i] = ac_build_intrinsic(&ctx->ac, load_name, data_type, params, 5, 0);
-
+ results[i > 0 ? 1 : 0] = ac_build_intrinsic(&ctx->ac, load_name, data_type, params, 5, 0);
}
+ assume(results[0]);
LLVMValueRef ret = results[0];
if (num_components > 4 || num_components == 3) {
LLVMValueRef masks[] = {
static LLVMValueRef visit_load_ubo_buffer(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
- LLVMValueRef results[8], ret;
+ LLVMValueRef ret;
LLVMValueRef rsrc = get_src(ctx, instr->src[0]);
LLVMValueRef offset = get_src(ctx, instr->src[1]);
int num_components = instr->num_components;
if (instr->dest.ssa.bit_size == 64)
num_components *= 2;
- for (unsigned i = 0; i < num_components; ++i) {
- LLVMValueRef params[] = {
- rsrc,
- 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->ac.f32,
- params, 2,
- AC_FUNC_ATTR_READNONE |
- AC_FUNC_ATTR_LEGACY);
- }
-
-
- ret = ac_build_gather_values(&ctx->ac, results, instr->num_components);
+ ret = ac_build_buffer_load(&ctx->ac, rsrc, num_components, NULL, offset,
+ NULL, 0, false, false, true, true);
+ ret = trim_vector(&ctx->ac, ret, num_components);
return LLVMBuildBitCast(ctx->ac.builder, ret,
get_def_type(ctx, &instr->dest.ssa), "");
}
*indir_out = offset;
}
-static LLVMValueRef
-lds_load(struct nir_to_llvm_context *ctx,
- LLVMValueRef dw_addr)
-{
- LLVMValueRef value;
- value = ac_build_load(&ctx->ac, ctx->lds, dw_addr);
- return value;
-}
-
-static void
-lds_store(struct nir_to_llvm_context *ctx,
- LLVMValueRef dw_addr, LLVMValueRef value)
-{
- value = LLVMBuildBitCast(ctx->builder, value, ctx->i32, "");
- ac_build_indexed_store(&ctx->ac, ctx->lds,
- dw_addr, value);
-}
/* The offchip buffer layout for TCS->TES is
*
total_vertices = LLVMBuildMul(ctx->builder, vertices_per_patch,
num_patches, "");
- constant16 = LLVMConstInt(ctx->i32, 16, false);
+ constant16 = LLVMConstInt(ctx->ac.i32, 16, false);
if (vertex_index) {
base_addr = LLVMBuildMul(ctx->builder, rel_patch_id,
vertices_per_patch, "");
LLVMValueRef param_index;
if (indir_index)
- param_index = LLVMBuildAdd(ctx->builder, LLVMConstInt(ctx->i32, param, false),
+ param_index = LLVMBuildAdd(ctx->builder, LLVMConstInt(ctx->ac.i32, param, false),
indir_index, "");
else {
if (const_index && !is_compact)
param += const_index;
- param_index = LLVMConstInt(ctx->i32, param, false);
+ param_index = LLVMConstInt(ctx->ac.i32, param, false);
}
return get_tcs_tes_buffer_address(ctx, vertex_index, param_index);
}
if (indir_index)
dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
LLVMBuildMul(ctx->builder, indir_index,
- LLVMConstInt(ctx->i32, 4, false), ""), "");
+ LLVMConstInt(ctx->ac.i32, 4, false), ""), "");
else if (const_index && !compact_const_index)
dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
- LLVMConstInt(ctx->i32, const_index, false), "");
+ LLVMConstInt(ctx->ac.i32, const_index, false), "");
dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
- LLVMConstInt(ctx->i32, param * 4, false), "");
+ LLVMConstInt(ctx->ac.i32, param * 4, false), "");
if (const_index && compact_const_index)
dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
- LLVMConstInt(ctx->i32, const_index, false), "");
+ LLVMConstInt(ctx->ac.i32, const_index, false), "");
return dw_addr;
}
static LLVMValueRef
-load_tcs_input(struct nir_to_llvm_context *ctx,
- nir_intrinsic_instr *instr)
+load_tcs_varyings(struct ac_shader_abi *abi,
+ LLVMValueRef vertex_index,
+ LLVMValueRef indir_index,
+ unsigned const_index,
+ unsigned location,
+ unsigned driver_location,
+ unsigned component,
+ unsigned num_components,
+ bool is_patch,
+ bool is_compact,
+ bool load_input)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
LLVMValueRef dw_addr, stride;
- unsigned const_index;
- LLVMValueRef vertex_index;
- LLVMValueRef indir_index;
- unsigned param;
LLVMValueRef value[4], result;
- 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);
- get_deref_offset(ctx->nir, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
-
- 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);
-
- for (unsigned i = 0; i < instr->num_components; i++) {
- value[i] = lds_load(ctx, dw_addr);
- dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
- ctx->i32one, "");
- }
- result = ac_build_gather_values(&ctx->ac, value, instr->num_components);
- result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), "");
- return result;
-}
-
-static LLVMValueRef
-load_tcs_output(struct nir_to_llvm_context *ctx,
- nir_intrinsic_instr *instr)
-{
- LLVMValueRef dw_addr;
- LLVMValueRef stride = NULL;
- LLVMValueRef value[4], result;
- LLVMValueRef vertex_index = NULL;
- LLVMValueRef indir_index = NULL;
- unsigned const_index = 0;
- unsigned param;
- 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);
- get_deref_offset(ctx->nir, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
+ unsigned param = shader_io_get_unique_index(location);
- if (!instr->variables[0]->var->data.patch) {
- stride = unpack_param(&ctx->ac, ctx->tcs_out_layout, 13, 8);
- dw_addr = get_tcs_out_current_patch_offset(ctx);
+ if (load_input) {
+ stride = unpack_param(&ctx->ac, ctx->tcs_in_layout, 13, 8);
+ dw_addr = get_tcs_in_current_patch_offset(ctx);
} else {
- dw_addr = get_tcs_out_current_patch_data_offset(ctx);
+ if (!is_patch) {
+ 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);
+ stride = NULL;
+ }
}
dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
indir_index);
- for (unsigned i = 0; i < instr->num_components; i++) {
- value[i] = lds_load(ctx, dw_addr);
+ for (unsigned i = 0; i < num_components + component; i++) {
+ value[i] = ac_lds_load(&ctx->ac, dw_addr);
dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
- ctx->i32one, "");
+ ctx->ac.i32_1, "");
}
- result = ac_build_gather_values(&ctx->ac, value, instr->num_components);
- result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), "");
+ result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
return result;
}
static void
-store_tcs_output(struct nir_to_llvm_context *ctx,
- nir_intrinsic_instr *instr,
+store_tcs_output(struct ac_shader_abi *abi,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned const_index,
+ unsigned location,
+ unsigned driver_location,
LLVMValueRef src,
+ unsigned component,
+ bool is_patch,
+ bool is_compact,
unsigned writemask)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
LLVMValueRef dw_addr;
LLVMValueRef stride = NULL;
LLVMValueRef buf_addr = NULL;
- LLVMValueRef vertex_index = NULL;
- LLVMValueRef indir_index = NULL;
- unsigned const_index = 0;
unsigned param;
- 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;
+ bool store_lds = true;
- get_deref_offset(ctx->nir, instr->variables[0],
- false, NULL, per_vertex ? &vertex_index : NULL,
- &const_index, &indir_index);
+ if (is_patch) {
+ if (!(ctx->tcs_patch_outputs_read & (1U << (location - VARYING_SLOT_PATCH0))))
+ store_lds = false;
+ } else {
+ if (!(ctx->tcs_outputs_read & (1ULL << location)))
+ store_lds = false;
+ }
- param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
- if (instr->variables[0]->var->data.location == VARYING_SLOT_CLIP_DIST0 &&
+ param = shader_io_get_unique_index(location);
+ if (location == VARYING_SLOT_CLIP_DIST0 &&
is_compact && const_index > 3) {
const_index -= 3;
param++;
}
- if (!instr->variables[0]->var->data.patch) {
+ if (!is_patch) {
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);
}
- mark_tess_output(ctx, instr->variables[0]->var->data.patch, param);
+ mark_tess_output(ctx, is_patch, param);
dw_addr = get_dw_address(ctx, dw_addr, param, const_index, is_compact, vertex_index, stride,
- indir_index);
+ param_index);
buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index, is_compact,
- vertex_index, indir_index);
+ vertex_index, param_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)
+ if (location == VARYING_SLOT_TESS_LEVEL_INNER ||
+ location == VARYING_SLOT_TESS_LEVEL_OUTER)
is_tess_factor = true;
unsigned base = is_compact ? const_index : 0;
for (unsigned chan = 0; chan < 8; chan++) {
if (!(writemask & (1 << chan)))
continue;
- LLVMValueRef value = llvm_extract_elem(&ctx->ac, src, chan);
+ LLVMValueRef value = ac_llvm_extract_elem(&ctx->ac, src, chan - component);
- lds_store(ctx, dw_addr, value);
+ if (store_lds || is_tess_factor) {
+ LLVMValueRef dw_addr_chan =
+ LLVMBuildAdd(ctx->builder, dw_addr,
+ LLVMConstInt(ctx->ac.i32, chan, false), "");
+ ac_lds_store(&ctx->ac, dw_addr_chan, value);
+ }
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,
4 * (base + chan), 1, 0, true, false);
-
- dw_addr = LLVMBuildAdd(ctx->builder, dw_addr,
- ctx->i32one, "");
}
if (writemask == 0xF) {
}
static LLVMValueRef
-load_tes_input(struct nir_to_llvm_context *ctx,
- const nir_intrinsic_instr *instr)
+load_tes_input(struct ac_shader_abi *abi,
+ LLVMValueRef vertex_index,
+ LLVMValueRef param_index,
+ unsigned const_index,
+ unsigned location,
+ unsigned driver_location,
+ unsigned component,
+ unsigned num_components,
+ bool is_patch,
+ bool is_compact,
+ bool load_input)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
LLVMValueRef buf_addr;
LLVMValueRef result;
- LLVMValueRef vertex_index = NULL;
- LLVMValueRef indir_index = NULL;
- unsigned const_index = 0;
- unsigned param;
- 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;
+ unsigned param = shader_io_get_unique_index(location);
- 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) {
+ if (location == VARYING_SLOT_CLIP_DIST0 && is_compact && const_index > 3) {
const_index -= 3;
param++;
}
+
buf_addr = get_tcs_tes_buffer_address_params(ctx, param, const_index,
- is_compact, vertex_index, indir_index);
+ is_compact, vertex_index, param_index);
+
+ LLVMValueRef comp_offset = LLVMConstInt(ctx->ac.i32, component * 4, false);
+ buf_addr = LLVMBuildAdd(ctx->builder, buf_addr, comp_offset, "");
- result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, instr->num_components, NULL,
+ result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, num_components, NULL,
buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false);
- result = trim_vector(&ctx->ac, result, instr->num_components);
- result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx->nir, &instr->dest.ssa), "");
+ result = trim_vector(&ctx->ac, result, num_components);
return result;
}
static LLVMValueRef
-load_gs_input(struct nir_to_llvm_context *ctx,
- nir_intrinsic_instr *instr)
+load_gs_input(struct ac_shader_abi *abi,
+ unsigned location,
+ unsigned driver_location,
+ unsigned component,
+ unsigned num_components,
+ unsigned vertex_index,
+ unsigned const_index,
+ LLVMTypeRef type)
{
- LLVMValueRef indir_index, vtx_offset;
- unsigned const_index;
- LLVMValueRef args[9];
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+ LLVMValueRef vtx_offset;
unsigned param, vtx_offset_param;
LLVMValueRef value[4], result;
- unsigned vertex_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],
- LLVMConstInt(ctx->i32, 4, false), "");
+ LLVMConstInt(ctx->ac.i32, 4, false), "");
- param = shader_io_get_unique_index(instr->variables[0]->var->data.location);
- for (unsigned i = 0; i < instr->num_components; i++) {
+ param = shader_io_get_unique_index(location);
+
+ for (unsigned i = component; i < num_components + component; i++) {
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);
+ LLVMConstInt(ctx->ac.i32, param * 4 + i + const_index, 0), "");
+ value[i] = ac_lds_load(&ctx->ac, 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);
+ LLVMValueRef soffset =
+ LLVMConstInt(ctx->ac.i32,
+ (param * 4 + i + const_index) * 256,
+ false);
+
+ value[i] = ac_build_buffer_load(&ctx->ac,
+ ctx->esgs_ring, 1,
+ ctx->ac.i32_0,
+ vtx_offset, soffset,
+ 0, 1, 0, true, false);
+
+ value[i] = LLVMBuildBitCast(ctx->builder, value[i],
+ type, "");
}
}
- result = ac_build_gather_values(&ctx->ac, value, instr->num_components);
-
+ result = ac_build_varying_gather_values(&ctx->ac, value, num_components, component);
+ result = ac_to_integer(&ctx->ac, result);
return result;
}
return val;
}
+static LLVMValueRef load_tess_varyings(struct ac_nir_context *ctx,
+ nir_intrinsic_instr *instr,
+ bool load_inputs)
+{
+ LLVMValueRef result;
+ LLVMValueRef vertex_index = NULL;
+ LLVMValueRef indir_index = NULL;
+ unsigned const_index = 0;
+ unsigned location = instr->variables[0]->var->data.location;
+ unsigned driver_location = instr->variables[0]->var->data.driver_location;
+ const bool is_patch = instr->variables[0]->var->data.patch;
+ const bool is_compact = instr->variables[0]->var->data.compact;
+
+ get_deref_offset(ctx, instr->variables[0],
+ false, NULL, is_patch ? NULL : &vertex_index,
+ &const_index, &indir_index);
+
+ result = ctx->abi->load_tess_varyings(ctx->abi, vertex_index, indir_index,
+ const_index, location, driver_location,
+ instr->variables[0]->var->data.location_frac,
+ instr->num_components,
+ is_patch, is_compact, load_inputs);
+ return LLVMBuildBitCast(ctx->ac.builder, result, get_def_type(ctx, &instr->dest.ssa), "");
+}
+
static LLVMValueRef visit_load_var(struct ac_nir_context *ctx,
nir_intrinsic_instr *instr)
{
LLVMValueRef values[8];
int idx = instr->variables[0]->var->data.driver_location;
int ve = instr->dest.ssa.num_components;
+ unsigned comp = instr->variables[0]->var->data.location_frac;
LLVMValueRef indir_index;
LLVMValueRef ret;
unsigned const_index;
+ unsigned stride = instr->variables[0]->var->data.compact ? 1 : 4;
bool vs_in = ctx->stage == MESA_SHADER_VERTEX &&
instr->variables[0]->var->data.mode == nir_var_shader_in;
get_deref_offset(ctx, instr->variables[0], vs_in, NULL, NULL,
switch (instr->variables[0]->var->data.mode) {
case nir_var_shader_in:
- if (ctx->stage == MESA_SHADER_TESS_CTRL)
- return load_tcs_input(ctx->nctx, instr);
- if (ctx->stage == MESA_SHADER_TESS_EVAL)
- return load_tes_input(ctx->nctx, instr);
+ if (ctx->stage == MESA_SHADER_TESS_CTRL ||
+ ctx->stage == MESA_SHADER_TESS_EVAL) {
+ return load_tess_varyings(ctx, instr, true);
+ }
+
if (ctx->stage == MESA_SHADER_GEOMETRY) {
- return load_gs_input(ctx->nctx, instr);
+ LLVMTypeRef type = LLVMIntTypeInContext(ctx->ac.context, instr->dest.ssa.bit_size);
+ LLVMValueRef indir_index;
+ unsigned const_index, vertex_index;
+ get_deref_offset(ctx, instr->variables[0],
+ false, &vertex_index, NULL,
+ &const_index, &indir_index);
+
+ return ctx->abi->load_inputs(ctx->abi, instr->variables[0]->var->data.location,
+ instr->variables[0]->var->data.driver_location,
+ instr->variables[0]->var->data.location_frac, ve,
+ vertex_index, const_index, type);
}
- for (unsigned chan = 0; chan < ve; chan++) {
+
+ for (unsigned chan = comp; chan < ve + comp; chan++) {
if (indir_index) {
unsigned count = glsl_count_attribute_slots(
instr->variables[0]->var->type,
count -= chan / 4;
LLVMValueRef tmp_vec = ac_build_gather_values_extended(
&ctx->ac, ctx->abi->inputs + idx + chan, count,
- 4, false, true);
+ stride, false, true);
values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
tmp_vec,
indir_index, "");
} else
- values[chan] = ctx->abi->inputs[idx + chan + const_index * 4];
+ values[chan] = ctx->abi->inputs[idx + chan + const_index * stride];
}
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, true);
+ stride, true, true);
values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
tmp_vec,
indir_index, "");
} else {
- values[chan] = LLVMBuildLoad(ctx->ac.builder, ctx->locals[idx + chan + const_index * 4], "");
+ values[chan] = LLVMBuildLoad(ctx->ac.builder, ctx->locals[idx + chan + const_index * stride], "");
}
}
break;
"");
}
case nir_var_shader_out:
- if (ctx->stage == MESA_SHADER_TESS_CTRL)
- return load_tcs_output(ctx->nctx, instr);
- for (unsigned chan = 0; chan < ve; chan++) {
+ if (ctx->stage == MESA_SHADER_TESS_CTRL) {
+ return load_tess_varyings(ctx, instr, false);
+ }
+
+ for (unsigned chan = comp; chan < ve + comp; 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->outputs + idx + chan, count,
- 4, true, true);
+ stride, true, true);
values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
tmp_vec,
indir_index, "");
} else {
values[chan] = LLVMBuildLoad(ctx->ac.builder,
- ctx->outputs[idx + chan + const_index * 4],
+ ctx->outputs[idx + chan + const_index * stride],
"");
}
}
default:
unreachable("unhandle variable mode");
}
- ret = ac_build_gather_values(&ctx->ac, values, ve);
+ ret = ac_build_varying_gather_values(&ctx->ac, values, ve, comp);
return LLVMBuildBitCast(ctx->ac.builder, ret, get_def_type(ctx, &instr->dest.ssa), "");
}
{
LLVMValueRef temp_ptr, value;
int idx = instr->variables[0]->var->data.driver_location;
+ unsigned comp = instr->variables[0]->var->data.location_frac;
LLVMValueRef src = ac_to_float(&ctx->ac, get_src(ctx, instr->src[0]));
- int writemask = instr->const_index[0];
+ int writemask = instr->const_index[0] << comp;
LLVMValueRef indir_index;
unsigned const_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->ac.builder, src,
- LLVMVectorType(ctx->ac.f32, get_llvm_num_components(src) * 2),
+ LLVMVectorType(ctx->ac.f32, ac_get_llvm_num_components(src) * 2),
"");
- writemask = 0;
- for (unsigned chan = 0; chan < 4; chan++) {
- if (old_writemask & (1 << chan))
- writemask |= 3u << (2 * chan);
- }
+ writemask = widen_mask(writemask, 2);
}
switch (instr->variables[0]->var->data.mode) {
case nir_var_shader_out:
if (ctx->stage == MESA_SHADER_TESS_CTRL) {
- store_tcs_output(ctx->nctx, instr, src, writemask);
+ LLVMValueRef vertex_index = NULL;
+ LLVMValueRef indir_index = NULL;
+ unsigned const_index = 0;
+ const unsigned location = instr->variables[0]->var->data.location;
+ const unsigned driver_location = instr->variables[0]->var->data.driver_location;
+ const unsigned comp = instr->variables[0]->var->data.location_frac;
+ const bool is_patch = instr->variables[0]->var->data.patch;
+ const bool is_compact = instr->variables[0]->var->data.compact;
+
+ get_deref_offset(ctx, instr->variables[0],
+ false, NULL, is_patch ? NULL : &vertex_index,
+ &const_index, &indir_index);
+
+ ctx->abi->store_tcs_outputs(ctx->abi, vertex_index, indir_index,
+ const_index, location, driver_location,
+ src, comp, is_patch, is_compact, writemask);
return;
}
if (!(writemask & (1 << chan)))
continue;
- value = llvm_extract_elem(&ctx->ac, src, chan);
+ value = ac_llvm_extract_elem(&ctx->ac, src, chan - comp);
if (instr->variables[0]->var->data.compact)
stride = 1;
if (!(writemask & (1 << chan)))
continue;
- value = llvm_extract_elem(&ctx->ac, src, chan);
+ value = ac_llvm_extract_elem(&ctx->ac, src, chan);
if (indir_index) {
unsigned count = glsl_count_attribute_slots(
instr->variables[0]->var->type, false);
LLVMValueRef ptr =
LLVMBuildStructGEP(ctx->ac.builder,
address, chan, "");
- LLVMValueRef src = llvm_extract_elem(&ctx->ac, val,
- chan);
+ LLVMValueRef src = ac_llvm_extract_elem(&ctx->ac, val,
+ chan);
src = LLVMBuildBitCast(
ctx->ac.builder, src,
LLVMGetElementType(LLVMTypeOf(ptr)), "");
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;
- if(instr->variables[0]->deref.child)
- type = instr->variables[0]->deref.child->type;
+ const struct glsl_type *type = glsl_without_array(instr->variables[0]->var->type);
LLVMValueRef src0 = get_src(ctx, instr->src[0]);
LLVMValueRef coords[4];
LLVMConstInt(ctx->ac.i32, 2, false), LLVMConstInt(ctx->ac.i32, 3, false),
};
LLVMValueRef res;
- LLVMValueRef sample_index = llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[1]), 0);
+ LLVMValueRef sample_index = ac_llvm_extract_elem(&ctx->ac, get_src(ctx, instr->src[1]), 0);
int count;
enum glsl_sampler_dim dim = glsl_get_sampler_dim(type);
fmask_load_address[1],
fmask_load_address[2],
sample_index,
- get_sampler_desc(ctx, instr->variables[0], AC_DESC_FMASK, true, false));
+ get_sampler_desc(ctx, instr->variables[0], AC_DESC_FMASK, NULL, true, false));
}
if (count == 1 && !gfx9_1d) {
if (instr->src[0].ssa->num_components)
if (is_ms)
count--;
for (chan = 0; chan < count; ++chan) {
- coords[chan] = llvm_extract_elem(&ctx->ac, src0, chan);
+ coords[chan] = ac_llvm_extract_elem(&ctx->ac, src0, chan);
}
if (add_frag_pos) {
for (chan = 0; chan < 2; ++chan)
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], 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);
+
+ const enum glsl_sampler_dim dim = glsl_get_sampler_dim(type);
+ if (dim == GLSL_SAMPLER_DIM_BUF) {
+ unsigned mask = nir_ssa_def_components_read(&instr->dest.ssa);
+ unsigned num_channels = util_last_bit(mask);
+ LLVMValueRef rsrc, vindex;
+
+ rsrc = get_sampler_desc(ctx, instr->variables[0], AC_DESC_BUFFER, NULL, true, false);
+ vindex = LLVMBuildExtractElement(ctx->ac.builder, get_src(ctx, instr->src[0]),
+ ctx->ac.i32_0, "");
+
+ /* TODO: set "glc" and "can_speculate" when OpenGL needs it. */
+ res = ac_build_buffer_load_format(&ctx->ac, rsrc, vindex,
+ ctx->ac.i32_0, num_channels,
+ false, false);
+ res = ac_build_expand_to_vec4(&ctx->ac, res, num_channels);
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 ||
- 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;
+ dim == GLSL_SAMPLER_DIM_CUBE ||
+ dim == GLSL_SAMPLER_DIM_3D ||
+ dim == GLSL_SAMPLER_DIM_SUBPASS ||
+ dim == GLSL_SAMPLER_DIM_SUBPASS_MS;
+ LLVMValueRef da = is_da ? ctx->ac.i1true : ctx->ac.i1false;
+ LLVMValueRef glc = ctx->ac.i1false;
+ LLVMValueRef slc = ctx->ac.i1false;
params[0] = get_image_coords(ctx, instr);
- params[1] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, false);
+ params[1] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, NULL, true, false);
params[2] = LLVMConstInt(ctx->ac.i32, 15, false); /* dmask */
- if (HAVE_LLVM <= 0x0309) {
- params[3] = i1false; /* r128 */
- params[4] = da;
- params[5] = glc;
- params[6] = slc;
- } else {
- LLVMValueRef lwe = i1false;
- params[3] = glc;
- params[4] = slc;
- params[5] = lwe;
- params[6] = da;
- }
+ params[3] = glc;
+ params[4] = slc;
+ params[5] = ctx->ac.i1false;
+ params[6] = da;
ac_get_image_intr_name("llvm.amdgcn.image.load",
ctx->ac.v4f32, /* vdata */
char intrinsic_name[64];
const nir_variable *var = instr->variables[0]->var;
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);
- LLVMValueRef glc = i1false;
+ const enum glsl_sampler_dim dim = glsl_get_sampler_dim(type);
+ LLVMValueRef glc = ctx->ac.i1false;
bool force_glc = ctx->ac.chip_class == SI;
if (force_glc)
- glc = i1true;
+ glc = ctx->ac.i1true;
- if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) {
+ if (dim == GLSL_SAMPLER_DIM_BUF) {
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[1] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_BUFFER, NULL, 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] = i1false; /* slc */
+ params[5] = ctx->ac.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 ? i1true : i1false;
- LLVMValueRef slc = i1false;
+ dim == GLSL_SAMPLER_DIM_CUBE ||
+ dim == GLSL_SAMPLER_DIM_3D;
+ LLVMValueRef da = is_da ? ctx->ac.i1true : ctx->ac.i1false;
+ LLVMValueRef slc = ctx->ac.i1false;
params[0] = ac_to_float(&ctx->ac, get_src(ctx, instr->src[2]));
params[1] = get_image_coords(ctx, instr); /* coords */
- params[2] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, true);
+ params[2] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, NULL, true, true);
params[3] = LLVMConstInt(ctx->ac.i32, 15, false); /* dmask */
- if (HAVE_LLVM <= 0x0309) {
- params[4] = i1false; /* r128 */
- params[5] = da;
- params[6] = glc;
- params[7] = slc;
- } else {
- LLVMValueRef lwe = i1false;
- params[4] = glc;
- params[5] = slc;
- params[6] = lwe;
- params[7] = da;
- }
+ params[4] = glc;
+ params[5] = slc;
+ params[6] = ctx->ac.i1false;
+ params[7] = da;
ac_get_image_intr_name("llvm.amdgcn.image.store",
LLVMTypeOf(params[0]), /* vdata */
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;
+ bool is_unsigned = glsl_get_sampler_result_type(type) == GLSL_TYPE_UINT;
+
switch (instr->intrinsic) {
case nir_intrinsic_image_atomic_add:
atomic_name = "add";
break;
case nir_intrinsic_image_atomic_min:
- atomic_name = "smin";
+ atomic_name = is_unsigned ? "umin" : "smin";
break;
case nir_intrinsic_image_atomic_max:
- atomic_name = "smax";
+ atomic_name = is_unsigned ? "umax" : "smax";
break;
case nir_intrinsic_image_atomic_and:
atomic_name = "and";
if (glsl_get_sampler_dim(type) == GLSL_SAMPLER_DIM_BUF) {
params[param_count++] = get_sampler_desc(ctx, instr->variables[0], AC_DESC_BUFFER,
- true, true);
+ NULL, 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 */
+ params[param_count++] = ctx->ac.i1false; /* slc */
length = snprintf(intrinsic_name, sizeof(intrinsic_name),
"llvm.amdgcn.buffer.atomic.%s", atomic_name);
LLVMValueRef coords = params[param_count++] = get_image_coords(ctx, instr);
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 */
+ NULL, true, true);
+ params[param_count++] = ctx->ac.i1false; /* r128 */
+ params[param_count++] = da ? ctx->ac.i1true : ctx->ac.i1false; /* da */
+ params[param_count++] = ctx->ac.i1false; /* slc */
build_int_type_name(LLVMTypeOf(coords),
coords_type, sizeof(coords_type));
const nir_variable *var = instr->variables[0]->var;
const struct glsl_type *type = instr->variables[0]->var->type;
bool da = glsl_sampler_type_is_array(var->type) ||
- glsl_get_sampler_dim(var->type) == GLSL_SAMPLER_DIM_CUBE;
+ glsl_get_sampler_dim(var->type) == GLSL_SAMPLER_DIM_CUBE ||
+ glsl_get_sampler_dim(var->type) == GLSL_SAMPLER_DIM_3D;
if(instr->variables[0]->deref.child)
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],
- AC_DESC_BUFFER, true, false), true);
+ AC_DESC_BUFFER, NULL, true, false), true);
struct ac_image_args args = { 0 };
args.da = da;
args.dmask = 0xf;
- args.resource = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, true, false);
+ args.resource = get_sampler_desc(ctx, instr->variables[0], AC_DESC_IMAGE, NULL, true, false);
args.opcode = ac_image_get_resinfo;
args.addr = ctx->ac.i32_0;
#define LGKM_CNT 0x07f
#define VM_CNT 0xf70
-static void emit_waitcnt(struct nir_to_llvm_context *ctx,
- unsigned simm16)
+static void emit_membar(struct nir_to_llvm_context *ctx,
+ const nir_intrinsic_instr *instr)
{
- LLVMValueRef args[1] = {
- LLVMConstInt(ctx->i32, simm16, false),
- };
- ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.s.waitcnt",
- ctx->voidt, args, 1, 0);
+ unsigned waitcnt = NOOP_WAITCNT;
+
+ switch (instr->intrinsic) {
+ case nir_intrinsic_memory_barrier:
+ case nir_intrinsic_group_memory_barrier:
+ waitcnt &= VM_CNT & LGKM_CNT;
+ break;
+ case nir_intrinsic_memory_barrier_atomic_counter:
+ case nir_intrinsic_memory_barrier_buffer:
+ case nir_intrinsic_memory_barrier_image:
+ waitcnt &= VM_CNT;
+ break;
+ case nir_intrinsic_memory_barrier_shared:
+ waitcnt &= LGKM_CNT;
+ break;
+ default:
+ break;
+ }
+ if (waitcnt != NOOP_WAITCNT)
+ ac_build_waitcnt(&ctx->ac, waitcnt);
}
-static void emit_barrier(struct nir_to_llvm_context *ctx)
+static void emit_barrier(struct ac_llvm_context *ac, gl_shader_stage stage)
{
/* SI only (thanks to a hw bug workaround):
* The real barrier instruction isn’t needed, because an entire patch
* always fits into a single wave.
*/
- if (ctx->options->chip_class == SI &&
- ctx->stage == MESA_SHADER_TESS_CTRL) {
- emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
+ if (ac->chip_class == SI && stage == MESA_SHADER_TESS_CTRL) {
+ ac_build_waitcnt(ac, LGKM_CNT & VM_CNT);
return;
}
- ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.s.barrier",
- ctx->voidt, NULL, 0, AC_FUNC_ATTR_CONVERGENT);
+ ac_build_intrinsic(ac, "llvm.amdgcn.s.barrier",
+ ac->voidt, NULL, 0, AC_FUNC_ATTR_CONVERGENT);
}
-static void emit_discard_if(struct ac_nir_context *ctx,
- const nir_intrinsic_instr *instr)
+static void emit_discard(struct ac_nir_context *ctx,
+ const nir_intrinsic_instr *instr)
{
LLVMValueRef cond;
- cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE,
- get_src(ctx, instr->src[0]),
- ctx->ac.i32_0, "");
+ if (instr->intrinsic == nir_intrinsic_discard_if) {
+ cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
+ get_src(ctx, instr->src[0]),
+ ctx->ac.i32_0, "");
+ } else {
+ assert(instr->intrinsic == nir_intrinsic_discard);
+ cond = LLVMConstInt(ctx->ac.i1, false, 0);
+ }
+
+ ac_build_kill_if_false(&ctx->ac, cond);
+}
- cond = LLVMBuildSelect(ctx->ac.builder, cond,
- LLVMConstReal(ctx->ac.f32, -1.0f),
- ctx->ac.f32_0, "");
- ac_build_kill(&ctx->ac, cond);
+static LLVMValueRef
+visit_load_helper_invocation(struct ac_nir_context *ctx)
+{
+ LLVMValueRef result = ac_build_intrinsic(&ctx->ac,
+ "llvm.amdgcn.ps.live",
+ ctx->ac.i1, NULL, 0,
+ AC_FUNC_ATTR_READNONE);
+ result = LLVMBuildNot(ctx->ac.builder, result, "");
+ return LLVMBuildSExt(ctx->ac.builder, result, ctx->ac.i32, "");
}
static LLVMValueRef
LLVMValueRef result;
LLVMValueRef thread_id = ac_get_thread_id(&ctx->ac);
result = LLVMBuildAnd(ctx->builder, ctx->tg_size,
- LLVMConstInt(ctx->i32, 0xfc0, false), "");
+ LLVMConstInt(ctx->ac.i32, 0xfc0, false), "");
return LLVMBuildAdd(ctx->builder, result, thread_id, "");
}
return result;
}
-#define INTERP_CENTER 0
-#define INTERP_CENTROID 1
-#define INTERP_SAMPLE 2
-
-static LLVMValueRef lookup_interp_param(struct nir_to_llvm_context *ctx,
+static LLVMValueRef lookup_interp_param(struct ac_shader_abi *abi,
enum glsl_interp_mode interp, unsigned location)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+
switch (interp) {
case INTERP_MODE_FLAT:
default:
return NULL;
}
-static LLVMValueRef load_sample_position(struct nir_to_llvm_context *ctx,
+static LLVMValueRef load_sample_position(struct ac_shader_abi *abi,
LLVMValueRef sample_id)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+
LLVMValueRef result;
- LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_PS_SAMPLE_POSITIONS, false));
+ LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false));
ptr = LLVMBuildBitCast(ctx->builder, ptr,
- const_array(ctx->v2f32, 64), "");
+ ac_array_in_const_addr_space(ctx->ac.v2f32), "");
sample_id = LLVMBuildAdd(ctx->builder, sample_id, ctx->sample_pos_offset, "");
result = ac_build_load_invariant(&ctx->ac, ptr, sample_id);
{
LLVMValueRef values[2];
- values[0] = emit_ffract(&ctx->ac, ctx->abi->frag_pos[0]);
- values[1] = emit_ffract(&ctx->ac, ctx->abi->frag_pos[1]);
+ values[0] = emit_ffract(&ctx->ac, ctx->abi->frag_pos[0], 32);
+ values[1] = emit_ffract(&ctx->ac, ctx->abi->frag_pos[1], 32);
return ac_build_gather_values(&ctx->ac, values, 2);
}
-static LLVMValueRef visit_interp(struct nir_to_llvm_context *ctx,
+static LLVMValueRef load_sample_mask_in(struct ac_nir_context *ctx)
+{
+ uint8_t log2_ps_iter_samples = ctx->nctx->shader_info->info.ps.force_persample ? ctx->nctx->options->key.fs.log2_num_samples : ctx->nctx->options->key.fs.log2_ps_iter_samples;
+
+ /* The bit pattern matches that used by fixed function fragment
+ * processing. */
+ static const uint16_t ps_iter_masks[] = {
+ 0xffff, /* not used */
+ 0x5555,
+ 0x1111,
+ 0x0101,
+ 0x0001,
+ };
+ assert(log2_ps_iter_samples < ARRAY_SIZE(ps_iter_masks));
+
+ uint32_t ps_iter_mask = ps_iter_masks[log2_ps_iter_samples];
+
+ LLVMValueRef result, sample_id;
+ sample_id = unpack_param(&ctx->ac, ctx->abi->ancillary, 8, 4);
+ sample_id = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->ac.i32, ps_iter_mask, false), sample_id, "");
+ result = LLVMBuildAnd(ctx->ac.builder, sample_id, ctx->abi->sample_coverage, "");
+ return result;
+}
+
+static LLVMValueRef visit_interp(struct ac_nir_context *ctx,
const nir_intrinsic_instr *instr)
{
- LLVMValueRef result[2];
+ LLVMValueRef result[4];
LLVMValueRef interp_param, attr_number;
unsigned location;
unsigned chan;
case nir_intrinsic_interp_var_at_sample:
case nir_intrinsic_interp_var_at_offset:
location = INTERP_CENTER;
- src0 = get_src(ctx->nir, instr->src[0]);
+ src0 = get_src(ctx, instr->src[0]);
break;
default:
break;
}
if (instr->intrinsic == nir_intrinsic_interp_var_at_offset) {
- 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, ""));
+ src_c0 = ac_to_float(&ctx->ac, LLVMBuildExtractElement(ctx->ac.builder, src0, ctx->ac.i32_0, ""));
+ src_c1 = ac_to_float(&ctx->ac, LLVMBuildExtractElement(ctx->ac.builder, src0, ctx->ac.i32_1, ""));
} else if (instr->intrinsic == nir_intrinsic_interp_var_at_sample) {
LLVMValueRef sample_position;
- LLVMValueRef halfval = LLVMConstReal(ctx->f32, 0.5f);
+ LLVMValueRef halfval = LLVMConstReal(ctx->ac.f32, 0.5f);
/* fetch sample ID */
- sample_position = load_sample_position(ctx, src0);
+ sample_position = ctx->abi->load_sample_position(ctx->abi, src0);
- src_c0 = LLVMBuildExtractElement(ctx->builder, sample_position, ctx->i32zero, "");
- src_c0 = LLVMBuildFSub(ctx->builder, src_c0, halfval, "");
- src_c1 = LLVMBuildExtractElement(ctx->builder, sample_position, ctx->i32one, "");
- src_c1 = LLVMBuildFSub(ctx->builder, src_c1, halfval, "");
+ src_c0 = LLVMBuildExtractElement(ctx->ac.builder, sample_position, ctx->ac.i32_0, "");
+ src_c0 = LLVMBuildFSub(ctx->ac.builder, src_c0, halfval, "");
+ src_c1 = LLVMBuildExtractElement(ctx->ac.builder, sample_position, ctx->ac.i32_1, "");
+ src_c1 = LLVMBuildFSub(ctx->ac.builder, src_c1, halfval, "");
}
- interp_param = lookup_interp_param(ctx, instr->variables[0]->var->data.interpolation, location);
- attr_number = LLVMConstInt(ctx->i32, input_index, false);
+ interp_param = ctx->abi->lookup_interp_param(ctx->abi, instr->variables[0]->var->data.interpolation, location);
+ attr_number = LLVMConstInt(ctx->ac.i32, input_index, false);
if (location == INTERP_CENTER) {
LLVMValueRef ij_out[2];
- LLVMValueRef ddxy_out = emit_ddxy_interp(ctx->nir, interp_param);
+ LLVMValueRef ddxy_out = emit_ddxy_interp(ctx, interp_param);
/*
* take the I then J parameters, and the DDX/Y for it, and
* interp_param.J = ddy * offset/sample.y + temp1;
*/
for (unsigned i = 0; i < 2; i++) {
- LLVMValueRef ix_ll = LLVMConstInt(ctx->i32, i, false);
- LLVMValueRef iy_ll = LLVMConstInt(ctx->i32, i + 2, false);
- LLVMValueRef ddx_el = LLVMBuildExtractElement(ctx->builder,
+ LLVMValueRef ix_ll = LLVMConstInt(ctx->ac.i32, i, false);
+ LLVMValueRef iy_ll = LLVMConstInt(ctx->ac.i32, i + 2, false);
+ LLVMValueRef ddx_el = LLVMBuildExtractElement(ctx->ac.builder,
ddxy_out, ix_ll, "");
- LLVMValueRef ddy_el = LLVMBuildExtractElement(ctx->builder,
+ LLVMValueRef ddy_el = LLVMBuildExtractElement(ctx->ac.builder,
ddxy_out, iy_ll, "");
- LLVMValueRef interp_el = LLVMBuildExtractElement(ctx->builder,
+ LLVMValueRef interp_el = LLVMBuildExtractElement(ctx->ac.builder,
interp_param, ix_ll, "");
LLVMValueRef temp1, temp2;
- interp_el = LLVMBuildBitCast(ctx->builder, interp_el,
- ctx->f32, "");
+ interp_el = LLVMBuildBitCast(ctx->ac.builder, interp_el,
+ ctx->ac.f32, "");
- temp1 = LLVMBuildFMul(ctx->builder, ddx_el, src_c0, "");
- temp1 = LLVMBuildFAdd(ctx->builder, temp1, interp_el, "");
+ temp1 = LLVMBuildFMul(ctx->ac.builder, ddx_el, src_c0, "");
+ temp1 = LLVMBuildFAdd(ctx->ac.builder, temp1, interp_el, "");
- temp2 = LLVMBuildFMul(ctx->builder, ddy_el, src_c1, "");
- temp2 = LLVMBuildFAdd(ctx->builder, temp2, temp1, "");
+ temp2 = LLVMBuildFMul(ctx->ac.builder, ddy_el, src_c1, "");
+ temp2 = LLVMBuildFAdd(ctx->ac.builder, temp2, temp1, "");
- ij_out[i] = LLVMBuildBitCast(ctx->builder,
- temp2, ctx->i32, "");
+ ij_out[i] = LLVMBuildBitCast(ctx->ac.builder,
+ temp2, ctx->ac.i32, "");
}
interp_param = ac_build_gather_values(&ctx->ac, ij_out, 2);
}
- for (chan = 0; chan < 2; chan++) {
- LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, false);
+ for (chan = 0; chan < 4; chan++) {
+ LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
if (interp_param) {
- interp_param = LLVMBuildBitCast(ctx->builder,
- interp_param, LLVMVectorType(ctx->f32, 2), "");
+ interp_param = LLVMBuildBitCast(ctx->ac.builder,
+ interp_param, ctx->ac.v2f32, "");
LLVMValueRef i = LLVMBuildExtractElement(
- ctx->builder, interp_param, ctx->i32zero, "");
+ ctx->ac.builder, interp_param, ctx->ac.i32_0, "");
LLVMValueRef j = LLVMBuildExtractElement(
- ctx->builder, interp_param, ctx->i32one, "");
+ ctx->ac.builder, interp_param, ctx->ac.i32_1, "");
result[chan] = ac_build_fs_interp(&ctx->ac,
llvm_chan, attr_number,
- ctx->prim_mask, i, j);
+ ctx->abi->prim_mask, i, j);
} else {
result[chan] = ac_build_fs_interp_mov(&ctx->ac,
- LLVMConstInt(ctx->i32, 2, false),
+ LLVMConstInt(ctx->ac.i32, 2, false),
llvm_chan, attr_number,
- ctx->prim_mask);
+ ctx->abi->prim_mask);
}
}
- return ac_build_gather_values(&ctx->ac, result, 2);
+ return ac_build_varying_gather_values(&ctx->ac, result, instr->num_components,
+ instr->variables[0]->var->data.location_frac);
}
static void
-visit_emit_vertex(struct nir_to_llvm_context *ctx,
- const nir_intrinsic_instr *instr)
+visit_emit_vertex(struct ac_shader_abi *abi, unsigned stream, LLVMValueRef *addrs)
{
LLVMValueRef gs_next_vertex;
- LLVMValueRef can_emit, kill;
+ LLVMValueRef can_emit;
int idx;
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+
+ assert(stream == 0);
- assert(instr->const_index[0] == 0);
/* Write vertex attribute values to GSVS ring */
gs_next_vertex = LLVMBuildLoad(ctx->builder,
ctx->gs_next_vertex,
* effects other than emitting vertices.
*/
can_emit = LLVMBuildICmp(ctx->builder, LLVMIntULT, gs_next_vertex,
- LLVMConstInt(ctx->i32, ctx->gs_max_out_vertices, false), "");
-
- kill = LLVMBuildSelect(ctx->builder, can_emit,
- LLVMConstReal(ctx->f32, 1.0f),
- LLVMConstReal(ctx->f32, -1.0f), "");
- ac_build_kill(&ctx->ac, kill);
+ LLVMConstInt(ctx->ac.i32, ctx->gs_max_out_vertices, false), "");
+ ac_build_kill_if_false(&ctx->ac, can_emit);
/* loop num outputs */
idx = 0;
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
- LLVMValueRef *out_ptr = &ctx->nir->outputs[i * 4];
+ LLVMValueRef *out_ptr = &addrs[i * 4];
int length = 4;
int slot = idx;
int slot_inc = 1;
for (unsigned j = 0; j < length; j++) {
LLVMValueRef out_val = LLVMBuildLoad(ctx->builder,
out_ptr[j], "");
- LLVMValueRef voffset = LLVMConstInt(ctx->i32, (slot * 4 + j) * ctx->gs_max_out_vertices, false);
+ LLVMValueRef voffset = LLVMConstInt(ctx->ac.i32, (slot * 4 + j) * ctx->gs_max_out_vertices, false);
voffset = LLVMBuildAdd(ctx->builder, voffset, gs_next_vertex, "");
- voffset = LLVMBuildMul(ctx->builder, voffset, LLVMConstInt(ctx->i32, 4, false), "");
+ voffset = LLVMBuildMul(ctx->builder, voffset, LLVMConstInt(ctx->ac.i32, 4, false), "");
- out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->i32, "");
+ out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->ac.i32, "");
ac_build_buffer_store_dword(&ctx->ac, ctx->gsvs_ring,
out_val, 1,
}
gs_next_vertex = LLVMBuildAdd(ctx->builder, gs_next_vertex,
- ctx->i32one, "");
+ ctx->ac.i32_1, "");
LLVMBuildStore(ctx->builder, gs_next_vertex, ctx->gs_next_vertex);
ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (0 << 8), ctx->gs_wave_id);
}
static void
-visit_end_primitive(struct nir_to_llvm_context *ctx,
- const nir_intrinsic_instr *instr)
+visit_end_primitive(struct ac_shader_abi *abi, unsigned stream)
{
- ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (0 << 8), ctx->gs_wave_id);
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+ ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8), ctx->gs_wave_id);
}
static LLVMValueRef
-visit_load_tess_coord(struct nir_to_llvm_context *ctx,
- const nir_intrinsic_instr *instr)
+load_tess_coord(struct ac_shader_abi *abi, LLVMTypeRef type,
+ unsigned num_components)
{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+
LLVMValueRef coord[4] = {
ctx->tes_u,
ctx->tes_v,
- ctx->f32zero,
- ctx->f32zero,
+ ctx->ac.f32_0,
+ ctx->ac.f32_0,
};
if (ctx->tes_primitive_mode == GL_TRIANGLES)
- coord[2] = LLVMBuildFSub(ctx->builder, ctx->f32one,
+ coord[2] = LLVMBuildFSub(ctx->builder, ctx->ac.f32_1,
LLVMBuildFAdd(ctx->builder, coord[0], coord[1], ""), "");
- LLVMValueRef result = ac_build_gather_values(&ctx->ac, coord, instr->num_components);
- return LLVMBuildBitCast(ctx->builder, result,
- get_def_type(ctx->nir, &instr->dest.ssa), "");
+ LLVMValueRef result = ac_build_gather_values(&ctx->ac, coord, num_components);
+ return LLVMBuildBitCast(ctx->builder, result, type, "");
+}
+
+static LLVMValueRef
+load_patch_vertices_in(struct ac_shader_abi *abi)
+{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+ return LLVMConstInt(ctx->ac.i32, ctx->options->key.tcs.input_vertices, false);
}
static void visit_intrinsic(struct ac_nir_context *ctx,
LLVMValueRef result = NULL;
switch (instr->intrinsic) {
+ case nir_intrinsic_ballot:
+ result = ac_build_ballot(&ctx->ac, get_src(ctx, instr->src[0]));
+ break;
+ case nir_intrinsic_read_invocation:
+ case nir_intrinsic_read_first_invocation: {
+ LLVMValueRef args[2];
+
+ /* Value */
+ args[0] = get_src(ctx, instr->src[0]);
+
+ unsigned num_args;
+ const char *intr_name;
+ if (instr->intrinsic == nir_intrinsic_read_invocation) {
+ num_args = 2;
+ intr_name = "llvm.amdgcn.readlane";
+
+ /* Invocation */
+ args[1] = get_src(ctx, instr->src[1]);
+ } else {
+ num_args = 1;
+ intr_name = "llvm.amdgcn.readfirstlane";
+ }
+
+ /* We currently have no other way to prevent LLVM from lifting the icmp
+ * calls to a dominating basic block.
+ */
+ ac_build_optimization_barrier(&ctx->ac, &args[0]);
+
+ result = ac_build_intrinsic(&ctx->ac, intr_name,
+ ctx->ac.i32, args, num_args,
+ AC_FUNC_ATTR_READNONE |
+ AC_FUNC_ATTR_CONVERGENT);
+ break;
+ }
+ case nir_intrinsic_load_subgroup_invocation:
+ result = ac_get_thread_id(&ctx->ac);
+ break;
case nir_intrinsic_load_work_group_id: {
- result = ctx->nctx->workgroup_ids;
+ LLVMValueRef values[3];
+
+ for (int i = 0; i < 3; i++) {
+ values[i] = ctx->abi->workgroup_ids[i] ?
+ ctx->abi->workgroup_ids[i] : ctx->ac.i32_0;
+ }
+
+ result = ac_build_gather_values(&ctx->ac, values, 3);
break;
}
case nir_intrinsic_load_base_vertex: {
result = ctx->abi->base_vertex;
break;
}
+ case nir_intrinsic_load_local_group_size:
+ result = ctx->abi->load_local_group_size(ctx->abi);
+ break;
case nir_intrinsic_load_vertex_id_zero_base: {
result = ctx->abi->vertex_id;
break;
}
case nir_intrinsic_load_local_invocation_id: {
- result = ctx->nctx->local_invocation_ids;
+ result = ctx->abi->local_invocation_ids;
break;
}
case nir_intrinsic_load_base_instance:
break;
case nir_intrinsic_load_invocation_id:
if (ctx->stage == MESA_SHADER_TESS_CTRL)
- result = unpack_param(&ctx->ac, ctx->nctx->tcs_rel_ids, 8, 5);
+ result = unpack_param(&ctx->ac, ctx->abi->tcs_rel_ids, 8, 5);
else
- result = ctx->nctx->gs_invocation_id;
+ result = ctx->abi->gs_invocation_id;
break;
case nir_intrinsic_load_primitive_id:
if (ctx->stage == MESA_SHADER_GEOMETRY) {
- ctx->nctx->shader_info->gs.uses_prim_id = true;
- result = ctx->nctx->gs_prim_id;
+ result = ctx->abi->gs_prim_id;
} else if (ctx->stage == MESA_SHADER_TESS_CTRL) {
- ctx->nctx->shader_info->tcs.uses_prim_id = true;
- result = ctx->nctx->tcs_patch_id;
+ result = ctx->abi->tcs_patch_id;
} else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
- ctx->nctx->shader_info->tcs.uses_prim_id = true;
- result = ctx->nctx->tes_patch_id;
+ result = ctx->abi->tes_patch_id;
} else
fprintf(stderr, "Unknown primitive id intrinsic: %d", ctx->stage);
break;
result = load_sample_pos(ctx);
break;
case nir_intrinsic_load_sample_mask_in:
- result = ctx->abi->sample_coverage;
+ if (ctx->nctx)
+ result = load_sample_mask_in(ctx);
+ else
+ result = ctx->abi->sample_coverage;
break;
case nir_intrinsic_load_frag_coord: {
LLVMValueRef values[4] = {
case nir_intrinsic_load_front_face:
result = ctx->abi->front_face;
break;
+ case nir_intrinsic_load_helper_invocation:
+ result = visit_load_helper_invocation(ctx);
+ break;
case nir_intrinsic_load_instance_id:
result = ctx->abi->instance_id;
break;
case nir_intrinsic_load_num_work_groups:
- result = ctx->nctx->num_work_groups;
+ result = ctx->abi->num_work_groups;
break;
case nir_intrinsic_load_local_invocation_index:
result = visit_load_local_invocation_index(ctx->nctx);
case nir_intrinsic_vulkan_resource_index:
result = visit_vulkan_resource_index(ctx->nctx, instr);
break;
+ case nir_intrinsic_vulkan_resource_reindex:
+ result = visit_vulkan_resource_reindex(ctx->nctx, instr);
+ break;
case nir_intrinsic_store_ssbo:
visit_store_ssbo(ctx, instr);
break;
case nir_intrinsic_image_size:
result = visit_image_size(ctx, instr);
break;
- case nir_intrinsic_discard:
- ac_build_intrinsic(&ctx->ac, "llvm.AMDGPU.kilp",
- LLVMVoidTypeInContext(ctx->ac.context),
- NULL, 0, AC_FUNC_ATTR_LEGACY);
+ case nir_intrinsic_shader_clock:
+ result = ac_build_shader_clock(&ctx->ac);
break;
+ case nir_intrinsic_discard:
case nir_intrinsic_discard_if:
- emit_discard_if(ctx, instr);
+ emit_discard(ctx, instr);
break;
case nir_intrinsic_memory_barrier:
- emit_waitcnt(ctx->nctx, VM_CNT);
+ case nir_intrinsic_group_memory_barrier:
+ case nir_intrinsic_memory_barrier_atomic_counter:
+ case nir_intrinsic_memory_barrier_buffer:
+ case nir_intrinsic_memory_barrier_image:
+ case nir_intrinsic_memory_barrier_shared:
+ emit_membar(ctx->nctx, instr);
break;
case nir_intrinsic_barrier:
- emit_barrier(ctx->nctx);
+ emit_barrier(&ctx->ac, ctx->stage);
break;
case nir_intrinsic_var_atomic_add:
case nir_intrinsic_var_atomic_imin:
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->nctx, instr);
+ result = visit_interp(ctx, instr);
break;
case nir_intrinsic_emit_vertex:
- visit_emit_vertex(ctx->nctx, instr);
+ ctx->abi->emit_vertex(ctx->abi, nir_intrinsic_stream_id(instr), ctx->outputs);
break;
case nir_intrinsic_end_primitive:
- visit_end_primitive(ctx->nctx, instr);
+ ctx->abi->emit_primitive(ctx->abi, nir_intrinsic_stream_id(instr));
+ break;
+ case nir_intrinsic_load_tess_coord: {
+ LLVMTypeRef type = ctx->nctx ?
+ get_def_type(ctx->nctx->nir, &instr->dest.ssa) :
+ NULL;
+ result = ctx->abi->load_tess_coord(ctx->abi, type, instr->num_components);
+ break;
+ }
+ case nir_intrinsic_load_tess_level_outer:
+ result = ctx->abi->load_tess_level(ctx->abi, VARYING_SLOT_TESS_LEVEL_OUTER);
break;
- case nir_intrinsic_load_tess_coord:
- result = visit_load_tess_coord(ctx->nctx, instr);
+ case nir_intrinsic_load_tess_level_inner:
+ result = ctx->abi->load_tess_level(ctx->abi, VARYING_SLOT_TESS_LEVEL_INNER);
break;
case nir_intrinsic_load_patch_vertices_in:
- result = LLVMConstInt(ctx->ac.i32, ctx->nctx->options->key.tcs.input_vertices, false);
+ result = ctx->abi->load_patch_vertices_in(ctx->abi);
+ break;
+ case nir_intrinsic_vote_all: {
+ LLVMValueRef tmp = ac_build_vote_all(&ctx->ac, get_src(ctx, instr->src[0]));
+ result = LLVMBuildSExt(ctx->ac.builder, tmp, ctx->ac.i32, "");
+ break;
+ }
+ case nir_intrinsic_vote_any: {
+ LLVMValueRef tmp = ac_build_vote_any(&ctx->ac, get_src(ctx, instr->src[0]));
+ result = LLVMBuildSExt(ctx->ac.builder, tmp, ctx->ac.i32, "");
break;
+ }
+ case nir_intrinsic_vote_eq: {
+ LLVMValueRef tmp = ac_build_vote_eq(&ctx->ac, get_src(ctx, instr->src[0]));
+ result = LLVMBuildSExt(ctx->ac.builder, tmp, ctx->ac.i32, "");
+ break;
+ }
default:
fprintf(stderr, "Unknown intrinsic: ");
nir_print_instr(&instr->instr, stderr);
}
static LLVMValueRef radv_load_ssbo(struct ac_shader_abi *abi,
- LLVMValueRef buffer, bool write)
+ LLVMValueRef buffer_ptr, bool write)
+{
+ struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+ LLVMValueRef result;
+
+ LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
+
+ result = LLVMBuildLoad(ctx->builder, buffer_ptr, "");
+ LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
+
+ return result;
+}
+
+static LLVMValueRef radv_load_ubo(struct ac_shader_abi *abi, LLVMValueRef buffer_ptr)
{
struct nir_to_llvm_context *ctx = nir_to_llvm_context_from_abi(abi);
+ LLVMValueRef result;
+
+ LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
- if (write && ctx->stage == MESA_SHADER_FRAGMENT)
- ctx->shader_info->fs.writes_memory = true;
+ result = LLVMBuildLoad(ctx->builder, buffer_ptr, "");
+ LLVMSetMetadata(result, ctx->ac.invariant_load_md_kind, ctx->ac.empty_md);
- return buffer;
+ return result;
}
static LLVMValueRef radv_get_sampler_desc(struct ac_shader_abi *abi,
assert(base_index < layout->binding_count);
- if (write && ctx->stage == MESA_SHADER_FRAGMENT)
- ctx->shader_info->fs.writes_memory = true;
-
switch (desc_type) {
case AC_DESC_IMAGE:
- type = ctx->v8i32;
+ type = ctx->ac.v8i32;
type_size = 32;
break;
case AC_DESC_FMASK:
- type = ctx->v8i32;
+ type = ctx->ac.v8i32;
offset += 32;
type_size = 32;
break;
case AC_DESC_SAMPLER:
- type = ctx->v4i32;
+ type = ctx->ac.v4i32;
if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER)
offset += 64;
type_size = 16;
break;
case AC_DESC_BUFFER:
- type = ctx->v4i32;
+ type = ctx->ac.v4i32;
type_size = 16;
break;
default:
const uint32_t *samplers = radv_immutable_samplers(layout, binding);
LLVMValueRef constants[] = {
- LLVMConstInt(ctx->i32, samplers[constant_index * 4 + 0], 0),
- LLVMConstInt(ctx->i32, samplers[constant_index * 4 + 1], 0),
- LLVMConstInt(ctx->i32, samplers[constant_index * 4 + 2], 0),
- LLVMConstInt(ctx->i32, samplers[constant_index * 4 + 3], 0),
+ LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 0], 0),
+ LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 1], 0),
+ LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 2], 0),
+ LLVMConstInt(ctx->ac.i32, samplers[constant_index * 4 + 3], 0),
};
return ac_build_gather_values(&ctx->ac, constants, 4);
}
assert(stride % type_size == 0);
if (!index)
- index = ctx->i32zero;
+ index = ctx->ac.i32_0;
- index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->i32, stride / type_size, 0), "");
+ index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
- list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->i32, offset, 0));
- list = LLVMBuildPointerCast(builder, list, const_array(type, 0), "");
+ list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->ac.i32, offset, 0));
+ list = LLVMBuildPointerCast(builder, list, ac_array_in_const_addr_space(type), "");
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,
+ const nir_tex_instr *tex_instr,
bool image, bool write)
{
LLVMValueRef index = NULL;
unsigned constant_index = 0;
- const nir_deref *tail = &deref->deref;
+ unsigned descriptor_set;
+ unsigned base_index;
- 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 (!deref) {
+ assert(tex_instr && !image);
+ descriptor_set = 0;
+ base_index = tex_instr->sampler_index;
+ } else {
+ 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;
+ if (!array_size)
+ array_size = 1;
- assert(child->deref_array_type != nir_deref_array_type_wildcard);
+ 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);
+ 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), "");
+ 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, "");
- }
+ if (!index)
+ index = indirect;
+ else
+ index = LLVMBuildAdd(ctx->ac.builder, index, indirect, "");
+ }
- constant_index += child->base_offset * array_size;
+ constant_index += child->base_offset * array_size;
- tail = &child->deref;
+ tail = &child->deref;
+ }
+ descriptor_set = deref->var->data.descriptor_set;
+ base_index = deref->var->data.binding;
}
return ctx->abi->load_sampler_desc(ctx->abi,
- deref->var->data.descriptor_set,
- deref->var->data.binding,
+ descriptor_set,
+ base_index,
constant_index, index,
desc_type, image, write);
}
LLVMValueRef *fmask_ptr)
{
if (instr->sampler_dim == GLSL_SAMPLER_DIM_BUF)
- *res_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_BUFFER, false, false);
+ *res_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_BUFFER, instr, false, false);
else
- *res_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_IMAGE, false, false);
+ *res_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_IMAGE, instr, false, false);
if (samp_ptr) {
if (instr->sampler)
- *samp_ptr = get_sampler_desc(ctx, instr->sampler, AC_DESC_SAMPLER, false, false);
+ *samp_ptr = get_sampler_desc(ctx, instr->sampler, AC_DESC_SAMPLER, instr, false, false);
else
- *samp_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_SAMPLER, false, false);
+ *samp_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_SAMPLER, instr, 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, AC_DESC_FMASK, false, false);
+ *fmask_ptr = get_sampler_desc(ctx, instr->texture, AC_DESC_FMASK, instr, false, false);
}
static LLVMValueRef apply_round_slice(struct ac_llvm_context *ctx,
if (coord)
for (chan = 0; chan < instr->coord_components; chan++)
- coords[chan] = llvm_extract_elem(&ctx->ac, coord, chan);
+ coords[chan] = ac_llvm_extract_elem(&ctx->ac, coord, chan);
if (offsets && instr->op != nir_texop_txf) {
LLVMValueRef offset[3], pack;
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->ac, offsets, chan);
+ for (chan = 0; chan < ac_get_llvm_num_components(offsets); chan++) {
+ offset[chan] = ac_llvm_extract_elem(&ctx->ac, offsets, chan);
offset[chan] = LLVMBuildAnd(ctx->ac.builder, offset[chan],
LLVMConstInt(ctx->ac.i32, 0x3f, false), "");
if (chan)
/* Pack depth comparison value */
if (instr->is_shadow && comparator) {
LLVMValueRef z = ac_to_float(&ctx->ac,
- llvm_extract_elem(&ctx->ac, comparator, 0));
+ ac_llvm_extract_elem(&ctx->ac, comparator, 0));
- /* TC-compatible HTILE promotes Z16 and Z24 to Z32_FLOAT,
+ /* TC-compatible HTILE on radeonsi 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)
+ if (ctx->ac.chip_class == VI && ctx->abi->clamp_shadow_reference)
z = ac_build_clamp(&ctx->ac, z);
address[count++] = z;
}
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));
+ derivs[i] = ac_to_float(&ctx->ac, ac_llvm_extract_elem(&ctx->ac, ddx, i));
+ derivs[num_dest_deriv_channels + i] = ac_to_float(&ctx->ac, 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;
filler = LLVMConstReal(ctx->ac.f32, 0.5);
if (instr->sampler_dim == GLSL_SAMPLER_DIM_1D) {
- if (instr->is_array) {
+ /* No nir_texop_lod, because it does not take a slice
+ * even with array textures. */
+ if (instr->is_array && instr->op != nir_texop_lod ) {
address[count] = address[count - 1];
address[count - 1] = filler;
count++;
const nir_ssa_undef_instr *instr)
{
unsigned num_components = instr->def.num_components;
+ LLVMTypeRef type = LLVMIntTypeInContext(ctx->ac.context, instr->def.bit_size);
LLVMValueRef undef;
if (num_components == 1)
- undef = LLVMGetUndef(ctx->ac.i32);
+ undef = LLVMGetUndef(type);
else {
- undef = LLVMGetUndef(LLVMVectorType(ctx->ac.i32, num_components));
+ undef = LLVMGetUndef(LLVMVectorType(type, num_components));
}
_mesa_hash_table_insert(ctx->defs, &instr->def, undef);
}
ctx->ac.context, fn, "");
LLVMValueRef cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE, value,
- LLVMConstInt(ctx->ac.i32, 0, false), "");
+ ctx->ac.i32_0, "");
LLVMBuildCondBr(ctx->ac.builder, cond, if_block, else_block);
LLVMPositionBuilderAtEnd(ctx->ac.builder, if_block);
int index = variable->data.location - VERT_ATTRIB_GENERIC0;
int idx = variable->data.location;
unsigned attrib_count = glsl_count_attribute_slots(variable->type, true);
+ uint8_t input_usage_mask =
+ ctx->shader_info->info.vs.input_usage_mask[variable->data.location];
+ unsigned num_channels = util_last_bit(input_usage_mask);
variable->data.driver_location = idx * 4;
- if (ctx->options->key.vs.instance_rate_inputs & (1u << index)) {
- buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.instance_id,
- ctx->abi.start_instance, "");
- ctx->shader_info->vs.vgpr_comp_cnt = MAX2(3,
- ctx->shader_info->vs.vgpr_comp_cnt);
- } else
- buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.vertex_id,
- ctx->abi.base_vertex, "");
-
for (unsigned i = 0; i < attrib_count; ++i, ++idx) {
- t_offset = LLVMConstInt(ctx->i32, index + i, false);
+ if (ctx->options->key.vs.instance_rate_inputs & (1u << (index + i))) {
+ buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.instance_id,
+ ctx->abi.start_instance, "");
+ if (ctx->options->key.vs.as_ls) {
+ ctx->shader_info->vs.vgpr_comp_cnt =
+ MAX2(2, ctx->shader_info->vs.vgpr_comp_cnt);
+ } else {
+ ctx->shader_info->vs.vgpr_comp_cnt =
+ MAX2(1, ctx->shader_info->vs.vgpr_comp_cnt);
+ }
+ } else
+ buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.vertex_id,
+ ctx->abi.base_vertex, "");
+ t_offset = LLVMConstInt(ctx->ac.i32, index + i, false);
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,
- LLVMConstInt(ctx->i32, 0, false),
- true);
+ ctx->ac.i32_0,
+ num_channels, false, true);
+
+ input = ac_build_expand_to_vec4(&ctx->ac, input, num_channels);
for (unsigned chan = 0; chan < 4; chan++) {
- LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, false);
+ LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
ctx->inputs[radeon_llvm_reg_index_soa(idx, chan)] =
ac_to_integer(&ctx->ac, LLVMBuildExtractElement(ctx->builder,
input, llvm_chan, ""));
LLVMValueRef i, j;
bool interp = interp_param != NULL;
- attr_number = LLVMConstInt(ctx->i32, attr, false);
+ attr_number = LLVMConstInt(ctx->ac.i32, attr, false);
/* fs.constant returns the param from the middle vertex, so it's not
* really useful for flat shading. It's meant to be used for custom
*/
if (interp) {
interp_param = LLVMBuildBitCast(ctx->builder, interp_param,
- LLVMVectorType(ctx->f32, 2), "");
+ ctx->ac.v2f32, "");
i = LLVMBuildExtractElement(ctx->builder, interp_param,
- ctx->i32zero, "");
+ ctx->ac.i32_0, "");
j = LLVMBuildExtractElement(ctx->builder, interp_param,
- ctx->i32one, "");
+ ctx->ac.i32_1, "");
}
for (chan = 0; chan < 4; chan++) {
- LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, false);
+ LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, chan, false);
if (interp) {
result[chan] = ac_build_fs_interp(&ctx->ac,
prim_mask, i, j);
} else {
result[chan] = ac_build_fs_interp_mov(&ctx->ac,
- LLVMConstInt(ctx->i32, 2, false),
+ LLVMConstInt(ctx->ac.i32, 2, false),
llvm_chan,
attr_number,
prim_mask);
else
interp_type = INTERP_CENTER;
- interp = lookup_interp_param(ctx, variable->data.interpolation, interp_type);
+ interp = lookup_interp_param(&ctx->abi, variable->data.interpolation, interp_type);
} else
interp = NULL;
}
if (uses_center && uses_centroid) {
- LLVMValueRef sel = LLVMBuildICmp(ctx->builder, LLVMIntSLT, ctx->prim_mask, ctx->ac.i32_0, "");
+ LLVMValueRef sel = LLVMBuildICmp(ctx->builder, LLVMIntSLT, ctx->abi.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, "");
}
if (i >= VARYING_SLOT_VAR0 || i == VARYING_SLOT_PNTC ||
i == VARYING_SLOT_PRIMITIVE_ID || i == VARYING_SLOT_LAYER) {
interp_param = *inputs;
- interp_fs_input(ctx, index, interp_param, ctx->prim_mask,
+ interp_fs_input(ctx, index, interp_param, ctx->abi.prim_mask,
inputs);
if (!interp_param)
for(int i = 0; i < 3; ++i)
inputs[i] = ctx->abi.frag_pos[i];
- inputs[3] = ac_build_fdiv(&ctx->ac, ctx->f32one,
+ inputs[3] = ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
ctx->abi.frag_pos[3]);
}
}
case GLSL_TYPE_UINT:
case GLSL_TYPE_BOOL:
case GLSL_TYPE_SUBROUTINE:
- return ctx->i32;
+ return ctx->ac.i32;
case GLSL_TYPE_FLOAT: /* TODO handle mediump */
- return ctx->f32;
+ return ctx->ac.f32;
case GLSL_TYPE_INT64:
case GLSL_TYPE_UINT64:
- return ctx->i64;
+ return ctx->ac.i64;
case GLSL_TYPE_DOUBLE:
- return ctx->f64;
+ return ctx->ac.f64;
default:
unreachable("unknown GLSL type");
}
nir_foreach_variable(variable, &func->impl->locals) {
unsigned attrib_count = glsl_count_attribute_slots(variable->type, false);
variable->data.driver_location = ctx->num_locals * 4;
+ variable->data.location_frac = 0;
ctx->num_locals += attrib_count;
}
ctx->locals = malloc(4 * ctx->num_locals * sizeof(LLVMValueRef));
LLVMAddGlobalInAddressSpace(
ctx->ac.module, glsl_to_llvm_type(ctx->nctx, variable->type),
variable->name ? variable->name : "",
- LOCAL_ADDR_SPACE);
+ AC_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 = 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));
-}
-
-
-static LLVMValueRef emit_pack_int16(struct nir_to_llvm_context *ctx,
- LLVMValueRef src0, LLVMValueRef src1)
-{
- LLVMValueRef const16 = LLVMConstInt(ctx->i32, 16, false);
- LLVMValueRef comp[2];
-
- comp[0] = LLVMBuildAnd(ctx->builder, src0, LLVMConstInt(ctx-> i32, 65535, 0), "");
- comp[1] = LLVMBuildAnd(ctx->builder, src1, LLVMConstInt(ctx-> i32, 65535, 0), "");
- comp[1] = LLVMBuildShl(ctx->builder, comp[1], const16, "");
- return LLVMBuildOr(ctx->builder, comp[0], comp[1], "");
-}
-
/* Initialize arguments for the shader export intrinsic */
static void
si_llvm_init_export_args(struct nir_to_llvm_context *ctx,
args->target = target;
args->compr = false;
- args->out[0] = LLVMGetUndef(ctx->f32);
- args->out[1] = LLVMGetUndef(ctx->f32);
- args->out[2] = LLVMGetUndef(ctx->f32);
- args->out[3] = LLVMGetUndef(ctx->f32);
-
- if (!values)
- return;
+ args->out[0] = LLVMGetUndef(ctx->ac.f32);
+ args->out[1] = LLVMGetUndef(ctx->ac.f32);
+ args->out[2] = LLVMGetUndef(ctx->ac.f32);
+ args->out[3] = LLVMGetUndef(ctx->ac.f32);
if (ctx->stage == MESA_SHADER_FRAGMENT && target >= V_008DFC_SQ_EXP_MRT) {
- LLVMValueRef val[4];
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;
+ unsigned chan;
+
+ LLVMValueRef (*packf)(struct ac_llvm_context *ctx, LLVMValueRef args[2]) = NULL;
+ LLVMValueRef (*packi)(struct ac_llvm_context *ctx, LLVMValueRef args[2],
+ unsigned bits, bool hi) = NULL;
switch(col_format) {
case V_028714_SPI_SHADER_ZERO:
break;
case V_028714_SPI_SHADER_FP16_ABGR:
- args->compr = 1;
-
- for (unsigned chan = 0; chan < 2; chan++) {
- LLVMValueRef pack_args[2] = {
- values[2 * chan],
- values[2 * chan + 1]
- };
- LLVMValueRef packed;
-
- packed = ac_build_cvt_pkrtz_f16(&ctx->ac, pack_args);
- args->out[chan] = packed;
- }
+ packf = ac_build_cvt_pkrtz_f16;
break;
case V_028714_SPI_SHADER_UNORM16_ABGR:
- for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = ac_build_clamp(&ctx->ac, values[chan]);
- val[chan] = LLVMBuildFMul(ctx->builder, val[chan],
- LLVMConstReal(ctx->f32, 65535), "");
- val[chan] = LLVMBuildFAdd(ctx->builder, val[chan],
- LLVMConstReal(ctx->f32, 0.5), "");
- val[chan] = LLVMBuildFPToUI(ctx->builder, val[chan],
- ctx->i32, "");
- }
-
- args->compr = 1;
- args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
- args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
+ packf = ac_build_cvt_pknorm_u16;
break;
case V_028714_SPI_SHADER_SNORM16_ABGR:
- for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = emit_float_saturate(&ctx->ac, values[chan], -1, 1);
- val[chan] = LLVMBuildFMul(ctx->builder, val[chan],
- LLVMConstReal(ctx->f32, 32767), "");
-
- /* If positive, add 0.5, else add -0.5. */
- val[chan] = LLVMBuildFAdd(ctx->builder, val[chan],
- LLVMBuildSelect(ctx->builder,
- LLVMBuildFCmp(ctx->builder, LLVMRealOGE,
- val[chan], ctx->f32zero, ""),
- LLVMConstReal(ctx->f32, 0.5),
- LLVMConstReal(ctx->f32, -0.5), ""), "");
- val[chan] = LLVMBuildFPToSI(ctx->builder, val[chan], ctx->i32, "");
- }
-
- args->compr = 1;
- args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
- args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
+ packf = ac_build_cvt_pknorm_i16;
break;
- case V_028714_SPI_SHADER_UINT16_ABGR: {
- 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);
+ case V_028714_SPI_SHADER_UINT16_ABGR:
+ packi = ac_build_cvt_pk_u16;
+ break;
- for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = ac_to_integer(&ctx->ac, values[chan]);
- val[chan] = emit_minmax_int(&ctx->ac, LLVMIntULT, val[chan], chan == 3 ? max_alpha : max_rgb);
- }
+ case V_028714_SPI_SHADER_SINT16_ABGR:
+ packi = ac_build_cvt_pk_i16;
+ break;
- args->compr = 1;
- args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
- args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
+ default:
+ case V_028714_SPI_SHADER_32_ABGR:
+ memcpy(&args->out[0], values, sizeof(values[0]) * 4);
break;
}
- case V_028714_SPI_SHADER_SINT16_ABGR: {
- 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);
+ /* Pack f16 or norm_i16/u16. */
+ if (packf) {
+ for (chan = 0; chan < 2; chan++) {
+ LLVMValueRef pack_args[2] = {
+ values[2 * chan],
+ values[2 * chan + 1]
+ };
+ LLVMValueRef packed;
- /* Clamp. */
- for (unsigned chan = 0; chan < 4; chan++) {
- val[chan] = ac_to_integer(&ctx->ac, values[chan]);
- val[chan] = emit_minmax_int(&ctx->ac, LLVMIntSLT, val[chan], chan == 3 ? max_alpha : max_rgb);
- val[chan] = emit_minmax_int(&ctx->ac, LLVMIntSGT, val[chan], chan == 3 ? min_alpha : min_rgb);
+ packed = packf(&ctx->ac, pack_args);
+ args->out[chan] = ac_to_float(&ctx->ac, packed);
}
-
- args->compr = 1;
- args->out[0] = emit_pack_int16(ctx, val[0], val[1]);
- args->out[1] = emit_pack_int16(ctx, val[2], val[3]);
- break;
+ args->compr = 1; /* COMPR flag */
}
- default:
- case V_028714_SPI_SHADER_32_ABGR:
- memcpy(&args->out[0], values, sizeof(values[0]) * 4);
- break;
+ /* Pack i16/u16. */
+ if (packi) {
+ for (chan = 0; chan < 2; chan++) {
+ LLVMValueRef pack_args[2] = {
+ ac_to_integer(&ctx->ac, values[2 * chan]),
+ ac_to_integer(&ctx->ac, values[2 * chan + 1])
+ };
+ LLVMValueRef packed;
+
+ packed = packi(&ctx->ac, pack_args,
+ is_int8 ? 8 : is_int10 ? 10 : 16,
+ chan == 1);
+ args->out[chan] = ac_to_float(&ctx->ac, packed);
+ }
+ args->compr = 1; /* COMPR flag */
}
- } else
- memcpy(&args->out[0], values, sizeof(values[0]) * 4);
+ return;
+ }
+
+ memcpy(&args->out[0], values, sizeof(values[0]) * 4);
for (unsigned i = 0; i < 4; ++i)
args->out[i] = ac_to_float(&ctx->ac, args->out[i]);
}
+static void
+radv_export_param(struct nir_to_llvm_context *ctx, unsigned index,
+ LLVMValueRef *values)
+{
+ struct ac_export_args args;
+
+ si_llvm_init_export_args(ctx, values,
+ V_008DFC_SQ_EXP_PARAM + index, &args);
+ ac_build_export(&ctx->ac, &args);
+}
+
+static LLVMValueRef
+radv_load_output(struct nir_to_llvm_context *ctx, unsigned index, unsigned chan)
+{
+ LLVMValueRef output =
+ ctx->nir->outputs[radeon_llvm_reg_index_soa(index, chan)];
+
+ return LLVMBuildLoad(ctx->builder, output, "");
+}
+
static void
handle_vs_outputs_post(struct nir_to_llvm_context *ctx,
bool export_prim_id,
i = VARYING_SLOT_CLIP_DIST0;
for (j = 0; j < ctx->num_output_clips + ctx->num_output_culls; j++)
- slots[j] = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
- ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
+ slots[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
for (i = ctx->num_output_clips + ctx->num_output_culls; i < 8; i++)
- slots[i] = LLVMGetUndef(ctx->f32);
+ slots[i] = LLVMGetUndef(ctx->ac.f32);
if (ctx->num_output_clips + ctx->num_output_culls > 4) {
target = V_008DFC_SQ_EXP_POS + 3;
}
- LLVMValueRef pos_values[4] = {ctx->f32zero, ctx->f32zero, ctx->f32zero, ctx->f32one};
+ LLVMValueRef pos_values[4] = {ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_0, ctx->ac.f32_1};
if (ctx->output_mask & (1ull << VARYING_SLOT_POS)) {
for (unsigned j = 0; j < 4; j++)
- pos_values[j] = LLVMBuildLoad(ctx->builder,
- ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_POS, j)], "");
+ pos_values[j] = radv_load_output(ctx, 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->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_PSIZ, 0)], "");
+ psize_value = radv_load_output(ctx, VARYING_SLOT_PSIZ, 0);
}
if (ctx->output_mask & (1ull << VARYING_SLOT_LAYER)) {
outinfo->writes_layer = true;
- layer_value = LLVMBuildLoad(ctx->builder,
- ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)], "");
+ layer_value = radv_load_output(ctx, VARYING_SLOT_LAYER, 0);
}
if (ctx->output_mask & (1ull << VARYING_SLOT_VIEWPORT)) {
outinfo->writes_viewport_index = true;
- viewport_index_value = LLVMBuildLoad(ctx->builder,
- ctx->nir->outputs[radeon_llvm_reg_index_soa(VARYING_SLOT_VIEWPORT, 0)], "");
+ viewport_index_value = radv_load_output(ctx, VARYING_SLOT_VIEWPORT, 0);
}
if (outinfo->writes_pointsize ||
pos_args[1].done = 0;
pos_args[1].target = V_008DFC_SQ_EXP_POS + 1;
pos_args[1].compr = 0;
- pos_args[1].out[0] = ctx->f32zero; /* X */
- pos_args[1].out[1] = ctx->f32zero; /* Y */
- pos_args[1].out[2] = ctx->f32zero; /* Z */
- pos_args[1].out[3] = ctx->f32zero; /* W */
+ pos_args[1].out[0] = ctx->ac.f32_0; /* X */
+ pos_args[1].out[1] = ctx->ac.f32_0; /* Y */
+ pos_args[1].out[2] = ctx->ac.f32_0; /* Z */
+ pos_args[1].out[3] = ctx->ac.f32_0; /* W */
if (outinfo->writes_pointsize == true)
pos_args[1].out[0] = psize_value;
LLVMValueRef v = viewport_index_value;
v = ac_to_integer(&ctx->ac, v);
v = LLVMBuildShl(ctx->builder, v,
- LLVMConstInt(ctx->i32, 16, false),
+ LLVMConstInt(ctx->ac.i32, 16, false),
"");
v = LLVMBuildOr(ctx->builder, v,
ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
if (!(ctx->output_mask & (1ull << i)))
continue;
- for (unsigned j = 0; j < 4; 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;
- outinfo->vs_output_param_offset[VARYING_SLOT_LAYER] = param_count;
- param_count++;
- } else if (i == VARYING_SLOT_PRIMITIVE_ID) {
- target = V_008DFC_SQ_EXP_PARAM + param_count;
- outinfo->vs_output_param_offset[VARYING_SLOT_PRIMITIVE_ID] = param_count;
- param_count++;
- } else if (i >= VARYING_SLOT_VAR0) {
- outinfo->export_mask |= 1u << (i - VARYING_SLOT_VAR0);
- target = V_008DFC_SQ_EXP_PARAM + param_count;
- outinfo->vs_output_param_offset[i] = param_count;
- param_count++;
- } else
+ if (i != VARYING_SLOT_LAYER &&
+ i != VARYING_SLOT_PRIMITIVE_ID &&
+ i < VARYING_SLOT_VAR0)
continue;
- si_llvm_init_export_args(ctx, values, target, &args);
+ for (unsigned j = 0; j < 4; j++)
+ values[j] = ac_to_float(&ctx->ac, radv_load_output(ctx, i, j));
- if (target >= V_008DFC_SQ_EXP_POS &&
- target <= (V_008DFC_SQ_EXP_POS + 3)) {
- memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
- &args, sizeof(args));
- } else {
- ac_build_export(&ctx->ac, &args);
- }
+ radv_export_param(ctx, param_count, values);
+
+ outinfo->vs_output_param_offset[i] = param_count++;
}
if (export_prim_id) {
LLVMValueRef values[4];
- target = V_008DFC_SQ_EXP_PARAM + param_count;
- outinfo->vs_output_param_offset[VARYING_SLOT_PRIMITIVE_ID] = param_count;
- param_count++;
values[0] = ctx->vs_prim_id;
ctx->shader_info->vs.vgpr_comp_cnt = MAX2(2,
ctx->shader_info->vs.vgpr_comp_cnt);
for (unsigned j = 1; j < 4; j++)
- values[j] = ctx->f32zero;
- si_llvm_init_export_args(ctx, values, target, &args);
- ac_build_export(&ctx->ac, &args);
+ values[j] = ctx->ac.f32_0;
+
+ radv_export_param(ctx, param_count, values);
+
+ outinfo->vs_output_param_offset[VARYING_SLOT_PRIMITIVE_ID] = param_count++;
outinfo->export_prim_id = true;
}
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), ""), "");
+ LLVMConstInt(ctx->ac.i32, 64, false), ""), "");
lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
- LLVMConstInt(ctx->i32, itemsize_dw, 0), "");
+ LLVMConstInt(ctx->ac.i32, itemsize_dw, 0), "");
}
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
if (lds_base) {
dw_addr = LLVMBuildAdd(ctx->builder, lds_base,
- LLVMConstInt(ctx->i32, param_index * 4, false),
+ LLVMConstInt(ctx->ac.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, "");
+ out_val = LLVMBuildBitCast(ctx->builder, out_val, ctx->ac.i32, "");
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, "");
+ ac_lds_store(&ctx->ac, dw_addr,
+ LLVMBuildLoad(ctx->builder, out_ptr[j], ""));
+ dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->ac.i32_1, "");
} else {
ac_build_buffer_store_dword(&ctx->ac,
ctx->esgs_ring,
if (length > 4)
mark_tess_output(ctx, false, param + 1);
LLVMValueRef dw_addr = LLVMBuildAdd(ctx->builder, base_dw_addr,
- LLVMConstInt(ctx->i32, param * 4, false),
+ LLVMConstInt(ctx->ac.i32, param * 4, false),
"");
for (unsigned j = 0; j < length; j++) {
- lds_store(ctx, dw_addr,
- LLVMBuildLoad(ctx->builder, out_ptr[j], ""));
- dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->i32one, "");
+ ac_lds_store(&ctx->ac, dw_addr,
+ LLVMBuildLoad(ctx->builder, out_ptr[j], ""));
+ dw_addr = LLVMBuildAdd(ctx->builder, dw_addr, ctx->ac.i32_1, "");
}
}
}
{
unsigned stride, outer_comps, inner_comps;
struct ac_build_if_state if_ctx, inner_if_ctx;
- 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);
+ LLVMValueRef invocation_id = unpack_param(&ctx->ac, ctx->abi.tcs_rel_ids, 8, 5);
+ LLVMValueRef rel_patch_id = unpack_param(&ctx->ac, ctx->abi.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];
int i;
- emit_barrier(ctx);
+ emit_barrier(&ctx->ac, ctx->stage);
switch (ctx->options->key.tcs.primitive_mode) {
case GL_ISOLINES:
ac_nir_build_if(&if_ctx, ctx,
LLVMBuildICmp(ctx->builder, LLVMIntEQ,
- invocation_id, ctx->i32zero, ""));
+ invocation_id, ctx->ac.i32_0, ""));
tess_inner_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
tess_outer_index = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
mark_tess_output(ctx, true, tess_outer_index);
lds_base = get_tcs_out_current_patch_data_offset(ctx);
lds_inner = LLVMBuildAdd(ctx->builder, lds_base,
- LLVMConstInt(ctx->i32, tess_inner_index * 4, false), "");
+ LLVMConstInt(ctx->ac.i32, tess_inner_index * 4, false), "");
lds_outer = LLVMBuildAdd(ctx->builder, lds_base,
- LLVMConstInt(ctx->i32, tess_outer_index * 4, false), "");
+ LLVMConstInt(ctx->ac.i32, tess_outer_index * 4, false), "");
for (i = 0; i < 4; i++) {
- inner[i] = LLVMGetUndef(ctx->i32);
- outer[i] = LLVMGetUndef(ctx->i32);
+ inner[i] = LLVMGetUndef(ctx->ac.i32);
+ outer[i] = LLVMGetUndef(ctx->ac.i32);
}
// LINES reverseal
if (ctx->options->key.tcs.primitive_mode == GL_ISOLINES) {
- outer[0] = out[1] = lds_load(ctx, lds_outer);
+ outer[0] = out[1] = ac_lds_load(&ctx->ac, lds_outer);
lds_outer = LLVMBuildAdd(ctx->builder, lds_outer,
- LLVMConstInt(ctx->i32, 1, false), "");
- outer[1] = out[0] = lds_load(ctx, lds_outer);
+ ctx->ac.i32_1, "");
+ outer[1] = out[0] = ac_lds_load(&ctx->ac, lds_outer);
} else {
for (i = 0; i < outer_comps; i++) {
outer[i] = out[i] =
- lds_load(ctx, lds_outer);
+ ac_lds_load(&ctx->ac, lds_outer);
lds_outer = LLVMBuildAdd(ctx->builder, lds_outer,
- LLVMConstInt(ctx->i32, 1, false), "");
+ ctx->ac.i32_1, "");
}
for (i = 0; i < inner_comps; i++) {
inner[i] = out[outer_comps+i] =
- lds_load(ctx, lds_inner);
+ ac_lds_load(&ctx->ac, lds_inner);
lds_inner = LLVMBuildAdd(ctx->builder, lds_inner,
- LLVMConstInt(ctx->i32, 1, false), "");
+ ctx->ac.i32_1, "");
}
}
buffer = ctx->hs_ring_tess_factor;
tf_base = ctx->tess_factor_offset;
byteoffset = LLVMBuildMul(ctx->builder, rel_patch_id,
- LLVMConstInt(ctx->i32, 4 * stride, false), "");
+ LLVMConstInt(ctx->ac.i32, 4 * stride, false), "");
unsigned tf_offset = 0;
if (ctx->options->chip_class <= VI) {
ac_nir_build_if(&inner_if_ctx, ctx,
LLVMBuildICmp(ctx->builder, LLVMIntEQ,
- rel_patch_id, ctx->i32zero, ""));
+ rel_patch_id, ctx->ac.i32_0, ""));
/* Store the dynamic HS control word. */
ac_build_buffer_store_dword(&ctx->ac, buffer,
- LLVMConstInt(ctx->i32, 0x80000000, false),
- 1, ctx->i32zero, tf_base,
+ LLVMConstInt(ctx->ac.i32, 0x80000000, false),
+ 1, ctx->ac.i32_0, tf_base,
0, 1, 0, true, false);
tf_offset += 4;
stride - 4, byteoffset, tf_base,
16 + tf_offset, 1, 0, true, false);
- //TODO store to offchip for TES to read - only if TES reads them
- if (1) {
+ //store to offchip for TES to read - only if TES reads them
+ if (ctx->options->key.tcs.tes_reads_tess_factors) {
LLVMValueRef inner_vec, outer_vec, tf_outer_offset;
LLVMValueRef tf_inner_offset;
unsigned param_outer, param_inner;
param_outer = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_OUTER);
tf_outer_offset = get_tcs_tes_buffer_address(ctx, NULL,
- LLVMConstInt(ctx->i32, param_outer, 0));
+ LLVMConstInt(ctx->ac.i32, param_outer, 0));
outer_vec = ac_build_gather_values(&ctx->ac, outer,
util_next_power_of_two(outer_comps));
if (inner_comps) {
param_inner = shader_io_get_unique_index(VARYING_SLOT_TESS_LEVEL_INNER);
tf_inner_offset = get_tcs_tes_buffer_address(ctx, NULL,
- LLVMConstInt(ctx->i32, param_inner, 0));
+ LLVMConstInt(ctx->ac.i32, param_inner, 0));
inner_vec = inner_comps == 1 ? inner[0] :
ac_build_gather_values(&ctx->ac, inner, inner_comps);
static bool
si_export_mrt_color(struct nir_to_llvm_context *ctx,
- LLVMValueRef *color, unsigned param, bool is_last,
+ LLVMValueRef *color, unsigned index, bool is_last,
struct ac_export_args *args)
{
/* Export */
- si_llvm_init_export_args(ctx, color, param,
- args);
+ si_llvm_init_export_args(ctx, color,
+ V_008DFC_SQ_EXP_MRT + index, args);
if (is_last) {
args->valid_mask = 1; /* whether the EXEC mask is valid */
}
static void
-si_export_mrt_z(struct nir_to_llvm_context *ctx,
- LLVMValueRef depth, LLVMValueRef stencil,
- LLVMValueRef samplemask)
+radv_export_mrt_z(struct nir_to_llvm_context *ctx,
+ LLVMValueRef depth, LLVMValueRef stencil,
+ LLVMValueRef samplemask)
{
struct ac_export_args args;
- args.enabled_channels = 0;
- args.valid_mask = 1;
- args.done = 1;
- args.target = V_008DFC_SQ_EXP_MRTZ;
- args.compr = false;
-
- args.out[0] = LLVMGetUndef(ctx->f32); /* R, depth */
- args.out[1] = LLVMGetUndef(ctx->f32); /* G, stencil test val[0:7], stencil op val[8:15] */
- args.out[2] = LLVMGetUndef(ctx->f32); /* B, sample mask */
- args.out[3] = LLVMGetUndef(ctx->f32); /* A, alpha to mask */
-
- if (depth) {
- args.out[0] = depth;
- args.enabled_channels |= 0x1;
- }
-
- if (stencil) {
- args.out[1] = stencil;
- args.enabled_channels |= 0x2;
- }
-
- if (samplemask) {
- args.out[2] = samplemask;
- args.enabled_channels |= 0x4;
- }
-
- /* SI (except OLAND and HAINAN) has a bug that it only looks
- * at the X writemask component. */
- if (ctx->options->chip_class == SI &&
- ctx->options->family != CHIP_OLAND &&
- ctx->options->family != CHIP_HAINAN)
- args.enabled_channels |= 0x1;
+ ac_export_mrt_z(&ctx->ac, depth, stencil, samplemask, &args);
ac_build_export(&ctx->ac, &args);
}
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
LLVMValueRef values[4];
+ bool last = false;
if (!(ctx->output_mask & (1ull << i)))
continue;
- if (i == FRAG_RESULT_DEPTH) {
- ctx->shader_info->fs.writes_z = true;
- 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 = 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 = 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] = ac_to_float(&ctx->ac, LLVMBuildLoad(ctx->builder,
- ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)], ""));
+ if (i < FRAG_RESULT_DATA0)
+ continue;
- 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);
+ for (unsigned j = 0; j < 4; j++)
+ values[j] = ac_to_float(&ctx->ac,
+ radv_load_output(ctx, i, j));
- bool ret = si_export_mrt_color(ctx, values, V_008DFC_SQ_EXP_MRT + (i - FRAG_RESULT_DATA0), last, &color_args[index]);
- if (ret)
- index++;
- }
+ if (!ctx->shader_info->info.ps.writes_z &&
+ !ctx->shader_info->info.ps.writes_stencil &&
+ !ctx->shader_info->info.ps.writes_sample_mask)
+ last = ctx->output_mask <= ((1ull << (i + 1)) - 1);
+
+ bool ret = si_export_mrt_color(ctx, values,
+ i - FRAG_RESULT_DATA0,
+ last, &color_args[index]);
+ if (ret)
+ index++;
+ }
+
+ /* Process depth, stencil, samplemask. */
+ if (ctx->shader_info->info.ps.writes_z) {
+ depth = ac_to_float(&ctx->ac,
+ radv_load_output(ctx, FRAG_RESULT_DEPTH, 0));
+ }
+ if (ctx->shader_info->info.ps.writes_stencil) {
+ stencil = ac_to_float(&ctx->ac,
+ radv_load_output(ctx, FRAG_RESULT_STENCIL, 0));
+ }
+ if (ctx->shader_info->info.ps.writes_sample_mask) {
+ samplemask = ac_to_float(&ctx->ac,
+ radv_load_output(ctx, FRAG_RESULT_SAMPLE_MASK, 0));
}
+ /* Export PS outputs. */
for (unsigned i = 0; i < index; i++)
ac_build_export(&ctx->ac, &color_args[i]);
- if (depth || stencil || samplemask)
- si_export_mrt_z(ctx, depth, stencil, samplemask);
- else if (!index) {
- si_export_mrt_color(ctx, NULL, V_008DFC_SQ_EXP_NULL, true, &color_args[0]);
- ac_build_export(&ctx->ac, &color_args[0]);
- }
- ctx->shader_info->fs.output_mask = index ? ((1ull << index) - 1) : 0;
+ if (depth || stencil || samplemask)
+ radv_export_mrt_z(ctx, depth, stencil, samplemask);
+ else if (!index)
+ ac_build_export_null(&ctx->ac);
}
static void
{
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_load_to_sgpr(&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->ac.i32, RING_ESGS_VS, false));
}
if (ctx->is_gs_copy_shader) {
- ctx->gsvs_ring = ac_build_load_to_sgpr(&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->ac.i32, RING_GSVS_VS, false));
}
if (ctx->stage == MESA_SHADER_GEOMETRY) {
LLVMValueRef tmp;
- 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->esgs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_ESGS_GS, false));
+ ctx->gsvs_ring = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_GS, false));
- ctx->gsvs_ring = LLVMBuildBitCast(ctx->builder, ctx->gsvs_ring, ctx->v4i32, "");
+ ctx->gsvs_ring = LLVMBuildBitCast(ctx->builder, ctx->gsvs_ring, ctx->ac.v4i32, "");
- ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->i32, 2, false), "");
- tmp = LLVMBuildExtractElement(ctx->builder, ctx->gsvs_ring, ctx->i32one, "");
+ ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, ctx->gsvs_num_entries, LLVMConstInt(ctx->ac.i32, 2, false), "");
+ tmp = LLVMBuildExtractElement(ctx->builder, ctx->gsvs_ring, ctx->ac.i32_1, "");
tmp = LLVMBuildOr(ctx->builder, tmp, ctx->gsvs_ring_stride, "");
- ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, tmp, ctx->i32one, "");
+ ctx->gsvs_ring = LLVMBuildInsertElement(ctx->builder, ctx->gsvs_ring, tmp, ctx->ac.i32_1, "");
}
if (ctx->stage == MESA_SHADER_TESS_CTRL ||
ctx->stage == MESA_SHADER_TESS_EVAL) {
- ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_OFFCHIP, false));
- ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->i32, RING_HS_TESS_FACTOR, false));
+ ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_OFFCHIP, false));
+ ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_FACTOR, false));
}
}
case MESA_SHADER_TESS_CTRL:
return chip_class >= CIK ? 128 : 64;
case MESA_SHADER_GEOMETRY:
- return 64;
+ return chip_class >= GFX9 ? 128 : 64;
case MESA_SHADER_COMPUTE:
break;
default:
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->ac.i32_0, "");
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, "");
+ ctx->rel_auto_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_rel_ids, ctx->rel_auto_id, "");
+ ctx->abi.vertex_id = LLVMBuildSelect(ctx->ac.builder, hs_empty, ctx->abi.tcs_patch_id, ctx->abi.vertex_id, "");
}
static void prepare_gs_input_vgprs(struct nir_to_llvm_context *ctx)
visit_cf_list(&ctx, &func->impl->body);
phi_post_pass(&ctx);
- ctx.abi->emit_outputs(ctx.abi, RADEON_LLVM_MAX_OUTPUTS,
- ctx.outputs);
+ if (nir->info.stage != MESA_SHADER_COMPUTE)
+ ctx.abi->emit_outputs(ctx.abi, RADEON_LLVM_MAX_OUTPUTS,
+ ctx.outputs);
free(ctx.locals);
ralloc_free(ctx.defs);
ctx.context = LLVMContextCreate();
ctx.module = LLVMModuleCreateWithNameInContext("shader", ctx.context);
- ac_llvm_context_init(&ctx.ac, ctx.context, options->chip_class);
+ ac_llvm_context_init(&ctx.ac, ctx.context, options->chip_class,
+ options->family);
ctx.ac.module = ctx.module;
LLVMSetTarget(ctx.module, options->supports_spill ? "amdgcn-mesa-mesa3d" : "amdgcn--");
LLVMDisposeTargetData(data_layout);
LLVMDisposeMessage(data_layout_str);
- setup_types(&ctx);
- ctx.builder = LLVMCreateBuilderInContext(ctx.context);
+ enum ac_float_mode float_mode =
+ options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
+ AC_FLOAT_MODE_DEFAULT;
+
+ ctx.builder = ac_create_builder(ctx.context, float_mode);
ctx.ac.builder = ctx.builder;
memset(shader_info, 0, sizeof(*shader_info));
for (i = 0; i < AC_UD_MAX_UD; i++)
shader_info->user_sgprs_locs.shader_data[i].sgpr_idx = -1;
- ctx.max_workgroup_size = ac_nir_get_max_workgroup_size(ctx.options->chip_class, shaders[0]);
+ ctx.max_workgroup_size = 0;
+ for (int i = 0; i < shader_count; ++i) {
+ ctx.max_workgroup_size = MAX2(ctx.max_workgroup_size,
+ ac_nir_get_max_workgroup_size(ctx.options->chip_class,
+ shaders[i]));
+ }
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.abi.inputs = &ctx.inputs[0];
ctx.abi.emit_outputs = handle_shader_outputs_post;
+ ctx.abi.emit_vertex = visit_emit_vertex;
+ ctx.abi.load_ubo = radv_load_ubo;
ctx.abi.load_ssbo = radv_load_ssbo;
ctx.abi.load_sampler_desc = radv_get_sampler_desc;
+ ctx.abi.clamp_shadow_reference = false;
if (shader_count >= 2)
ac_init_exec_full_mask(&ctx.ac);
ctx.num_output_culls = shaders[i]->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");
-
+ ctx.gs_next_vertex = ac_build_alloca(&ctx.ac, ctx.ac.i32, "gs_next_vertex");
ctx.gs_max_out_vertices = shaders[i]->info.gs.vertices_out;
+ ctx.abi.load_inputs = load_gs_input;
+ ctx.abi.emit_primitive = visit_end_primitive;
+ } else if (shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) {
+ ctx.tcs_outputs_read = shaders[i]->info.outputs_read;
+ ctx.tcs_patch_outputs_read = shaders[i]->info.patch_outputs_read;
+ ctx.abi.load_tess_varyings = load_tcs_varyings;
+ ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
+ ctx.abi.store_tcs_outputs = store_tcs_output;
} else if (shaders[i]->info.stage == MESA_SHADER_TESS_EVAL) {
ctx.tes_primitive_mode = shaders[i]->info.tess.primitive_mode;
+ ctx.abi.load_tess_varyings = load_tes_input;
+ ctx.abi.load_tess_coord = load_tess_coord;
+ ctx.abi.load_patch_vertices_in = load_patch_vertices_in;
} 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);
+ if (ctx.options->key.vs.as_ls) {
+ ctx.shader_info->vs.vgpr_comp_cnt =
+ MAX2(2, ctx.shader_info->vs.vgpr_comp_cnt);
+ } else {
+ ctx.shader_info->vs.vgpr_comp_cnt =
+ MAX2(1, 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.abi.lookup_interp_param = lookup_interp_param;
+ ctx.abi.load_sample_position = load_sample_position;
}
if (i)
- emit_barrier(&ctx);
+ emit_barrier(&ctx.ac, ctx.stage);
ac_setup_rings(&ctx);
LLVMBuildRetVoid(ctx.builder);
+ if (options->dump_preoptir)
+ ac_dump_module(ctx.module);
+
ac_llvm_finalize_module(&ctx);
if (shader_count == 1)
/* +3 for scratch wave offset and VCC */
config->num_sgprs = MAX2(config->num_sgprs,
shader_info->num_input_sgprs + 3);
+
+ /* Enable 64-bit and 16-bit denormals, because there is no performance
+ * cost.
+ *
+ * If denormals are enabled, all floating-point output modifiers are
+ * ignored.
+ *
+ * Don't enable denormals for 32-bit floats, because:
+ * - Floating-point output modifiers would be ignored by the hw.
+ * - Some opcodes don't support denormals, such as v_mad_f32. We would
+ * have to stop using those.
+ * - SI & CI would be very slow.
+ */
+ config->float_mode |= V_00B028_FP_64_DENORMS;
}
static void
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 */
+ /* in LS mode we need at least 1, invocation id needs 2, handled elsewhere */
if (options->key.vs.as_ls)
shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt);
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);
+
+ /* Determine the ES type (VS or TES) for the GS on GFX9. */
+ if (options->chip_class == GFX9) {
+ if (nir_count == 2 &&
+ nir[1]->info.stage == MESA_SHADER_GEOMETRY) {
+ shader_info->gs.es_type = nir[0]->info.stage;
+ }
+ }
}
static void
ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx)
{
- LLVMValueRef args[9];
- args[0] = ctx->gsvs_ring;
- args[1] = LLVMBuildMul(ctx->builder, ctx->abi.vertex_id, LLVMConstInt(ctx->i32, 4, false), "");
- args[3] = ctx->i32zero;
- args[4] = ctx->i32one; /* OFFEN */
- args[5] = ctx->i32zero; /* IDXEN */
- args[6] = ctx->i32one; /* GLC */
- args[7] = ctx->i32one; /* SLC */
- args[8] = ctx->i32zero; /* TFE */
-
+ LLVMValueRef vtx_offset =
+ LLVMBuildMul(ctx->builder, ctx->abi.vertex_id,
+ LLVMConstInt(ctx->ac.i32, 4, false), "");
int idx = 0;
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
}
for (unsigned j = 0; j < length; j++) {
- LLVMValueRef value;
- args[2] = LLVMConstInt(ctx->i32,
+ LLVMValueRef value, soffset;
+
+ soffset = LLVMConstInt(ctx->ac.i32,
(slot * 4 + j) *
ctx->gs_max_out_vertices * 16 * 4, false);
- value = ac_build_intrinsic(&ctx->ac,
- "llvm.SI.buffer.load.dword.i32.i32",
- ctx->i32, args, 9,
- AC_FUNC_ATTR_READONLY |
- AC_FUNC_ATTR_LEGACY);
+ value = ac_build_buffer_load(&ctx->ac, ctx->gsvs_ring,
+ 1, ctx->ac.i32_0,
+ vtx_offset, soffset,
+ 0, 1, 1, true, false);
LLVMBuildStore(ctx->builder,
ac_to_float(&ctx->ac, value), ctx->nir->outputs[radeon_llvm_reg_index_soa(i, j)]);
ctx.options = options;
ctx.shader_info = shader_info;
- ac_llvm_context_init(&ctx.ac, ctx.context, options->chip_class);
+ ac_llvm_context_init(&ctx.ac, ctx.context, options->chip_class,
+ options->family);
ctx.ac.module = ctx.module;
ctx.is_gs_copy_shader = true;
LLVMSetTarget(ctx.module, "amdgcn--");
- setup_types(&ctx);
- ctx.builder = LLVMCreateBuilderInContext(ctx.context);
+ enum ac_float_mode float_mode =
+ options->unsafe_math ? AC_FLOAT_MODE_UNSAFE_FP_MATH :
+ AC_FLOAT_MODE_DEFAULT;
+
+ ctx.builder = ac_create_builder(ctx.context, float_mode);
ctx.ac.builder = ctx.builder;
ctx.stage = MESA_SHADER_VERTEX;
projects / mesa.git / blobdiff