diff options
| author | Bas Vermeulen <[email protected]> | 2018-04-09 13:06:01 +0200 |
|---|---|---|
| committer | Marek Olšák <[email protected]> | 2018-04-09 13:47:52 -0400 |
| commit | c63bef15fc6f55ad9d5aa82076b94108f6b54236 (patch) | |
| tree | 16b6cef2a09c49554cb0a5fda31c50dfe46a8816 /src/gallium/drivers/radeonsi | |
| parent | be628e47494fb19db593a2006feb1c3688014d92 (diff) | |
radeonsi: convert dispatch packet to little endian
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 <[email protected]>
Signed-off-by: Marek Olšák <[email protected]>
Reviewed-by: Michel Dänzer <[email protected]>
Diffstat (limited to 'src/gallium/drivers/radeonsi')
| -rw-r--r-- | src/gallium/drivers/radeonsi/si_compute.c | 24 |
1 files changed, 12 insertions, 12 deletions
diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c index dfede47605f..8ac5b262c40 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]); } } |