[oclgrind] 02/03: Fix big-endian issues
James Price
jprice-guest at moszumanska.debian.org
Thu Jan 19 13:10:45 UTC 2017
This is an automated email from the git hooks/post-receive script.
jprice-guest pushed a commit to branch master
in repository oclgrind.
commit 9c433785079003ec0bcb5251a875c2b47f82dcb7
Author: James Price <j.price at bristol.ac.uk>
Date: Thu Jan 19 13:05:32 2017 +0000
Fix big-endian issues
---
debian/patches/big-endian.patch | 239 ++++++++++++++++++++++++++++++++++++++++
debian/patches/series | 1 +
2 files changed, 240 insertions(+)
diff --git a/debian/patches/big-endian.patch b/debian/patches/big-endian.patch
new file mode 100644
index 0000000..86ce95a
--- /dev/null
+++ b/debian/patches/big-endian.patch
@@ -0,0 +1,239 @@
+Author: James Price <j.price at bristol.ac.uk>
+Description: Fix big-endian issues
+Origin: upstream commits bccc640 and 2e138bc
+Last-Update: 2017-01-19
+--- a/src/core/WorkItem.cpp
++++ b/src/core/WorkItem.cpp
+@@ -1135,7 +1135,7 @@
+ TypedValue op = getOperand(instruction->getOperand(0));
+ for (unsigned i = 0; i < result.num; i++)
+ {
+- memcpy(result.data+i*result.size, op.data+i*op.size, result.size);
++ result.setUInt(op.getUInt(i), i);
+ }
+ }
+
+--- a/tests/kernels/atomics/atomic_cmpxchg_read_race.cl
++++ b/tests/kernels/atomics/atomic_cmpxchg_read_race.cl
+@@ -7,6 +7,6 @@
+ }
+ else
+ {
+- atomic_cmpxchg(data, 0, i);
++ atomic_cmpxchg(data, 0, 0x01000001);
+ }
+ }
+--- a/tests/kernels/atomics/atomic_cmpxchg_read_race.ref
++++ b/tests/kernels/atomics/atomic_cmpxchg_read_race.ref
+@@ -2,4 +2,4 @@
+ ERROR Write-write data race at global memory
+
+ EXACT Argument 'data': 4 bytes
+-EXACT data[0] = 1
++MATCH data[0] =
+--- a/tests/kernels/atomics/atomic_race_after.cl
++++ b/tests/kernels/atomics/atomic_race_after.cl
+@@ -1,8 +1,8 @@
+-kernel void atomic_race_after(global int *data)
++kernel void atomic_race_after(global int *data, global int *output)
+ {
+ atomic_inc(data);
+ if (get_global_id(0) == get_global_size(0)-1)
+ {
+- (*data)++;
++ *output = *data;
+ }
+ }
+--- a/tests/kernels/atomics/atomic_race_after.ref
++++ b/tests/kernels/atomics/atomic_race_after.ref
+@@ -1,6 +1,4 @@
+ ERROR Read-write data race at global memory
+-ERROR Read-write data race at global memory
+-ERROR Write-write data race at global memory
+
+-EXACT Argument 'data': 4 bytes
+-EXACT data[0] = 5
++EXACT Argument 'output': 4 bytes
++MATCH output[0] =
+--- a/tests/kernels/atomics/atomic_race_after.sim
++++ b/tests/kernels/atomics/atomic_race_after.sim
+@@ -3,4 +3,5 @@
+ 4 1 1
+ 4 1 1
+
++<size=4 fill=0>
+ <size=4 fill=0 dump>
+--- a/tests/kernels/atomics/atomic_race_before.cl
++++ b/tests/kernels/atomics/atomic_race_before.cl
+@@ -4,5 +4,5 @@
+ {
+ *data = 0;
+ }
+- atomic_inc(data);
++ atomic_dec(data);
+ }
+--- a/tests/kernels/atomics/atomic_race_before.ref
++++ b/tests/kernels/atomics/atomic_race_before.ref
+@@ -6,4 +6,4 @@
+ ERROR Write-write data race at global memory address
+
+ EXACT Argument 'data': 4 bytes
+-EXACT data[0] = 4
++EXACT data[0] = -4
+--- a/tests/kernels/uninitialized/padded_nested_struct_memcpy.cl
++++ b/tests/kernels/uninitialized/padded_nested_struct_memcpy.cl
+@@ -17,10 +17,10 @@
+ {
+ struct S s;
+ s.a = 1;
+- s.b = 2;
++ s.b = 0x02000002;
+ s.c = 3;
+ s.d.a = 4;
+- s.d.b = 5;
++ s.d.b = 0x05000005;
+ s.d.c = 6;
+
+ *output = s;
+--- a/tests/kernels/uninitialized/padded_nested_struct_memcpy.ref
++++ b/tests/kernels/uninitialized/padded_nested_struct_memcpy.ref
+@@ -6,7 +6,7 @@
+ EXACT output[4] = 2
+ EXACT output[5] = 0
+ EXACT output[6] = 0
+-EXACT output[7] = 0
++EXACT output[7] = 2
+ EXACT output[8] = 3
+ MATCH output[9] =
+ MATCH output[10] =
+@@ -18,7 +18,7 @@
+ EXACT output[16] = 5
+ EXACT output[17] = 0
+ EXACT output[18] = 0
+-EXACT output[19] = 0
++EXACT output[19] = 5
+ EXACT output[20] = 6
+ MATCH output[21] =
+ MATCH output[22] =
+--- a/tests/kernels/uninitialized/padded_struct_alloca_fp.cl
++++ b/tests/kernels/uninitialized/padded_struct_alloca_fp.cl
+@@ -9,7 +9,7 @@
+ {
+ struct S s;
+ s.a = 42;
+- s.b = -7;
++ s.b = 0xF9FFFFF9;
+ s.c = 127;
+
+ *output = s;
+--- a/tests/kernels/uninitialized/padded_struct_alloca_fp.ref
++++ b/tests/kernels/uninitialized/padded_struct_alloca_fp.ref
+@@ -1,4 +1,13 @@
+ EXACT Argument 'output': 12 bytes
+ EXACT output[0] = 42
+-EXACT output[1] = -7
+-EXACT output[2] = 127
++EXACT output[1] = 0
++EXACT output[2] = 0
++EXACT output[3] = 0
++EXACT output[4] = -7
++EXACT output[5] = -1
++EXACT output[6] = -1
++EXACT output[7] = -7
++EXACT output[8] = 127
++EXACT output[9] = 0
++EXACT output[10] = 0
++EXACT output[11] = 0
+--- a/tests/kernels/uninitialized/padded_struct_alloca_fp.sim
++++ b/tests/kernels/uninitialized/padded_struct_alloca_fp.sim
+@@ -3,4 +3,4 @@
+ 1 1 1
+ 1 1 1
+
+-<size=12 int fill=0 dump>
++<size=12 char fill=0 dump>
+--- a/tests/kernels/uninitialized/padded_struct_memcpy_fp.cl
++++ b/tests/kernels/uninitialized/padded_struct_memcpy_fp.cl
+@@ -12,7 +12,7 @@
+
+ struct S s;
+ s.a = 42;
+- s.b = -7;
++ s.b = 0xF9FFFFF9;
+ s.c = 127;
+
+ if (lid == 0)
+--- a/tests/kernels/uninitialized/padded_struct_memcpy_fp.ref
++++ b/tests/kernels/uninitialized/padded_struct_memcpy_fp.ref
+@@ -1,4 +1,13 @@
+ EXACT Argument 'output': 12 bytes
+ EXACT output[0] = 42
+-EXACT output[1] = -7
+-EXACT output[2] = 127
++EXACT output[1] = 0
++EXACT output[2] = 0
++EXACT output[3] = 0
++EXACT output[4] = -7
++EXACT output[5] = -1
++EXACT output[6] = -1
++EXACT output[7] = -7
++EXACT output[8] = 127
++EXACT output[9] = 0
++EXACT output[10] = 0
++EXACT output[11] = 0
+--- a/tests/kernels/uninitialized/padded_struct_memcpy_fp.sim
++++ b/tests/kernels/uninitialized/padded_struct_memcpy_fp.sim
+@@ -5,4 +5,4 @@
+
+ <size=12 char>
+
+-<size=12 int fill=0 dump>
++<size=12 char fill=0 dump>
+--- a/tests/kernels/uninitialized/uninitialized_padded_nested_struct_memcpy.cl
++++ b/tests/kernels/uninitialized/uninitialized_padded_nested_struct_memcpy.cl
+@@ -15,6 +15,6 @@
+
+ kernel void uninitialized_padded_nested_struct_memcpy(local int *scratch, global struct S *output)
+ {
+- struct S s = {1, 2, 3, {4, *scratch, 5}};
++ struct S s = {1, 0x02000002, 3, {4, *scratch, 5}};
+ *output = s;
+ }
+--- a/tests/kernels/uninitialized/uninitialized_padded_nested_struct_memcpy.ref
++++ b/tests/kernels/uninitialized/uninitialized_padded_nested_struct_memcpy.ref
+@@ -8,7 +8,7 @@
+ EXACT output[4] = 2
+ EXACT output[5] = 0
+ EXACT output[6] = 0
+-EXACT output[7] = 0
++EXACT output[7] = 2
+ EXACT output[8] = 3
+ MATCH output[9] =
+ MATCH output[10] =
+--- a/tests/kernels/uninitialized/uninitialized_padded_struct_memcpy.ref
++++ b/tests/kernels/uninitialized/uninitialized_padded_struct_memcpy.ref
+@@ -2,5 +2,14 @@
+
+ EXACT Argument 'output': 12 bytes
+ EXACT output[0] = 1
+-MATCH output[1] =
+-EXACT output[2] = 2
++EXACT output[1] = 0
++EXACT output[2] = 0
++EXACT output[3] = 0
++MATCH output[4] =
++MATCH output[5] =
++MATCH output[6] =
++MATCH output[7] =
++EXACT output[8] = 2
++EXACT output[9] = 0
++EXACT output[10] = 0
++EXACT output[11] = 0
+--- a/tests/kernels/uninitialized/uninitialized_padded_struct_memcpy.sim
++++ b/tests/kernels/uninitialized/uninitialized_padded_struct_memcpy.sim
+@@ -5,4 +5,4 @@
+
+ <size=8>
+
+-<size=12 int fill=0 dump>
++<size=12 char fill=0 dump>
diff --git a/debian/patches/series b/debian/patches/series
index 31d46c6..69b1ca1 100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@ -4,3 +4,4 @@ pch-location.patch
use-opencl-headers.patch
altivec-fix.patch
align-alloc.patch
+big-endian.patch
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/oclgrind.git
More information about the Pkg-opencl-commits
mailing list