From c9ab7d7525d653f50a764c4ba6bbec0a98e1d92a Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Thu, 29 Sep 2022 03:00:38 +0200 Subject: [PATCH] radeonsi: use default float mode for CL Signed-off-by: Karol Herbst Reviewed-by: Pierre-Eric Pelloux-Prayer Part-of: --- src/gallium/drivers/radeonsi/si_shader.c | 3 ++- src/gallium/drivers/radeonsi/si_shader_internal.h | 3 ++- src/gallium/drivers/radeonsi/si_shader_llvm.c | 9 ++++++--- 3 files changed, 10 insertions(+), 5 deletions(-) diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index c4f933d5b76..61de41fb42c 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -2461,7 +2461,8 @@ si_get_shader_part(struct si_screen *sscreen, struct si_shader_part **list, } struct si_shader_context ctx; - si_llvm_context_init(&ctx, sscreen, compiler, wave32 ? 32 : 64, exports_color_null, exports_mrtz); + si_llvm_context_init(&ctx, sscreen, compiler, wave32 ? 32 : 64, exports_color_null, exports_mrtz, + AC_FLOAT_MODE_DEFAULT_OPENGL); ctx.shader = &shader; ctx.stage = stage; diff --git a/src/gallium/drivers/radeonsi/si_shader_internal.h b/src/gallium/drivers/radeonsi/si_shader_internal.h index c03dad739e9..de1bc793dc2 100644 --- a/src/gallium/drivers/radeonsi/si_shader_internal.h +++ b/src/gallium/drivers/radeonsi/si_shader_internal.h @@ -182,7 +182,8 @@ bool si_compile_llvm(struct si_screen *sscreen, struct si_shader_binary *binary, gl_shader_stage stage, const char *name, bool less_optimized); void si_llvm_context_init(struct si_shader_context *ctx, struct si_screen *sscreen, struct ac_llvm_compiler *compiler, unsigned wave_size, - bool exports_color_null, bool exports_mrtz); + bool exports_color_null, bool exports_mrtz, + enum ac_float_mode float_mode); void si_llvm_create_func(struct si_shader_context *ctx, const char *name, LLVMTypeRef *return_types, unsigned num_return_elems, unsigned max_workgroup_size); void si_llvm_create_main_func(struct si_shader_context *ctx); diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c index a3c501ad4c5..68b3fd09edb 100644 --- a/src/gallium/drivers/radeonsi/si_shader_llvm.c +++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c @@ -126,14 +126,15 @@ bool si_compile_llvm(struct si_screen *sscreen, struct si_shader_binary *binary, void si_llvm_context_init(struct si_shader_context *ctx, struct si_screen *sscreen, struct ac_llvm_compiler *compiler, unsigned wave_size, - bool exports_color_null, bool exports_mrtz) + bool exports_color_null, bool exports_mrtz, + enum ac_float_mode float_mode) { memset(ctx, 0, sizeof(*ctx)); ctx->screen = sscreen; ctx->compiler = compiler; ac_llvm_context_init(&ctx->ac, compiler, sscreen->info.gfx_level, sscreen->info.family, - sscreen->info.has_3d_cube_border_color_mipmap, AC_FLOAT_MODE_DEFAULT_OPENGL, + sscreen->info.has_3d_cube_border_color_mipmap, float_mode, wave_size, 64, exports_color_null, exports_mrtz); } @@ -1091,13 +1092,15 @@ bool si_llvm_compile_shader(struct si_screen *sscreen, struct ac_llvm_compiler * { struct si_shader_selector *sel = shader->selector; struct si_shader_context ctx; + enum ac_float_mode float_mode = nir->info.stage == MESA_SHADER_KERNEL ? AC_FLOAT_MODE_DEFAULT : AC_FLOAT_MODE_DEFAULT_OPENGL; bool exports_color_null = sel->info.colors_written; bool exports_mrtz = sel->info.writes_z || sel->info.writes_stencil || sel->info.writes_samplemask; if (!exports_mrtz && !exports_color_null) exports_color_null = si_shader_uses_discard(shader) || sscreen->info.gfx_level < GFX10; - si_llvm_context_init(&ctx, sscreen, compiler, shader->wave_size, exports_color_null, exports_mrtz); + si_llvm_context_init(&ctx, sscreen, compiler, shader->wave_size, exports_color_null, exports_mrtz, + float_mode); ctx.args = args; if (!si_llvm_translate_nir(&ctx, shader, nir, false)) {