diff options
| author | Marek Olšák <[email protected]> | 2019-07-12 17:22:30 -0400 |
|---|---|---|
| committer | Marek Olšák <[email protected]> | 2019-07-19 20:16:39 -0400 |
| commit | e2c8ff009e5e077ae93ad19c63bd06d5c23d8376 (patch) | |
| tree | 2f0a1682bf1df56bdca4c937106fbb3dd9d7fa71 /src/gallium/drivers | |
| parent | a8a526c5cb8a160ef4b4a9db38359247986fe692 (diff) | |
radeonsi: set threadgroup size to 0 for threadgroups with only 1 wave
This has no effect on Wave64.
Reviewed-by: Pierre-Eric Pelloux-Prayer <[email protected]>
Acked-by: Samuel Pitoiset <[email protected]>
Diffstat (limited to 'src/gallium/drivers')
| -rw-r--r-- | src/gallium/drivers/radeonsi/si_shader.c | 6 |
1 files changed, 3 insertions, 3 deletions
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 3cfb1207b4a..8d06859f1f2 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -4484,10 +4484,10 @@ static unsigned si_get_max_workgroup_size(const struct si_shader *shader) case PIPE_SHADER_TESS_CTRL: /* Return this so that LLVM doesn't remove s_barrier * instructions on chips where we use s_barrier. */ - return shader->selector->screen->info.chip_class >= GFX7 ? 128 : 64; + return shader->selector->screen->info.chip_class >= GFX7 ? 128 : 0; case PIPE_SHADER_GEOMETRY: - return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 64; + return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 0; case PIPE_SHADER_COMPUTE: break; /* see below */ @@ -7626,7 +7626,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, /* Create the function. */ si_create_function(ctx, "tcs_epilog", NULL, 0, &fninfo, - ctx->screen->info.chip_class >= GFX7 ? 128 : 64); + ctx->screen->info.chip_class >= GFX7 ? 128 : 0); ac_declare_lds_as_pointer(&ctx->ac); func = ctx->main_fn; |