<div dir="ltr"><div>I pushed both patches, thanks.<br><br></div>Marek<br></div><div class="gmail_extra"><br><div class="gmail_quote">On Mon, Apr 9, 2018 at 7:06 AM, Bas Vermeulen <span dir="ltr"><<a href="mailto:bas@daedalean.ai" target="_blank">bas@daedalean.ai</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">The parameters for the compute engine are wrong when using<br>
an E8860 on a big endian machine.<br>
To fix this, convert the contents of struct dispatch_packet<br>
to little endian.<br>
<br>
This ensures that get_global_id(0) and similar functions<br>
in the OpenCL code get the correct endian values, and<br>
makes my simple OpenCL program work correctly.<br>
<br>
Signed-off-by: Bas Vermeulen <<a href="mailto:bas@daedalean.ai">bas@daedalean.ai</a>><br>
---<br>
 src/gallium/drivers/radeonsi/<wbr>si_compute.c | 24 ++++++++++++------------<br>
 1 file changed, 12 insertions(+), 12 deletions(-)<br>
<br>
diff --git a/src/gallium/drivers/<wbr>radeonsi/si_compute.c b/src/gallium/drivers/<wbr>radeonsi/si_compute.c<br>
index dfede47605..8ac5b262c4 100644<br>
--- a/src/gallium/drivers/<wbr>radeonsi/si_compute.c<br>
+++ b/src/gallium/drivers/<wbr>radeonsi/si_compute.c<br>
@@ -564,18 +564,18 @@ static void si_setup_user_sgprs_co_v2(<wbr>struct si_context *sctx,<br>
                /* Upload dispatch ptr */<br>
                memset(&dispatch, 0, sizeof(dispatch));<br>
<br>
-               dispatch.workgroup_size_x = info->block[0];<br>
-               dispatch.workgroup_size_y = info->block[1];<br>
-               dispatch.workgroup_size_z = info->block[2];<br>
+               dispatch.workgroup_size_x = util_cpu_to_le16(info->block[<wbr>0]);<br>
+               dispatch.workgroup_size_y = util_cpu_to_le16(info->block[<wbr>1]);<br>
+               dispatch.workgroup_size_z = util_cpu_to_le16(info->block[<wbr>2]);<br>
<br>
-               dispatch.grid_size_x = info->grid[0] * info->block[0];<br>
-               dispatch.grid_size_y = info->grid[1] * info->block[1];<br>
-               dispatch.grid_size_z = info->grid[2] * info->block[2];<br>
+               dispatch.grid_size_x = util_cpu_to_le32(info->grid[0] * info->block[0]);<br>
+               dispatch.grid_size_y = util_cpu_to_le32(info->grid[1] * info->block[1]);<br>
+               dispatch.grid_size_z = util_cpu_to_le32(info->grid[2] * info->block[2]);<br>
<br>
-               dispatch.private_segment_size = program->private_size;<br>
-               dispatch.group_segment_size = program->local_size;<br>
+               dispatch.private_segment_size = util_cpu_to_le32(program-><wbr>private_size);<br>
+               dispatch.group_segment_size = util_cpu_to_le32(program-><wbr>local_size);<br>
<br>
-               dispatch.kernarg_address = kernel_args_va;<br>
+               dispatch.kernarg_address = util_cpu_to_le64(kernel_args_<wbr>va);<br>
<br>
                u_upload_data(sctx->b.const_<wbr>uploader, 0, sizeof(dispatch),<br>
                               256, &dispatch, &dispatch_offset,<br>
@@ -652,9 +652,9 @@ static bool si_upload_compute_input(struct si_context *sctx,<br>
<br>
        if (!code_object) {<br>
                for (i = 0; i < 3; i++) {<br>
-                       kernel_args[i] = info->grid[i];<br>
-                       kernel_args[i + 3] = info->grid[i] * info->block[i];<br>
-                       kernel_args[i + 6] = info->block[i];<br>
+                       kernel_args[i] = util_cpu_to_le32(info->grid[i]<wbr>);<br>
+                       kernel_args[i + 3] = util_cpu_to_le32(info->grid[i] * info->block[i]);<br>
+                       kernel_args[i + 6] = util_cpu_to_le32(info->block[<wbr>i]);<br>
                }<br>
        }<br>
<span class="HOEnZb"><font color="#888888"><br>
--<br>
2.14.1<br>
<br>
<br>
--<br>
This message has been scanned for viruses and<br>
dangerous content by MailScanner, and is<br>
believed to be clean.<br>
<br>
</font></span></blockquote></div><br></div>