[PATCH 1/2] gallium: Add dimension parameter to launch_grid
Jan Vesely
jan.vesely at rutgers.edu
Tue Aug 5 16:30:03 PDT 2014
On Tue, 2014-08-05 at 13:45 +0300, Francisco Jerez wrote:
> Jan Vesely <jan.vesely at rutgers.edu> writes:
>
> > This is needed for OpenCL
> >
> > CC: Francisco Jerez <currojerez at riseup.net>
> > CC: Tom Stellard <tom at stellard.net>
> >
> > Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> > ---
> >
> > I tried to find another way how to get the information, but in the end
> > there is no way to distinguish between <1,1,1> dim 1, and <1,1,1> dim 2 (or 3).
> > So passing the work_dim information is required.
> >
> > this series fixes piglits cl-program-max-work-item-sizes, and get-work-dim.cl on my TURKS gpu.
> >
>
> My plan was to keep the pipe driver interface as simple as possible and
> pass the dimension as an additional *kernel* parameter directly from the
> clover front-end (along with other things that we currently don't handle
> and we could abstract from the pipe driver, like the grid offset).
I'm currently looking into Tom's suggestion to pass these values after
the kernel args, appending workdim in clover should not be a big
problem.
> I don't see any reason to bother the driver with this as e.g. <1> and
> <1,1,1> are equivalent grid definitions that will invariably result in
> the same hardware setup.
My original idea was to get the last useful ( > 1) dimension, or return
1. However, the specs say that it should return the value passed to
clEnqueueNDRangeKernel, and that's what the existing piglits tests.
jan
>
> P.S.: Sorry for taking so long to get back to you, been quite busy
> during the last couple of weeks.
no problem, there's no rush. I need to prioritize university work these
weeks anyway
>
> > src/gallium/drivers/ilo/ilo_gpgpu.c | 2 +-
> > src/gallium/drivers/nouveau/nvc0/nvc0_compute.c | 3 +-
> > src/gallium/drivers/nouveau/nvc0/nvc0_context.h | 4 +--
> > src/gallium/drivers/nouveau/nvc0/nve4_compute.c | 3 +-
> > src/gallium/drivers/r600/evergreen_compute.c | 2 +-
> > src/gallium/drivers/radeonsi/si_compute.c | 2 +-
> > src/gallium/include/pipe/p_context.h | 2 +-
> > src/gallium/state_trackers/clover/core/kernel.cpp | 3 +-
> > src/gallium/tests/trivial/compute.c | 38 +++++++++++------------
> > 9 files changed, 31 insertions(+), 28 deletions(-)
> >
> > diff --git a/src/gallium/drivers/ilo/ilo_gpgpu.c b/src/gallium/drivers/ilo/ilo_gpgpu.c
> > index b17a518..44e5c59 100644
> > --- a/src/gallium/drivers/ilo/ilo_gpgpu.c
> > +++ b/src/gallium/drivers/ilo/ilo_gpgpu.c
> > @@ -35,7 +35,7 @@
> > static void
> > ilo_launch_grid(struct pipe_context *pipe,
> > const uint *block_layout, const uint *grid_layout,
> > - uint32_t pc, const void *input)
> > + uint32_t pc, const void *input, uint dimensions)
> > {
> > }
> >
> > diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_compute.c b/src/gallium/drivers/nouveau/nvc0/nvc0_compute.c
> > index ad287a2..2340afe 100644
> > --- a/src/gallium/drivers/nouveau/nvc0/nvc0_compute.c
> > +++ b/src/gallium/drivers/nouveau/nvc0/nvc0_compute.c
> > @@ -197,7 +197,8 @@ void
> > nvc0_launch_grid(struct pipe_context *pipe,
> > const uint *block_layout, const uint *grid_layout,
> > uint32_t label,
> > - const void *input)
> > + const void *input,
> > + uint dimensions)
> > {
> > struct nvc0_context *nvc0 = nvc0_context(pipe);
> > struct nouveau_pushbuf *push = nvc0->base.pushbuf;
> > diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_context.h b/src/gallium/drivers/nouveau/nvc0/nvc0_context.h
> > index ebeb8c4..30b5311 100644
> > --- a/src/gallium/drivers/nouveau/nvc0/nvc0_context.h
> > +++ b/src/gallium/drivers/nouveau/nvc0/nvc0_context.h
> > @@ -353,10 +353,10 @@ void nvc0_push_vbo(struct nvc0_context *, const struct pipe_draw_info *);
> >
> > /* nve4_compute.c */
> > void nve4_launch_grid(struct pipe_context *,
> > - const uint *, const uint *, uint32_t, const void *);
> > + const uint *, const uint *, uint32_t, const void *, uint);
> >
> > /* nvc0_compute.c */
> > void nvc0_launch_grid(struct pipe_context *,
> > - const uint *, const uint *, uint32_t, const void *);
> > + const uint *, const uint *, uint32_t, const void *, uint);
> >
> > #endif
> > diff --git a/src/gallium/drivers/nouveau/nvc0/nve4_compute.c b/src/gallium/drivers/nouveau/nvc0/nve4_compute.c
> > index f243316..d1c7b08 100644
> > --- a/src/gallium/drivers/nouveau/nvc0/nve4_compute.c
> > +++ b/src/gallium/drivers/nouveau/nvc0/nve4_compute.c
> > @@ -432,7 +432,8 @@ void
> > nve4_launch_grid(struct pipe_context *pipe,
> > const uint *block_layout, const uint *grid_layout,
> > uint32_t label,
> > - const void *input)
> > + const void *input,
> > + uint dimensions)
> > {
> > struct nvc0_context *nvc0 = nvc0_context(pipe);
> > struct nouveau_pushbuf *push = nvc0->base.pushbuf;
> > diff --git a/src/gallium/drivers/r600/evergreen_compute.c b/src/gallium/drivers/r600/evergreen_compute.c
> > index 1970414..3928676 100644
> > --- a/src/gallium/drivers/r600/evergreen_compute.c
> > +++ b/src/gallium/drivers/r600/evergreen_compute.c
> > @@ -541,7 +541,7 @@ void evergreen_emit_cs_shader(
> > static void evergreen_launch_grid(
> > struct pipe_context *ctx_,
> > const uint *block_layout, const uint *grid_layout,
> > - uint32_t pc, const void *input)
> > + uint32_t pc, const void *input, uint dimensions)
> > {
> > struct r600_context *ctx = (struct r600_context *)ctx_;
> >
> > diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c
> > index 42e4fec..3e044b5 100644
> > --- a/src/gallium/drivers/radeonsi/si_compute.c
> > +++ b/src/gallium/drivers/radeonsi/si_compute.c
> > @@ -162,7 +162,7 @@ static unsigned compute_num_waves_for_scratch(
> > static void si_launch_grid(
> > struct pipe_context *ctx,
> > const uint *block_layout, const uint *grid_layout,
> > - uint32_t pc, const void *input)
> > + uint32_t pc, const void *input, uint dimensions)
> > {
> > struct si_context *sctx = (struct si_context*)ctx;
> > struct si_pipe_compute *program = sctx->cs_shader_state.program;
> > diff --git a/src/gallium/include/pipe/p_context.h b/src/gallium/include/pipe/p_context.h
> > index af5674f..70212aa 100644
> > --- a/src/gallium/include/pipe/p_context.h
> > +++ b/src/gallium/include/pipe/p_context.h
> > @@ -523,7 +523,7 @@ struct pipe_context {
> > */
> > void (*launch_grid)(struct pipe_context *context,
> > const uint *block_layout, const uint *grid_layout,
> > - uint32_t pc, const void *input);
> > + uint32_t pc, const void *input, uint dimensions);
> > /*@}*/
> >
> > /**
> > diff --git a/src/gallium/state_trackers/clover/core/kernel.cpp b/src/gallium/state_trackers/clover/core/kernel.cpp
> > index 5e5fe51..641f9b8 100644
> > --- a/src/gallium/state_trackers/clover/core/kernel.cpp
> > +++ b/src/gallium/state_trackers/clover/core/kernel.cpp
> > @@ -70,6 +70,7 @@ kernel::launch(command_queue &q,
> > const auto reduced_grid_size =
> > map(divides(), grid_size, block_size);
> > void *st = exec.bind(&q);
> > + const unsigned dims = block_size.size();
> >
> > // The handles are created during exec_context::bind(), so we need make
> > // sure to call exec_context::bind() before retrieving them.
> > @@ -93,7 +94,7 @@ kernel::launch(command_queue &q,
> > pad_vector(q, block_size, 1).data(),
> > pad_vector(q, reduced_grid_size, 1).data(),
> > find(name_equals(_name), m.syms).offset,
> > - exec.input.data());
> > + exec.input.data(), dims);
> >
> > q.pipe->set_global_binding(q.pipe, 0, exec.g_buffers.size(), NULL, NULL);
> > q.pipe->set_compute_resources(q.pipe, 0, exec.resources.size(), NULL);
> > diff --git a/src/gallium/tests/trivial/compute.c b/src/gallium/tests/trivial/compute.c
> > index 4edb857..555b8df 100644
> > --- a/src/gallium/tests/trivial/compute.c
> > +++ b/src/gallium/tests/trivial/compute.c
> > @@ -421,11 +421,11 @@ static void destroy_globals(struct context *ctx)
> >
> > static void launch_grid(struct context *ctx, const uint *block_layout,
> > const uint *grid_layout, uint32_t pc,
> > - const void *input)
> > + const void *input, uint dimensions)
> > {
> > struct pipe_context *pipe = ctx->pipe;
> >
> > - pipe->launch_grid(pipe, block_layout, grid_layout, pc, input);
> > + pipe->launch_grid(pipe, block_layout, grid_layout, pc, input, dimensions);
> > }
> >
> > static void test_system_values(struct context *ctx)
> > @@ -492,7 +492,7 @@ static void test_system_values(struct context *ctx)
> > init_tex(ctx, 0, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 76800, 0, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){4, 3, 5}, (uint []){5, 4, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){4, 3, 5}, (uint []){5, 4, 1}, 0, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -537,7 +537,7 @@ static void test_resource_access(struct context *ctx)
> > init_tex(ctx, 1, PIPE_TEXTURE_2D, true, PIPE_FORMAT_R32_FLOAT,
> > 60, 12, init1);
> > init_compute_resources(ctx, (int []) { 0, 1, -1 });
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){15, 12, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){15, 12, 1}, 0, NULL, 3);
> > check_tex(ctx, 1, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -598,7 +598,7 @@ static void test_function_calls(struct context *ctx)
> > init_tex(ctx, 0, PIPE_TEXTURE_2D, true, PIPE_FORMAT_R32_FLOAT,
> > 15, 12, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){3, 3, 3}, (uint []){5, 4, 1}, 15, NULL);
> > + launch_grid(ctx, (uint []){3, 3, 3}, (uint []){5, 4, 1}, 15, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -640,7 +640,7 @@ static void test_input_global(struct context *ctx)
> > init_globals(ctx, (int []){ 0, 1, 2, 3, -1 },
> > (uint32_t *[]){ &input[1], &input[3],
> > &input[5], &input[7] });
> > - launch_grid(ctx, (uint []){4, 1, 1}, (uint []){1, 1, 1}, 0, input);
> > + launch_grid(ctx, (uint []){4, 1, 1}, (uint []){1, 1, 1}, 0, input, 3);
> > check_tex(ctx, 0, expect, NULL);
> > check_tex(ctx, 1, expect, NULL);
> > check_tex(ctx, 2, expect, NULL);
> > @@ -704,7 +704,7 @@ static void test_private(struct context *ctx)
> > init_tex(ctx, 0, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 32768, 0, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){16, 1, 1}, (uint []){16, 1, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){16, 1, 1}, (uint []){16, 1, 1}, 0, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -791,7 +791,7 @@ static void test_local(struct context *ctx)
> > init_tex(ctx, 0, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 4096, 0, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){64, 1, 1}, (uint []){16, 1, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){64, 1, 1}, (uint []){16, 1, 1}, 0, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -846,7 +846,7 @@ static void test_sample(struct context *ctx)
> > init_compute_resources(ctx, (int []) { 1, -1 });
> > init_sampler_views(ctx, (int []) { 0, -1 });
> > init_sampler_states(ctx, 2);
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){128, 32, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){128, 32, 1}, 0, NULL, 3);
> > check_tex(ctx, 1, expect, NULL);
> > destroy_sampler_states(ctx);
> > destroy_sampler_views(ctx);
> > @@ -896,10 +896,10 @@ static void test_many_kern(struct context *ctx)
> > init_tex(ctx, 0, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 16, 0, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 0, NULL);
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 5, NULL);
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 10, NULL);
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 15, NULL);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 0, NULL, 3);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 5, NULL, 3);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 10, NULL, 3);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){1, 1, 1}, 15, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -937,7 +937,7 @@ static void test_constant(struct context *ctx)
> > init_tex(ctx, 1, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 256, 0, init);
> > init_compute_resources(ctx, (int []) { 0, 1, -1 });
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){64, 1, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){64, 1, 1}, 0, NULL, 3);
> > check_tex(ctx, 1, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -983,7 +983,7 @@ static void test_resource_indirect(struct context *ctx)
> > init_tex(ctx, 3, PIPE_BUFFER, false, PIPE_FORMAT_R32_FLOAT,
> > 256, 0, init);
> > init_compute_resources(ctx, (int []) { 0, 1, 2, 3, -1 });
> > - launch_grid(ctx, (uint []){1, 1, 1}, (uint []){64, 1, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){1, 1, 1}, (uint []){64, 1, 1}, 0, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -1091,7 +1091,7 @@ static void test_surface_ld(struct context *ctx)
> > init_compute_resources(ctx, (int []) { 0, 1, -1 });
> > init_sampler_states(ctx, 2);
> > launch_grid(ctx, (uint []){1, 1, 1}, (uint []){128, 32, 1}, 0,
> > - NULL);
> > + NULL, 3);
> > check_tex(ctx, 1, (is_int ? expecti : expectf), NULL);
> > destroy_sampler_states(ctx);
> > destroy_compute_resources(ctx);
> > @@ -1198,7 +1198,7 @@ static void test_surface_st(struct context *ctx)
> > init_compute_resources(ctx, (int []) { 0, 1, -1 });
> > init_sampler_states(ctx, 2);
> > launch_grid(ctx, (uint []){1, 1, 1}, (uint []){128, 32, 1}, 0,
> > - NULL);
> > + NULL, 3);
> > check_tex(ctx, 1, (is_int && is_signed ? expects :
> > is_int && !is_signed ? expectu :
> > expectf), check);
> > @@ -1426,7 +1426,7 @@ static void test_atom_ops(struct context *ctx, bool global)
> > init_tex(ctx, 0, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 40, 0, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){10, 1, 1}, (uint []){1, 1, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){10, 1, 1}, (uint []){1, 1, 1}, 0, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > @@ -1565,7 +1565,7 @@ static void test_atom_race(struct context *ctx, bool global)
> > init_tex(ctx, 0, PIPE_BUFFER, true, PIPE_FORMAT_R32_FLOAT,
> > 4096, 0, init);
> > init_compute_resources(ctx, (int []) { 0, -1 });
> > - launch_grid(ctx, (uint []){64, 1, 1}, (uint []){16, 1, 1}, 0, NULL);
> > + launch_grid(ctx, (uint []){64, 1, 1}, (uint []){16, 1, 1}, 0, NULL, 3);
> > check_tex(ctx, 0, expect, NULL);
> > destroy_compute_resources(ctx);
> > destroy_tex(ctx);
> > --
> > 1.9.3
--
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20140805/977c591e/attachment.sig>
More information about the llvm-commits
mailing list