[Mesa-dev] [PATCH 2/2] radeonsi: convert dispatch packet to little endian

[Marek Olšák]

I pushed both patches, thanks.

Marek

On Mon, Apr 9, 2018 at 7:06 AM, Bas Vermeulen <bas at daedalean.ai> wrote:

> The parameters for the compute engine are wrong when using
> an E8860 on a big endian machine.
> To fix this, convert the contents of struct dispatch_packet
> to little endian.
>
> This ensures that get_global_id(0) and similar functions
> in the OpenCL code get the correct endian values, and
> makes my simple OpenCL program work correctly.
>
> Signed-off-by: Bas Vermeulen <bas at daedalean.ai>
> ---
>  src/gallium/drivers/radeonsi/si_compute.c | 24 ++++++++++++------------
>  1 file changed, 12 insertions(+), 12 deletions(-)
>
> diff --git a/src/gallium/drivers/radeonsi/si_compute.c
> b/src/gallium/drivers/radeonsi/si_compute.c
> index dfede47605..8ac5b262c4 100644
> --- a/src/gallium/drivers/radeonsi/si_compute.c
> +++ b/src/gallium/drivers/radeonsi/si_compute.c
> @@ -564,18 +564,18 @@ static void si_setup_user_sgprs_co_v2(struct
> si_context *sctx,
>                 /* Upload dispatch ptr */
>                 memset(&dispatch, 0, sizeof(dispatch));
>
> -               dispatch.workgroup_size_x = info->block[0];
> -               dispatch.workgroup_size_y = info->block[1];
> -               dispatch.workgroup_size_z = info->block[2];
> +               dispatch.workgroup_size_x = util_cpu_to_le16(info->block[
> 0]);
> +               dispatch.workgroup_size_y = util_cpu_to_le16(info->block[
> 1]);
> +               dispatch.workgroup_size_z = util_cpu_to_le16(info->block[
> 2]);
>
> -               dispatch.grid_size_x = info->grid[0] * info->block[0];
> -               dispatch.grid_size_y = info->grid[1] * info->block[1];
> -               dispatch.grid_size_z = info->grid[2] * info->block[2];
> +               dispatch.grid_size_x = util_cpu_to_le32(info->grid[0] *
> info->block[0]);
> +               dispatch.grid_size_y = util_cpu_to_le32(info->grid[1] *
> info->block[1]);
> +               dispatch.grid_size_z = util_cpu_to_le32(info->grid[2] *
> info->block[2]);
>
> -               dispatch.private_segment_size = program->private_size;
> -               dispatch.group_segment_size = program->local_size;
> +               dispatch.private_segment_size = util_cpu_to_le32(program->
> private_size);
> +               dispatch.group_segment_size = util_cpu_to_le32(program->
> local_size);
>
> -               dispatch.kernarg_address = kernel_args_va;
> +               dispatch.kernarg_address = util_cpu_to_le64(kernel_args_
> va);
>
>                 u_upload_data(sctx->b.const_uploader, 0, sizeof(dispatch),
>                                256, &dispatch, &dispatch_offset,
> @@ -652,9 +652,9 @@ static bool si_upload_compute_input(struct si_context
> *sctx,
>
>         if (!code_object) {
>                 for (i = 0; i < 3; i++) {
> -                       kernel_args[i] = info->grid[i];
> -                       kernel_args[i + 3] = info->grid[i] *
> info->block[i];
> -                       kernel_args[i + 6] = info->block[i];
> +                       kernel_args[i] = util_cpu_to_le32(info->grid[i]);
> +                       kernel_args[i + 3] =
> util_cpu_to_le32(info->grid[i] * info->block[i]);
> +                       kernel_args[i + 6] = util_cpu_to_le32(info->block[
> i]);
>                 }
>         }
>
> --
> 2.14.1
>
>
> --
> This message has been scanned for viruses and
> dangerous content by MailScanner, and is
> believed to be clean.
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20180409/15b2b4bf/attachment-0001.html>