struct ac_llvm_context ac;
const struct ac_nir_compiler_options *options;
struct ac_shader_variant_info *shader_info;
-
+ unsigned max_workgroup_size;
LLVMContextRef context;
LLVMModuleRef module;
LLVMBuilderRef builder;
LLVMBuilderRef builder, LLVMTypeRef *return_types,
unsigned num_return_elems, LLVMTypeRef *param_types,
unsigned param_count, unsigned array_params_mask,
- unsigned sgpr_params, bool unsafe_math)
+ unsigned sgpr_params, unsigned max_workgroup_size,
+ bool unsafe_math)
{
LLVMTypeRef main_function_type, ret_type;
LLVMBasicBlockRef main_function_body;
}
}
+ if (max_workgroup_size) {
+ ac_llvm_add_target_dep_function_attr(main_function,
+ "amdgpu-max-work-group-size",
+ max_workgroup_size);
+ }
if (unsafe_math) {
/* These were copied from some LLVM test. */
LLVMAddTargetDependentFunctionAttr(main_function,
ctx->main_function = create_llvm_function(
ctx->context, ctx->module, ctx->builder, NULL, 0, arg_types,
- arg_idx, array_params_mask, sgpr_count, ctx->options->unsafe_math);
+ arg_idx, array_params_mask, sgpr_count, ctx->max_workgroup_size,
+ ctx->options->unsafe_math);
set_llvm_calling_convention(ctx->main_function, ctx->stage);
ctx->shader_info->num_input_sgprs = 0;
case nir_op_fmax:
result = emit_intrin_2f_param(ctx, "llvm.maxnum",
to_float_type(ctx, def_type), src[0], src[1]);
+ if (instr->dest.dest.ssa.bit_size == 32)
+ result = emit_intrin_1f_param(ctx, "llvm.canonicalize",
+ to_float_type(ctx, def_type),
+ result);
break;
case nir_op_fmin:
result = emit_intrin_2f_param(ctx, "llvm.minnum",
to_float_type(ctx, def_type), src[0], src[1]);
+ if (instr->dest.dest.ssa.bit_size == 32)
+ result = emit_intrin_1f_param(ctx, "llvm.canonicalize",
+ to_float_type(ctx, def_type),
+ result);
break;
case nir_op_ffma:
result = emit_intrin_3f_param(ctx, "llvm.fma",
is_compact, vertex_index, indir_index);
result = ac_build_buffer_load(&ctx->ac, ctx->hs_ring_tess_offchip, instr->num_components, NULL,
- buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true);
+ buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, 1, 0, true, false);
result = trim_vector(ctx, result, instr->num_components);
result = LLVMBuildBitCast(ctx->builder, result, get_def_type(ctx, &instr->dest.ssa), "");
return result;
outinfo = &ctx->shader_info->tes.outinfo;
}
- ac_eliminate_const_vs_outputs(&ctx->ac,
- ctx->main_function,
- outinfo->vs_output_param_offset,
- VARYING_SLOT_MAX,
- &outinfo->param_exports);
+ ac_optimize_vs_outputs(&ctx->ac,
+ ctx->main_function,
+ outinfo->vs_output_param_offset,
+ VARYING_SLOT_MAX,
+ &outinfo->param_exports);
}
static void
}
}
+static unsigned
+ac_nir_get_max_workgroup_size(enum chip_class chip_class,
+ struct nir_shader *nir)
+{
+ switch (nir->stage) {
+ case MESA_SHADER_TESS_CTRL:
+ return chip_class >= CIK ? 128 : 64;
+ case MESA_SHADER_GEOMETRY:
+ return 64;
+ case MESA_SHADER_COMPUTE:
+ break;
+ default:
+ return 0;
+ }
+
+ unsigned max_workgroup_size = nir->info.cs.local_size[0] *
+ nir->info.cs.local_size[1] *
+ nir->info.cs.local_size[2];
+ return max_workgroup_size;
+}
+
static
LLVMModuleRef ac_translate_nir_to_llvm(LLVMTargetMachineRef tm,
struct nir_shader *nir,
ctx.builder = LLVMCreateBuilderInContext(ctx.context);
ctx.ac.builder = ctx.builder;
ctx.stage = nir->stage;
+ ctx.max_workgroup_size = ac_nir_get_max_workgroup_size(ctx.options->chip_class, nir);
for (i = 0; i < AC_UD_MAX_SETS; i++)
shader_info->user_sgprs_locs.descriptor_sets[i].sgpr_idx = -1;
} else if (nir->stage == MESA_SHADER_GEOMETRY) {
ctx.gs_next_vertex = ac_build_alloca(&ctx, ctx.i32, "gs_next_vertex");
- ctx.gs_max_out_vertices = nir->info->gs.vertices_out;
+ ctx.gs_max_out_vertices = nir->info.gs.vertices_out;
} else if (nir->stage == MESA_SHADER_TESS_EVAL) {
- ctx.tes_primitive_mode = nir->info->tess.primitive_mode;
+ ctx.tes_primitive_mode = nir->info.tess.primitive_mode;
}
ac_setup_rings(&ctx);
if (nir->stage == MESA_SHADER_FRAGMENT)
handle_fs_inputs_pre(&ctx, nir);
- ctx.num_output_clips = nir->info->clip_distance_array_size;
- ctx.num_output_culls = nir->info->cull_distance_array_size;
+ ctx.num_output_clips = nir->info.clip_distance_array_size;
+ ctx.num_output_culls = nir->info.cull_distance_array_size;
nir_foreach_variable(variable, &nir->outputs)
handle_shader_output_decl(&ctx, variable);
unsigned addclip = ctx.num_output_clips + ctx.num_output_culls > 4;
shader_info->gs.gsvs_vertex_size = (util_bitcount64(ctx.output_mask) + addclip) * 16;
shader_info->gs.max_gsvs_emit_size = shader_info->gs.gsvs_vertex_size *
- nir->info->gs.vertices_out;
+ nir->info.gs.vertices_out;
} else if (nir->stage == MESA_SHADER_TESS_CTRL) {
shader_info->tcs.outputs_written = ctx.tess_outputs_written;
shader_info->tcs.patch_outputs_written = ctx.tess_patch_outputs_written;
switch (nir->stage) {
case MESA_SHADER_COMPUTE:
for (int i = 0; i < 3; ++i)
- shader_info->cs.block_size[i] = nir->info->cs.local_size[i];
+ shader_info->cs.block_size[i] = nir->info.cs.local_size[i];
break;
case MESA_SHADER_FRAGMENT:
- shader_info->fs.early_fragment_test = nir->info->fs.early_fragment_tests;
+ shader_info->fs.early_fragment_test = nir->info.fs.early_fragment_tests;
break;
case MESA_SHADER_GEOMETRY:
- shader_info->gs.vertices_in = nir->info->gs.vertices_in;
- shader_info->gs.vertices_out = nir->info->gs.vertices_out;
- shader_info->gs.output_prim = nir->info->gs.output_primitive;
- shader_info->gs.invocations = nir->info->gs.invocations;
+ shader_info->gs.vertices_in = nir->info.gs.vertices_in;
+ shader_info->gs.vertices_out = nir->info.gs.vertices_out;
+ shader_info->gs.output_prim = nir->info.gs.output_primitive;
+ shader_info->gs.invocations = nir->info.gs.invocations;
break;
case MESA_SHADER_TESS_EVAL:
- shader_info->tes.primitive_mode = nir->info->tess.primitive_mode;
- shader_info->tes.spacing = nir->info->tess.spacing;
- shader_info->tes.ccw = nir->info->tess.ccw;
- shader_info->tes.point_mode = nir->info->tess.point_mode;
+ shader_info->tes.primitive_mode = nir->info.tess.primitive_mode;
+ shader_info->tes.spacing = nir->info.tess.spacing;
+ shader_info->tes.ccw = nir->info.tess.ccw;
+ shader_info->tes.point_mode = nir->info.tess.point_mode;
shader_info->tes.as_es = options->key.tes.as_es;
break;
case MESA_SHADER_TESS_CTRL:
- shader_info->tcs.tcs_vertices_out = nir->info->tess.tcs_vertices_out;
+ shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
break;
case MESA_SHADER_VERTEX:
shader_info->vs.as_es = options->key.vs.as_es;
create_function(&ctx);
- ctx.gs_max_out_vertices = geom_shader->info->gs.vertices_out;
+ ctx.gs_max_out_vertices = geom_shader->info.gs.vertices_out;
ac_setup_rings(&ctx);
- ctx.num_output_clips = geom_shader->info->clip_distance_array_size;
- ctx.num_output_culls = geom_shader->info->cull_distance_array_size;
+ ctx.num_output_clips = geom_shader->info.clip_distance_array_size;
+ ctx.num_output_culls = geom_shader->info.cull_distance_array_size;
nir_foreach_variable(variable, &geom_shader->outputs)
handle_shader_output_decl(&ctx, variable);
projects / mesa.git / blobdiff