#include <llvm-c/Transforms/Utils.h>
#include "sid.h"
-#include "gfx9d.h"
#include "ac_binary.h"
#include "ac_llvm_util.h"
#include "ac_llvm_build.h"
*
* Test: dEQP-VK.tessellation.shader_input_output.barrier
*/
- if (ctx->options->chip_class >= CIK && ctx->options->family != CHIP_STONEY)
+ if (ctx->options->chip_class >= GFX7 && ctx->options->family != CHIP_STONEY)
hardware_lds_size = 65536;
num_patches = MIN2(num_patches, hardware_lds_size / (input_patch_size + output_patch_size));
*/
num_patches = MIN2(num_patches, 40);
- /* SI bug workaround - limit LS-HS threadgroups to only one wave. */
- if (ctx->options->chip_class == SI) {
+ /* GFX6 bug workaround - limit LS-HS threadgroups to only one wave. */
+ if (ctx->options->chip_class == GFX6) {
unsigned one_wave = 64 / MAX2(num_tcs_input_cp, num_tcs_output_cp);
num_patches = MIN2(num_patches, one_wave);
}
options->address32_hi);
}
- if (max_workgroup_size) {
- ac_llvm_add_target_dep_function_attr(main_function,
- "amdgpu-max-work-group-size",
- max_workgroup_size);
- }
+ ac_llvm_set_workgroup_size(main_function, max_workgroup_size);
+
if (options->unsafe_math) {
/* These were copied from some LLVM test. */
LLVMAddTargetDependentFunctionAttr(main_function,
} else
stride = LLVMConstInt(ctx->ac.i32, layout->binding[binding].size, false);
- offset = ac_build_imad(&ctx->ac, index, stride,
- LLVMConstInt(ctx->ac.i32, base_offset, false));
+ offset = LLVMConstInt(ctx->ac.i32, base_offset, false);
- desc_ptr = ac_build_gep0(&ctx->ac, desc_ptr, offset);
+ if (layout->binding[binding].type != VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
+ offset = ac_build_imad(&ctx->ac, index, stride, offset);
+ }
+
+ desc_ptr = LLVMBuildGEP(ctx->ac.builder, desc_ptr, &offset, 1, "");
desc_ptr = ac_cast_ptr(&ctx->ac, desc_ptr, ctx->ac.v4i32);
LLVMSetMetadata(desc_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
+ if (layout->binding[binding].type == VK_DESCRIPTOR_TYPE_INLINE_UNIFORM_BLOCK_EXT) {
+ uint32_t desc_type = S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
+ S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
+ S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
+ S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) |
+ S_008F0C_NUM_FORMAT(V_008F0C_BUF_NUM_FORMAT_FLOAT) |
+ S_008F0C_DATA_FORMAT(V_008F0C_BUF_DATA_FORMAT_32);
+
+ LLVMValueRef desc_components[4] = {
+ LLVMBuildPtrToInt(ctx->ac.builder, desc_ptr, ctx->ac.intptr, ""),
+ LLVMConstInt(ctx->ac.i32, S_008F04_BASE_ADDRESS_HI(ctx->options->address32_hi), false),
+ /* High limit to support variable sizes. */
+ LLVMConstInt(ctx->ac.i32, 0xffffffff, false),
+ LLVMConstInt(ctx->ac.i32, desc_type, false),
+ };
+
+ return ac_build_gather_values(&ctx->ac, desc_components, 4);
+ }
+
return desc_ptr;
}
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);
+ 4 * (base + chan), ac_glc, false);
}
if (writemask == 0xF) {
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, src, 4,
buf_addr, ctx->oc_lds,
- (base * 4), 1, 0, true, false);
+ (base * 4), ac_glc, false);
}
}
buf_addr = LLVMBuildAdd(ctx->ac.builder, buf_addr, comp_offset, "");
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);
+ buf_addr, ctx->oc_lds, is_compact ? (4 * const_index) : 0, ac_glc, true, false);
result = ac_trim_vector(&ctx->ac, result, num_components);
return result;
}
ctx->esgs_ring, 1,
ctx->ac.i32_0,
vtx_offset, soffset,
- 0, 1, 0, true, false);
+ 0, ac_glc, true, false);
}
if (ac_get_type_size(type) == 2) {
struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
LLVMValueRef result;
- LLVMValueRef ptr = ac_build_gep0(&ctx->ac, ctx->ring_offsets, LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false));
+ LLVMValueRef index = LLVMConstInt(ctx->ac.i32, RING_PS_SAMPLE_POSITIONS, false);
+ LLVMValueRef ptr = LLVMBuildGEP(ctx->ac.builder, ctx->ring_offsets, &index, 1, "");
ptr = LLVMBuildBitCast(ctx->ac.builder, ptr,
ac_array_in_const_addr_space(ctx->ac.v2f32), "");
ctx->gsvs_ring[stream],
out_val, 1,
voffset, ctx->gs2vs_offset, 0,
- 1, 1, true, true);
+ ac_glc | ac_slc, true);
}
}
struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
LLVMValueRef result;
+ if (LLVMGetTypeKind(LLVMTypeOf(buffer_ptr)) != LLVMPointerTypeKind) {
+ /* Do not load the descriptor for inlined uniform blocks. */
+ return buffer_ptr;
+ }
+
LLVMSetMetadata(buffer_ptr, ctx->ac.uniform_md_kind, ctx->ac.empty_md);
result = LLVMBuildLoad(ctx->ac.builder, buffer_ptr, "");
break;
case AC_DESC_SAMPLER:
type = ctx->ac.v4i32;
- if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER)
- offset += 64;
+ if (binding->type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) {
+ offset += radv_combined_image_descriptor_sampler_offset(binding);
+ }
type_size = 16;
break;
type = ctx->ac.v4i32;
type_size = 16;
break;
+ case AC_DESC_PLANE_0:
+ case AC_DESC_PLANE_1:
+ case AC_DESC_PLANE_2:
+ type = ctx->ac.v8i32;
+ type_size = 32;
+ offset += 32 * (desc_type - AC_DESC_PLANE_0);
+ break;
default:
unreachable("invalid desc_type\n");
}
assert(stride % type_size == 0);
- if (!index)
- index = ctx->ac.i32_0;
+ LLVMValueRef adjusted_index = index;
+ if (!adjusted_index)
+ adjusted_index = ctx->ac.i32_0;
- index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
+ adjusted_index = LLVMBuildMul(builder, adjusted_index, LLVMConstInt(ctx->ac.i32, stride / type_size, 0), "");
- list = ac_build_gep0(&ctx->ac, list, LLVMConstInt(ctx->ac.i32, offset, 0));
+ LLVMValueRef val_offset = LLVMConstInt(ctx->ac.i32, offset, 0);
+ list = LLVMBuildGEP(builder, list, &val_offset, 1, "");
list = LLVMBuildPointerCast(builder, list,
ac_array_in_const32_addr_space(type), "");
- return ac_build_load_to_sgpr(&ctx->ac, list, index);
+ LLVMValueRef descriptor = ac_build_load_to_sgpr(&ctx->ac, list, adjusted_index);
+
+ /* 3 plane formats always have same size and format for plane 1 & 2, so
+ * use the tail from plane 1 so that we can store only the first 16 bytes
+ * of the last plane. */
+ if (desc_type == AC_DESC_PLANE_2) {
+ LLVMValueRef descriptor2 = radv_get_sampler_desc(abi, descriptor_set, base_index, constant_index, index, AC_DESC_PLANE_1,image, write, bindless);
+
+ LLVMValueRef components[8];
+ for (unsigned i = 0; i < 4; ++i)
+ components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor, i);
+
+ for (unsigned i = 4; i < 8; ++i)
+ components[i] = ac_llvm_extract_elem(&ctx->ac, descriptor2, i);
+ descriptor = ac_build_gather_values(&ctx->ac, components, 8);
+ }
+
+ return descriptor;
}
/* For 2_10_10_10 formats the alpha is handled as unsigned by pre-vega HW.
LLVMValueRef c30 = LLVMConstInt(ctx->ac.i32, 30, 0);
+ alpha = LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.f32, "");
+
if (adjustment == RADV_ALPHA_ADJUST_SSCALED)
alpha = LLVMBuildFPToUI(ctx->ac.builder, alpha, ctx->ac.i32, "");
else
alpha = LLVMBuildSIToFP(ctx->ac.builder, alpha, ctx->ac.f32, "");
}
- return alpha;
+ return LLVMBuildBitCast(ctx->ac.builder, alpha, ctx->ac.i32, "");
}
static unsigned
for (unsigned i = num_channels; i < 4; i++) {
chan[i] = i == 3 ? one : zero;
- chan[i] = ac_to_float(&ctx->ac, chan[i]);
+ chan[i] = ac_to_integer(&ctx->ac, chan[i]);
}
return ac_build_gather_values(&ctx->ac, chan, 4);
buffer_index = LLVMBuildUDiv(ctx->ac.builder, buffer_index,
LLVMConstInt(ctx->ac.i32, divisor, 0), "");
}
-
- 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 = ctx->ac.i32_0;
}
} else
buffer_index = LLVMBuildAdd(ctx->ac.builder, ctx->abi.vertex_id,
ctx->abi.base_vertex, "");
- t_offset = LLVMConstInt(ctx->ac.i32, attrib_index, false);
-
- t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
/* Adjust the number of channels to load based on the vertex
* attribute format.
*/
unsigned num_format_channels = get_num_channels_from_data_format(data_format);
unsigned num_channels = MIN2(num_input_channels, num_format_channels);
+ unsigned attrib_binding = ctx->options->key.vs.vertex_attribute_bindings[attrib_index];
+ unsigned attrib_offset = ctx->options->key.vs.vertex_attribute_offsets[attrib_index];
+ unsigned attrib_stride = ctx->options->key.vs.vertex_attribute_strides[attrib_index];
+
+ if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
+ /* Always load, at least, 3 channels for formats that
+ * need to be shuffled because X<->Z.
+ */
+ num_channels = MAX2(num_channels, 3);
+ }
- input = ac_build_buffer_load_format(&ctx->ac, t_list,
+ if (attrib_stride != 0 && attrib_offset > attrib_stride) {
+ LLVMValueRef buffer_offset =
+ LLVMConstInt(ctx->ac.i32,
+ attrib_offset / attrib_stride, false);
+
+ buffer_index = LLVMBuildAdd(ctx->ac.builder,
buffer_index,
- ctx->ac.i32_0,
- num_channels, false, true);
+ buffer_offset, "");
+
+ attrib_offset = attrib_offset % attrib_stride;
+ }
+
+ t_offset = LLVMConstInt(ctx->ac.i32, attrib_binding, false);
+ t_list = ac_build_load_to_sgpr(&ctx->ac, t_list_ptr, t_offset);
+
+ input = ac_build_struct_tbuffer_load(&ctx->ac, t_list,
+ buffer_index,
+ LLVMConstInt(ctx->ac.i32, attrib_offset, false),
+ ctx->ac.i32_0, ctx->ac.i32_0,
+ num_channels,
+ data_format, num_format, 0, true);
+
+ if (ctx->options->key.vs.post_shuffle & (1 << attrib_index)) {
+ LLVMValueRef c[4];
+ c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2);
+ c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1);
+ c[2] = ac_llvm_extract_elem(&ctx->ac, input, 0);
+ c[3] = ac_llvm_extract_elem(&ctx->ac, input, 3);
+
+ input = ac_build_gather_values(&ctx->ac, c, 4);
+ }
input = radv_fixup_vertex_input_fetches(ctx, input, num_channels,
is_float);
mark_16bit_fs_input(ctx, glsl_get_array_element(type), location + i * stride);
}
} else {
- assert(glsl_type_is_struct(type));
+ assert(glsl_type_is_struct_or_ifc(type));
for (unsigned i = 0; i < glsl_get_length(type); i++) {
mark_16bit_fs_input(ctx, glsl_get_struct_field(type, i), location);
location += glsl_count_attribute_slots(glsl_get_struct_field(type, i), false);
break;
case V_028714_SPI_SHADER_32_AR:
- args->enabled_channels = 0x9;
- args->out[0] = values[0];
- args->out[3] = values[3];
+ if (ctx->ac.chip_class >= GFX10) {
+ args->enabled_channels = 0x3;
+ args->out[0] = values[0];
+ args->out[1] = values[3];
+ } else {
+ args->enabled_channels = 0x9;
+ args->out[0] = values[0];
+ args->out[3] = values[3];
+ }
break;
case V_028714_SPI_SHADER_FP16_ABGR:
/* fall through */
case 4: /* as v4i32 */
vdata = ac_build_gather_values(&ctx->ac, out,
- util_next_power_of_two(num_comps));
+ !ac_has_vec3_support(ctx->ac.chip_class, false) ?
+ util_next_power_of_two(num_comps) :
+ num_comps);
break;
}
ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf],
vdata, num_comps, so_write_offsets[buf],
ctx->ac.i32_0, offset,
- 1, 1, true, false);
+ ac_glc | ac_slc, false);
}
static void
static void
handle_vs_outputs_post(struct radv_shader_context *ctx,
bool export_prim_id, bool export_layer_id,
+ bool export_clip_dists,
struct radv_vs_output_info *outinfo)
{
uint32_t param_count = 0;
memcpy(&pos_args[target - V_008DFC_SQ_EXP_POS],
&args, sizeof(args));
- /* Export the clip/cull distances values to the next stage. */
- radv_export_param(ctx, param_count, &slots[0], 0xf);
- outinfo->vs_output_param_offset[location] = param_count++;
+ if (export_clip_dists) {
+ /* Export the clip/cull distances values to the next stage. */
+ radv_export_param(ctx, param_count, &slots[0], 0xf);
+ outinfo->vs_output_param_offset[location] = param_count++;
+ }
}
}
num_pos_exports++;
}
+ /* Navi10-14 skip POS0 exports if EXEC=0 and DONE=0, causing a hang.
+ * Setting valid_mask=1 prevents it and has no other effect.
+ */
+ if (ctx->ac.family == CHIP_NAVI10 ||
+ ctx->ac.family == CHIP_NAVI12 ||
+ ctx->ac.family == CHIP_NAVI14)
+ pos_args[0].valid_mask = 1;
+
pos_idx = 0;
for (i = 0; i < 4; i++) {
if (!pos_args[i].out[0])
LLVMValueRef values[4];
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->ac.f32_0;
out_val, 1,
NULL, ctx->es2gs_offset,
(4 * param_index + j) * 4,
- 1, 1, true, true);
+ ac_glc | ac_slc, true);
}
}
}
LLVMConstInt(ctx->ac.i32, 4 * stride, false), "");
unsigned tf_offset = 0;
- if (ctx->options->chip_class <= VI) {
+ if (ctx->options->chip_class <= GFX8) {
ac_nir_build_if(&inner_if_ctx, ctx,
LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
rel_patch_id, ctx->ac.i32_0, ""));
ac_build_buffer_store_dword(&ctx->ac, buffer,
LLVMConstInt(ctx->ac.i32, 0x80000000, false),
1, ctx->ac.i32_0, tf_base,
- 0, 1, 0, true, false);
+ 0, ac_glc, false);
tf_offset += 4;
ac_nir_build_endif(&inner_if_ctx);
/* Store the tessellation factors. */
ac_build_buffer_store_dword(&ctx->ac, buffer, vec0,
MIN2(stride, 4), byteoffset, tf_base,
- tf_offset, 1, 0, true, false);
+ tf_offset, ac_glc, false);
if (vec1)
ac_build_buffer_store_dword(&ctx->ac, buffer, vec1,
stride - 4, byteoffset, tf_base,
- 16 + tf_offset, 1, 0, true, false);
+ 16 + tf_offset, ac_glc, false);
//store to offchip for TES to read - only if TES reads them
if (ctx->options->key.tcs.tes_reads_tess_factors) {
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, outer_vec,
outer_comps, tf_outer_offset,
- ctx->oc_lds, 0, 1, 0, true, false);
+ ctx->oc_lds, 0, ac_glc, false);
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,
ac_build_gather_values(&ctx->ac, inner, inner_comps);
ac_build_buffer_store_dword(&ctx->ac, ctx->hs_ring_tess_offchip, inner_vec,
inner_comps, tf_inner_offset,
- ctx->oc_lds, 0, 1, 0, true, false);
+ ctx->oc_lds, 0, ac_glc, false);
}
}
ac_nir_build_endif(&if_ctx);
else
handle_vs_outputs_post(ctx, ctx->options->key.vs.export_prim_id,
ctx->options->key.vs.export_layer_id,
+ ctx->options->key.vs.export_clip_dists,
&ctx->shader_info->vs.outinfo);
break;
case MESA_SHADER_FRAGMENT:
else
handle_vs_outputs_post(ctx, ctx->options->key.tes.export_prim_id,
ctx->options->key.tes.export_layer_id,
+ ctx->options->key.tes.export_clip_dists,
&ctx->shader_info->tes.outinfo);
break;
default:
static void
ac_setup_rings(struct radv_shader_context *ctx)
{
- if (ctx->options->chip_class <= VI &&
+ if (ctx->options->chip_class <= GFX8 &&
(ctx->stage == MESA_SHADER_GEOMETRY ||
ctx->options->key.vs.as_es || ctx->options->key.tes.as_es)) {
unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS
stride = 4 * num_components * ctx->gs_max_out_vertices;
- /* Limit on the stride field for <= CIK. */
+ /* Limit on the stride field for <= GFX7. */
assert(stride < (1 << 14));
ring = LLVMBuildBitCast(ctx->ac.builder,
{
switch (nir->info.stage) {
case MESA_SHADER_TESS_CTRL:
- return chip_class >= CIK ? 128 : 64;
+ return chip_class >= GFX7 ? 128 : 64;
case MESA_SHADER_GEOMETRY:
return chip_class >= GFX9 ? 128 : 64;
case MESA_SHADER_COMPUTE:
ctx.abi.clamp_shadow_reference = false;
ctx.abi.gfx9_stride_size_workaround = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x800;
+ /* Because the new raw/struct atomic intrinsics are buggy with LLVM 8,
+ * we fallback to the old intrinsics for atomic buffer image operations
+ * and thus we need to apply the indexing workaround...
+ */
+ ctx.abi.gfx9_stride_size_workaround_for_atomic = ctx.ac.chip_class == GFX9 && HAVE_LLVM < 0x900;
+
if (shader_count >= 2)
ac_init_exec_full_mask(&ctx.ac);
- if (ctx.ac.chip_class == GFX9 &&
+ if ((ctx.ac.family == CHIP_VEGA10 ||
+ ctx.ac.family == CHIP_RAVEN) &&
shaders[shader_count - 1]->info.stage == MESA_SHADER_TESS_CTRL)
ac_nir_fixup_ls_hs_input_vgprs(&ctx);
ctx.tcs_vertices_per_patch = shaders[i]->info.tess.tcs_vertices_out;
ctx.tcs_num_patches = ctx.options->key.tes.num_patches;
} else if (shaders[i]->info.stage == MESA_SHADER_VERTEX) {
- if (shader_info->info.vs.needs_instance_id) {
- 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);
- }
- }
ctx.abi.load_base_vertex = radv_load_base_vertex;
} else if (shaders[i]->info.stage == MESA_SHADER_FRAGMENT) {
shader_info->fs.can_discard = shaders[i]->info.fs.uses_discard;
LLVMDisposeMessage(description);
}
-static unsigned ac_llvm_compile(LLVMModuleRef M,
- struct ac_shader_binary *binary,
- struct ac_llvm_compiler *ac_llvm)
+static unsigned radv_llvm_compile(LLVMModuleRef M,
+ char **pelf_buffer, size_t *pelf_size,
+ struct ac_llvm_compiler *ac_llvm)
{
unsigned retval = 0;
LLVMContextRef llvm_ctx;
&retval);
/* Compile IR*/
- if (!radv_compile_to_binary(ac_llvm, M, binary))
+ if (!radv_compile_to_elf(ac_llvm, M, pelf_buffer, pelf_size))
retval = 1;
return retval;
}
static void ac_compile_llvm_module(struct ac_llvm_compiler *ac_llvm,
LLVMModuleRef llvm_module,
- struct ac_shader_binary *binary,
- struct ac_shader_config *config,
+ struct radv_shader_binary **rbinary,
struct radv_shader_variant_info *shader_info,
gl_shader_stage stage,
const struct radv_nir_compiler_options *options)
{
+ char *elf_buffer = NULL;
+ size_t elf_size = 0;
+ char *llvm_ir_string = NULL;
if (options->dump_shader)
ac_dump_module(llvm_module);
- memset(binary, 0, sizeof(*binary));
-
if (options->record_llvm_ir) {
char *llvm_ir = LLVMPrintModuleToString(llvm_module);
- binary->llvm_ir_string = strdup(llvm_ir);
+ llvm_ir_string = strdup(llvm_ir);
LLVMDisposeMessage(llvm_ir);
}
- int v = ac_llvm_compile(llvm_module, binary, ac_llvm);
+ int v = radv_llvm_compile(llvm_module, &elf_buffer, &elf_size, ac_llvm);
if (v) {
fprintf(stderr, "compile failed\n");
}
- if (options->dump_shader)
- fprintf(stderr, "disasm:\n%s\n", binary->disasm_string);
-
- ac_shader_binary_read_config(binary, config, 0, options->supports_spill);
-
LLVMContextRef ctx = LLVMGetModuleContext(llvm_module);
LLVMDisposeModule(llvm_module);
LLVMContextDispose(ctx);
- if (stage == MESA_SHADER_FRAGMENT) {
- shader_info->num_input_vgprs = 0;
- if (G_0286CC_PERSP_SAMPLE_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 2;
- if (G_0286CC_PERSP_CENTER_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 2;
- if (G_0286CC_PERSP_CENTROID_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 2;
- if (G_0286CC_PERSP_PULL_MODEL_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 3;
- if (G_0286CC_LINEAR_SAMPLE_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 2;
- if (G_0286CC_LINEAR_CENTER_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 2;
- if (G_0286CC_LINEAR_CENTROID_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 2;
- if (G_0286CC_LINE_STIPPLE_TEX_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_POS_X_FLOAT_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_POS_Y_FLOAT_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_POS_Z_FLOAT_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_POS_W_FLOAT_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_FRONT_FACE_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_ANCILLARY_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_SAMPLE_COVERAGE_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- if (G_0286CC_POS_FIXED_PT_ENA(config->spi_ps_input_addr))
- shader_info->num_input_vgprs += 1;
- }
- config->num_vgprs = MAX2(config->num_vgprs, shader_info->num_input_vgprs);
-
- /* +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;
+ size_t llvm_ir_size = llvm_ir_string ? strlen(llvm_ir_string) : 0;
+ size_t alloc_size = sizeof(struct radv_shader_binary_rtld) + elf_size + llvm_ir_size + 1;
+ struct radv_shader_binary_rtld *rbin = calloc(1, alloc_size);
+ memcpy(rbin->data, elf_buffer, elf_size);
+ if (llvm_ir_string)
+ memcpy(rbin->data + elf_size, llvm_ir_string, llvm_ir_size + 1);
+
+ rbin->base.type = RADV_BINARY_TYPE_RTLD;
+ rbin->base.stage = stage;
+ rbin->base.total_size = alloc_size;
+ rbin->elf_size = elf_size;
+ rbin->llvm_ir_size = llvm_ir_size;
+ *rbinary = &rbin->base;
+
+ free(llvm_ir_string);
+ free(elf_buffer);
}
static void
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;
+ shader_info->tes.export_prim_id = options->key.tes.export_prim_id;
break;
case MESA_SHADER_TESS_CTRL:
shader_info->tcs.tcs_vertices_out = nir->info.tess.tcs_vertices_out;
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 2, handled elsewhere */
- if (options->key.vs.as_ls)
- shader_info->vs.vgpr_comp_cnt = MAX2(1, shader_info->vs.vgpr_comp_cnt);
+ shader_info->vs.export_prim_id = options->key.vs.export_prim_id;
break;
default:
break;
void
radv_compile_nir_shader(struct ac_llvm_compiler *ac_llvm,
- struct ac_shader_binary *binary,
- struct ac_shader_config *config,
+ struct radv_shader_binary **rbinary,
struct radv_shader_variant_info *shader_info,
struct nir_shader *const *nir,
int nir_count,
llvm_module = ac_translate_nir_to_llvm(ac_llvm, nir, nir_count, shader_info,
options);
- ac_compile_llvm_module(ac_llvm, llvm_module, binary, config, shader_info,
- nir[0]->info.stage, options);
+ ac_compile_llvm_module(ac_llvm, llvm_module, rbinary, shader_info,
+ nir[nir_count - 1]->info.stage, options);
for (int i = 0; i < nir_count; ++i)
ac_fill_shader_info(shader_info, nir[i], options);
ctx->gsvs_ring[0],
1, ctx->ac.i32_0,
vtx_offset, soffset,
- 0, 1, 1, true, false);
+ 0, ac_glc | ac_slc, true, false);
LLVMTypeRef type = LLVMGetAllocatedType(ctx->abi.outputs[ac_llvm_reg_index_soa(i, j)]);
if (ac_get_type_size(type) == 2) {
radv_emit_streamout(ctx, stream);
if (stream == 0) {
- handle_vs_outputs_post(ctx, false, false,
+ handle_vs_outputs_post(ctx, false, false, true,
&ctx->shader_info->vs.outinfo);
}
void
radv_compile_gs_copy_shader(struct ac_llvm_compiler *ac_llvm,
struct nir_shader *geom_shader,
- struct ac_shader_binary *binary,
- struct ac_shader_config *config,
+ struct radv_shader_binary **rbinary,
struct radv_shader_variant_info *shader_info,
const struct radv_nir_compiler_options *options)
{
ac_llvm_finalize_module(&ctx, ac_llvm->passmgr, options);
- ac_compile_llvm_module(ac_llvm, ctx.ac.module, binary, config, shader_info,
+ ac_compile_llvm_module(ac_llvm, ctx.ac.module, rbinary, shader_info,
MESA_SHADER_VERTEX, options);
+ (*rbinary)->is_gs_copy_shader = true;
+
}
projects / mesa.git / blobdiff