From f0b2944a0bf63abde981e742984ff10a0c034c70 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sun, 7 May 2023 01:44:35 +0200 Subject: [PATCH] gallium: add PIPE_COMPUTE_CAP_MAX_SUBGROUPS This is required for enabling OpenCL subgroups. Signed-off-by: Karol Herbst Reviewed-by: Alyssa Rosenzweig Part-of: --- docs/gallium/screen.rst | 4 ++++ src/gallium/drivers/asahi/agx_pipe.c | 3 +++ src/gallium/drivers/crocus/crocus_screen.c | 1 + src/gallium/drivers/freedreno/freedreno_screen.c | 3 +++ src/gallium/drivers/iris/iris_screen.c | 3 +++ src/gallium/drivers/llvmpipe/lp_screen.c | 6 ++++++ src/gallium/drivers/nouveau/nv50/nv50_screen.c | 2 ++ src/gallium/drivers/nouveau/nvc0/nvc0_screen.c | 2 ++ src/gallium/drivers/panfrost/pan_screen.c | 3 +++ src/gallium/drivers/r600/r600_pipe_common.c | 1 + src/gallium/drivers/radeonsi/si_get.c | 15 +++++++++++++++ src/gallium/drivers/softpipe/sp_screen.c | 1 + src/gallium/drivers/v3d/v3d_screen.c | 3 +++ src/gallium/drivers/zink/zink_screen.c | 1 + src/gallium/include/pipe/p_defines.h | 1 + 15 files changed, 49 insertions(+) diff --git a/docs/gallium/screen.rst b/docs/gallium/screen.rst index 087b29e9f44..7e106531917 100644 --- a/docs/gallium/screen.rst +++ b/docs/gallium/screen.rst @@ -798,6 +798,10 @@ pipe_screen::get_compute_param. clock in MHz. Value type: ``uint32_t`` * ``PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS``: Maximum number of compute units Value type: ``uint32_t`` +* ``PIPE_COMPUTE_CAP_MAX_SUBGROUPS``: The max amount of subgroups there can be + inside a block. Non 0 indicates support for OpenCL subgroups including + implementing ``get_compute_state_subgroup_size`` if multiple subgroup sizes + are supported. * ``PIPE_COMPUTE_CAP_IMAGES_SUPPORTED``: Whether images are supported non-zero means yes, zero means no. Value type: ``uint32_t`` * ``PIPE_COMPUTE_CAP_SUBGROUP_SIZE``: The size of a basic execution unit in diff --git a/src/gallium/drivers/asahi/agx_pipe.c b/src/gallium/drivers/asahi/agx_pipe.c index b0a2547f17b..587d07208fc 100644 --- a/src/gallium/drivers/asahi/agx_pipe.c +++ b/src/gallium/drivers/asahi/agx_pipe.c @@ -1858,6 +1858,9 @@ agx_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type, case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t[]){32}); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t[]){0 /* TODO */}); + case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK: RET((uint64_t[]){1024}); // TODO } diff --git a/src/gallium/drivers/crocus/crocus_screen.c b/src/gallium/drivers/crocus/crocus_screen.c index cf8ac7064bd..a40914dd721 100644 --- a/src/gallium/drivers/crocus/crocus_screen.c +++ b/src/gallium/drivers/crocus/crocus_screen.c @@ -601,6 +601,7 @@ crocus_get_compute_param(struct pipe_screen *pscreen, case PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE: case PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE: case PIPE_COMPUTE_CAP_MAX_INPUT_SIZE: + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: // XXX: I think these are for Clover... return 0; diff --git a/src/gallium/drivers/freedreno/freedreno_screen.c b/src/gallium/drivers/freedreno/freedreno_screen.c index 2b5481693cd..78bc3d4b027 100644 --- a/src/gallium/drivers/freedreno/freedreno_screen.c +++ b/src/gallium/drivers/freedreno/freedreno_screen.c @@ -846,6 +846,9 @@ fd_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type, case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t[]){32}); // TODO + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t[]){0}); // TODO + case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK: RET((uint64_t[]){ compiler->max_variable_workgroup_size }); } diff --git a/src/gallium/drivers/iris/iris_screen.c b/src/gallium/drivers/iris/iris_screen.c index 9b588c7cb1a..c418c4dba94 100644 --- a/src/gallium/drivers/iris/iris_screen.c +++ b/src/gallium/drivers/iris/iris_screen.c @@ -620,6 +620,9 @@ iris_get_compute_param(struct pipe_screen *pscreen, case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t []) { BRW_SUBGROUP_SIZE }); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t []) { devinfo->max_cs_workgroup_threads }); + case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE: case PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE: RET((uint64_t []) { 1 << 30 }); /* TODO */ diff --git a/src/gallium/drivers/llvmpipe/lp_screen.c b/src/gallium/drivers/llvmpipe/lp_screen.c index c95563d4548..f4d356be1dd 100644 --- a/src/gallium/drivers/llvmpipe/lp_screen.c +++ b/src/gallium/drivers/llvmpipe/lp_screen.c @@ -537,6 +537,12 @@ llvmpipe_get_compute_param(struct pipe_screen *_screen, *subgroup_size = lp_native_vector_width / 32; } return sizeof(uint32_t); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + if (ret) { + uint32_t *subgroup_size = ret; + *subgroup_size = 1024 / (lp_native_vector_width / 32); + } + return sizeof(uint32_t); case PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS: if (ret) { uint32_t *max_compute_units = ret; diff --git a/src/gallium/drivers/nouveau/nv50/nv50_screen.c b/src/gallium/drivers/nouveau/nv50/nv50_screen.c index 8083f2a5b0e..38ef520bb2d 100644 --- a/src/gallium/drivers/nouveau/nv50/nv50_screen.c +++ b/src/gallium/drivers/nouveau/nv50/nv50_screen.c @@ -451,6 +451,8 @@ nv50_screen_get_compute_param(struct pipe_screen *pscreen, RET((uint64_t []) { 4096 }); case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t []) { 32 }); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t []) { 0 }); case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE: RET((uint64_t []) { nouveau_device_get_global_mem_size(dev) }); case PIPE_COMPUTE_CAP_IMAGES_SUPPORTED: diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c b/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c index 746a9df7ec2..7bd9928f637 100644 --- a/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c +++ b/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c @@ -558,6 +558,8 @@ nvc0_screen_get_compute_param(struct pipe_screen *pscreen, RET((uint64_t []) { 4096 }); case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t []) { 32 }); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t []) { 0 }); case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE: RET((uint64_t []) { nouveau_device_get_global_mem_size(dev) }); case PIPE_COMPUTE_CAP_IMAGES_SUPPORTED: diff --git a/src/gallium/drivers/panfrost/pan_screen.c b/src/gallium/drivers/panfrost/pan_screen.c index c0bff78d4ae..01d32f21f1a 100644 --- a/src/gallium/drivers/panfrost/pan_screen.c +++ b/src/gallium/drivers/panfrost/pan_screen.c @@ -754,6 +754,9 @@ panfrost_get_compute_param(struct pipe_screen *pscreen, case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t[]){pan_subgroup_size(dev->arch)}); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t[]){0 /* TODO */}); + case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK: RET((uint64_t[]){1024}); // TODO } diff --git a/src/gallium/drivers/r600/r600_pipe_common.c b/src/gallium/drivers/r600/r600_pipe_common.c index c47e1dfe551..01f9d7fbec2 100644 --- a/src/gallium/drivers/r600/r600_pipe_common.c +++ b/src/gallium/drivers/r600/r600_pipe_common.c @@ -1033,6 +1033,7 @@ static int r600_get_compute_param(struct pipe_screen *screen, } return sizeof(uint32_t); case PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE: + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: break; /* unused */ case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: if (ret) { diff --git a/src/gallium/drivers/radeonsi/si_get.c b/src/gallium/drivers/radeonsi/si_get.c index f9ce8caee97..643fd8d06a7 100644 --- a/src/gallium/drivers/radeonsi/si_get.c +++ b/src/gallium/drivers/radeonsi/si_get.c @@ -1102,6 +1102,21 @@ static int si_get_compute_param(struct pipe_screen *screen, enum pipe_shader_ir return sizeof(uint32_t); case PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE: break; /* unused */ + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: { + if (ret) { + uint32_t *max_subgroups = ret; + unsigned threads = get_max_threads_per_block(sscreen, ir_type); + unsigned subgroup_size; + + if (sscreen->debug_flags & DBG(W64_CS) || sscreen->info.gfx_level < GFX10) + subgroup_size = 64; + else + subgroup_size = 32; + + *max_subgroups = threads / subgroup_size; + } + return sizeof(uint32_t); + } case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: if (ret) { uint32_t *subgroup_size = ret; diff --git a/src/gallium/drivers/softpipe/sp_screen.c b/src/gallium/drivers/softpipe/sp_screen.c index 3b6b7175e40..77bd2e0498a 100644 --- a/src/gallium/drivers/softpipe/sp_screen.c +++ b/src/gallium/drivers/softpipe/sp_screen.c @@ -557,6 +557,7 @@ softpipe_get_compute_param(struct pipe_screen *_screen, case PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS: case PIPE_COMPUTE_CAP_IMAGES_SUPPORTED: case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: case PIPE_COMPUTE_CAP_ADDRESS_BITS: case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK: break; diff --git a/src/gallium/drivers/v3d/v3d_screen.c b/src/gallium/drivers/v3d/v3d_screen.c index bce1eeafcd9..57111eb5f66 100644 --- a/src/gallium/drivers/v3d/v3d_screen.c +++ b/src/gallium/drivers/v3d/v3d_screen.c @@ -548,6 +548,9 @@ v3d_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type, case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t []) { 16 }); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: + RET((uint32_t []) { 0 }); + } return 0; diff --git a/src/gallium/drivers/zink/zink_screen.c b/src/gallium/drivers/zink/zink_screen.c index ff6867f09c5..a67dcd9c5d1 100644 --- a/src/gallium/drivers/zink/zink_screen.c +++ b/src/gallium/drivers/zink/zink_screen.c @@ -440,6 +440,7 @@ zink_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type, case PIPE_COMPUTE_CAP_SUBGROUP_SIZE: RET((uint32_t []) { screen->info.props11.subgroupSize }); + case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE: case PIPE_COMPUTE_CAP_MAX_CLOCK_FREQUENCY: case PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS: diff --git a/src/gallium/include/pipe/p_defines.h b/src/gallium/include/pipe/p_defines.h index a96b04072b1..6b5cfc10f79 100644 --- a/src/gallium/include/pipe/p_defines.h +++ b/src/gallium/include/pipe/p_defines.h @@ -1139,6 +1139,7 @@ enum pipe_compute_cap PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE, PIPE_COMPUTE_CAP_MAX_CLOCK_FREQUENCY, PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS, + PIPE_COMPUTE_CAP_MAX_SUBGROUPS, PIPE_COMPUTE_CAP_IMAGES_SUPPORTED, PIPE_COMPUTE_CAP_SUBGROUP_SIZE, PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK,