#include "si_compute.h"
#include "sid.h"
#include "sid_tables.h"
+#include "tgsi/tgsi_from_mesa.h"
#include "driver_ddebug/dd_util.h"
#include "util/u_dump.h"
#include "util/u_log.h"
}
static void si_dump_shader(struct si_screen *sscreen,
- enum pipe_shader_type processor,
struct si_shader *shader, FILE *f)
{
if (shader->shader_log)
fwrite(shader->shader_log, shader->shader_log_size, 1, f);
else
- si_shader_dump(sscreen, shader, NULL, processor, f, false);
+ si_shader_dump(sscreen, shader, NULL, f, false);
if (shader->bo && sscreen->options.dump_shader_binary) {
unsigned size = shader->bo->b.b.width0;
*/
struct si_context *ctx;
struct si_shader *shader;
- enum pipe_shader_type processor;
/* For keep-alive reference counts */
struct si_shader_selector *sel;
{
struct si_log_chunk_shader *chunk = data;
struct si_screen *sscreen = chunk->ctx->screen;
- si_dump_shader(sscreen, chunk->processor,
- chunk->shader, f);
+ si_dump_shader(sscreen, chunk->shader, f);
}
static struct u_log_chunk_type si_log_chunk_type_shader = {
struct si_log_chunk_shader *chunk = CALLOC_STRUCT(si_log_chunk_shader);
chunk->ctx = ctx;
- chunk->processor = state->cso->info.processor;
chunk->shader = current;
si_shader_selector_reference(ctx, &chunk->sel, current->selector);
u_log_chunk(log, &si_log_chunk_type_shader, chunk);
struct si_log_chunk_shader *chunk = CALLOC_STRUCT(si_log_chunk_shader);
chunk->ctx = ctx;
- chunk->processor = PIPE_SHADER_COMPUTE;
chunk->shader = &state->program->shader;
si_compute_reference(&chunk->program, state->program);
u_log_chunk(log, &si_log_chunk_type_shader, chunk);
struct si_shader_binary *binary,
uint64_t *addr,
unsigned *num,
- struct si_shader_inst *instructions)
+ struct si_shader_inst *instructions,
+ enum pipe_shader_type shader_type)
{
if (!ac_rtld_open(rtld_binary, (struct ac_rtld_open_info){
.info = &screen->info,
+ .shader_type = tgsi_processor_to_shader_stage(shader_type),
.num_parts = 1,
.elf_ptrs = &binary->elf_buffer,
.elf_sizes = &binary->elf_size }))
return;
struct si_screen *screen = shader->selector->screen;
+ enum pipe_shader_type shader_type = shader->selector->type;
uint64_t start_addr = shader->bo->gpu_address;
uint64_t end_addr = start_addr + shader->bo->b.b.width0;
unsigned i;
if (shader->prolog) {
si_add_split_disasm(screen, &rtld_binaries[0], &shader->prolog->binary,
- &inst_addr, &num_inst, instructions);
+ &inst_addr, &num_inst, instructions, shader_type);
}
if (shader->previous_stage) {
si_add_split_disasm(screen, &rtld_binaries[1], &shader->previous_stage->binary,
- &inst_addr, &num_inst, instructions);
+ &inst_addr, &num_inst, instructions, shader_type);
}
if (shader->prolog2) {
si_add_split_disasm(screen, &rtld_binaries[2], &shader->prolog2->binary,
- &inst_addr, &num_inst, instructions);
+ &inst_addr, &num_inst, instructions, shader_type);
}
si_add_split_disasm(screen, &rtld_binaries[3], &shader->binary,
- &inst_addr, &num_inst, instructions);
+ &inst_addr, &num_inst, instructions, shader_type);
if (shader->epilog) {
si_add_split_disasm(screen, &rtld_binaries[4], &shader->epilog->binary,
- &inst_addr, &num_inst, instructions);
+ &inst_addr, &num_inst, instructions, shader_type);
}
fprintf(f, COLOR_YELLOW "%s - annotated disassembly:" COLOR_RESET "\n",
- si_get_shader_name(shader, shader->selector->type));
+ si_get_shader_name(shader));
/* Print instructions with annotations. */
for (i = 0; i < num_inst; i++) {
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data);
-static void si_dump_shader_key(unsigned processor, const struct si_shader *shader,
- FILE *f);
+static void si_dump_shader_key(const struct si_shader *shader, FILE *f);
static void si_build_vs_prolog_function(struct si_shader_context *ctx,
union si_shader_part_key *key);
static void si_shader_dump_disassembly(struct si_screen *screen,
const struct si_shader_binary *binary,
+ enum pipe_shader_type shader_type,
struct pipe_debug_callback *debug,
const char *name, FILE *file)
{
if (!ac_rtld_open(&rtld_binary, (struct ac_rtld_open_info){
.info = &screen->info,
+ .shader_type = tgsi_processor_to_shader_stage(shader_type),
.num_parts = 1,
.elf_ptrs = &binary->elf_buffer,
.elf_sizes = &binary->elf_size }))
DIV_ROUND_UP(max_workgroup_size, 64);
}
break;
+ default:;
}
/* Compute the per-SIMD wave counts. */
const struct ac_shader_config *conf = &shader->config;
if (screen->options.debug_disassembly)
- si_shader_dump_disassembly(screen, &shader->binary, debug, "main", NULL);
+ si_shader_dump_disassembly(screen, &shader->binary,
+ shader->selector->type,
+ debug, "main", NULL);
pipe_debug_message(debug, SHADER_INFO,
"Shader Stats: SGPRS: %d VGPRS: %d Code Size: %d "
static void si_shader_dump_stats(struct si_screen *sscreen,
struct si_shader *shader,
- unsigned processor,
FILE *file,
bool check_debug_option)
{
const struct ac_shader_config *conf = &shader->config;
+ enum pipe_shader_type shader_type =
+ shader->selector ? shader->selector->type : PIPE_SHADER_COMPUTE;
if (!check_debug_option ||
- si_can_dump_shader(sscreen, processor)) {
- if (processor == PIPE_SHADER_FRAGMENT) {
+ si_can_dump_shader(sscreen, shader_type)) {
+ if (shader_type == PIPE_SHADER_FRAGMENT) {
fprintf(file, "*** SHADER CONFIG ***\n"
"SPI_PS_INPUT_ADDR = 0x%04x\n"
"SPI_PS_INPUT_ENA = 0x%04x\n",
}
}
-const char *si_get_shader_name(const struct si_shader *shader, unsigned processor)
+const char *si_get_shader_name(const struct si_shader *shader)
{
- switch (processor) {
+ enum pipe_shader_type shader_type =
+ shader->selector ? shader->selector->type : PIPE_SHADER_COMPUTE;
+
+ switch (shader_type) {
case PIPE_SHADER_VERTEX:
if (shader->key.as_es)
return "Vertex Shader as ES";
}
void si_shader_dump(struct si_screen *sscreen, struct si_shader *shader,
- struct pipe_debug_callback *debug, unsigned processor,
+ struct pipe_debug_callback *debug,
FILE *file, bool check_debug_option)
{
+ enum pipe_shader_type shader_type =
+ shader->selector ? shader->selector->type : PIPE_SHADER_COMPUTE;
+
if (!check_debug_option ||
- si_can_dump_shader(sscreen, processor))
- si_dump_shader_key(processor, shader, file);
+ si_can_dump_shader(sscreen, shader_type))
+ si_dump_shader_key(shader, file);
if (!check_debug_option && shader->binary.llvm_ir_string) {
if (shader->previous_stage &&
shader->previous_stage->binary.llvm_ir_string) {
fprintf(file, "\n%s - previous stage - LLVM IR:\n\n",
- si_get_shader_name(shader, processor));
+ si_get_shader_name(shader));
fprintf(file, "%s\n", shader->previous_stage->binary.llvm_ir_string);
}
fprintf(file, "\n%s - main shader part - LLVM IR:\n\n",
- si_get_shader_name(shader, processor));
+ si_get_shader_name(shader));
fprintf(file, "%s\n", shader->binary.llvm_ir_string);
}
if (!check_debug_option ||
- (si_can_dump_shader(sscreen, processor) &&
+ (si_can_dump_shader(sscreen, shader_type) &&
!(sscreen->debug_flags & DBG(NO_ASM)))) {
- fprintf(file, "\n%s:\n", si_get_shader_name(shader, processor));
+ fprintf(file, "\n%s:\n", si_get_shader_name(shader));
if (shader->prolog)
si_shader_dump_disassembly(sscreen, &shader->prolog->binary,
- debug, "prolog", file);
+ shader_type, debug, "prolog", file);
if (shader->previous_stage)
si_shader_dump_disassembly(sscreen, &shader->previous_stage->binary,
- debug, "previous stage", file);
+ shader_type, debug, "previous stage", file);
if (shader->prolog2)
si_shader_dump_disassembly(sscreen, &shader->prolog2->binary,
- debug, "prolog2", file);
+ shader_type, debug, "prolog2", file);
- si_shader_dump_disassembly(sscreen, &shader->binary, debug, "main", file);
+ si_shader_dump_disassembly(sscreen, &shader->binary, shader_type, debug, "main", file);
if (shader->epilog)
si_shader_dump_disassembly(sscreen, &shader->epilog->binary,
- debug, "epilog", file);
+ shader_type, debug, "epilog", file);
fprintf(file, "\n");
}
- si_shader_dump_stats(sscreen, shader, processor, file,
- check_debug_option);
+ si_shader_dump_stats(sscreen, shader, file, check_debug_option);
}
static int si_compile_llvm(struct si_screen *sscreen,
struct ac_llvm_compiler *compiler,
LLVMModuleRef mod,
struct pipe_debug_callback *debug,
- unsigned processor,
+ enum pipe_shader_type shader_type,
const char *name,
bool less_optimized)
{
unsigned count = p_atomic_inc_return(&sscreen->num_compilations);
- if (si_can_dump_shader(sscreen, processor)) {
+ if (si_can_dump_shader(sscreen, shader_type)) {
fprintf(stderr, "radeonsi: Compiling shader %d\n", count);
if (!(sscreen->debug_flags & (DBG(NO_IR) | DBG(PREOPT_IR)))) {
struct ac_rtld_binary rtld;
if (!ac_rtld_open(&rtld, (struct ac_rtld_open_info){
.info = &sscreen->info,
+ .shader_type = tgsi_processor_to_shader_stage(shader_type),
.num_parts = 1,
.elf_ptrs = &binary->elf_buffer,
.elf_sizes = &binary->elf_size }))
"GS Copy Shader", false) == 0) {
if (si_can_dump_shader(sscreen, PIPE_SHADER_GEOMETRY))
fprintf(stderr, "GS Copy Shader:\n");
- si_shader_dump(sscreen, ctx.shader, debug,
- PIPE_SHADER_GEOMETRY, stderr, true);
+ si_shader_dump(sscreen, ctx.shader, debug, stderr, true);
if (!ctx.shader->config.scratch_bytes_per_wave)
ok = si_shader_binary_upload(sscreen, ctx.shader, 0);
fprintf(f, "}\n");
}
-static void si_dump_shader_key(unsigned processor, const struct si_shader *shader,
- FILE *f)
+static void si_dump_shader_key(const struct si_shader *shader, FILE *f)
{
const struct si_shader_key *key = &shader->key;
+ enum pipe_shader_type shader_type =
+ shader->selector ? shader->selector->type : PIPE_SHADER_COMPUTE;
fprintf(f, "SHADER KEY\n");
- switch (processor) {
+ switch (shader_type) {
case PIPE_SHADER_VERTEX:
si_dump_shader_key_vs(key, &key->part.vs.prolog,
"part.vs.prolog", f);
assert(0);
}
- if ((processor == PIPE_SHADER_GEOMETRY ||
- processor == PIPE_SHADER_TESS_EVAL ||
- processor == PIPE_SHADER_VERTEX) &&
+ if ((shader_type == PIPE_SHADER_GEOMETRY ||
+ shader_type == PIPE_SHADER_TESS_EVAL ||
+ shader_type == PIPE_SHADER_VERTEX) &&
!key->as_es && !key->as_ls) {
fprintf(f, " opt.kill_outputs = 0x%"PRIx64"\n", key->opt.kill_outputs);
fprintf(f, " opt.clip_disable = %u\n", key->opt.clip_disable);
/* Dump TGSI code before doing TGSI->LLVM conversion in case the
* conversion fails. */
- if (si_can_dump_shader(sscreen, sel->info.processor) &&
+ if (si_can_dump_shader(sscreen, sel->type) &&
!(sscreen->debug_flags & DBG(NO_TGSI))) {
if (sel->tokens)
tgsi_dump(sel->tokens, 0);
/* Compile to bytecode. */
r = si_compile_llvm(sscreen, &shader->binary, &shader->config, compiler,
ctx.ac.module, debug, ctx.type,
- si_get_shader_name(shader, ctx.type),
+ si_get_shader_name(shader),
si_should_optimize_less(compiler, shader->selector));
si_llvm_dispose(&ctx);
if (r) {
shader->config.num_vgprs = MAX2(shader->config.num_vgprs,
shader->info.num_input_vgprs);
break;
+ default:;
}
/* Update SGPR and VGPR counts. */
}
si_fix_resource_usage(sscreen, shader);
- si_shader_dump(sscreen, shader, debug, sel->info.processor,
- stderr, true);
+ si_shader_dump(sscreen, shader, debug, stderr, true);
/* Upload. */
if (!si_shader_binary_upload(sscreen, shader, 0)) {