From f84fcb055beb3403fe2ce669b1e7fc6bf285f465 Mon Sep 17 00:00:00 2001
From: julienhascoet <jhascoet@kalray.eu>
Date: Fri, 20 Nov 2020 15:07:14 +0100
Subject: [PATCH] Fix aliasing issue for test_half (#45) (#1004)

OpenCL-C code must respect aliasing rules as in C99.
Therefore, accessing data from multiple pointer sizes
must all point to character type to keep strict aliasing
rules. If not, when the compiler performs optimization,
the dataflow was broken leading to failure.
We also fix the formatting and refactor few things.
---
 test_conformance/half/Test_vLoadHalf.cpp | 33 +++++++++++++++++++-----
 1 file changed, 27 insertions(+), 6 deletions(-)

diff --git a/test_conformance/half/Test_vLoadHalf.cpp b/test_conformance/half/Test_vLoadHalf.cpp
index 1c53492a..52867c25 100644
--- a/test_conformance/half/Test_vLoadHalf.cpp
+++ b/test_conformance/half/Test_vLoadHalf.cpp
@@ -170,17 +170,38 @@ int Test_vLoadHalf_private( cl_device_id device, bool aligned )
         };
 
         const char *source_local2[] = {
-            "__kernel void test( const __global half *p, __global float", vector_size_name_extensions[vectorSize], " *f )\n"
+            "#define VECTOR_LEN (",
+            vector_size_names[vectorSize],
+            "/",
+            align_divisors[vectorSize],
+            ")\n"
+            "#define ALIGN_TYPE ",
+            align_types[vectorSize],
+            "\n"
+            "__kernel void test( const __global half *p, __global float",
+            vector_size_name_extensions[vectorSize],
+            " *f )\n"
             "{\n"
-            "   __local ", align_types[vectorSize], " data[", local_buf_size, "/", align_divisors[vectorSize], "];\n"
+            "   __local uchar data[",
+            local_buf_size,
+            "/",
+            align_divisors[vectorSize],
+            "*sizeof(ALIGN_TYPE)] ",
+            "__attribute__((aligned(sizeof(ALIGN_TYPE))));\n"
             "   __local half* hdata_p = (__local half*) data;\n"
-            "   __global ", align_types[vectorSize], "* i_p = (__global ", align_types[vectorSize],"*)p;\n"
+            "   __global ALIGN_TYPE* i_p = (__global ALIGN_TYPE*)p;\n"
             "   size_t i = get_global_id(0);\n"
             "   size_t lid = get_local_id(0);\n"
             "   int k;\n"
-            "   for (k=0; k<",vector_size_names[vectorSize],"/",align_divisors[vectorSize],"; k++)\n"
-            "     data[lid*",vector_size_names[vectorSize],"/",align_divisors[vectorSize],"+k] = i_p[i*",vector_size_names[vectorSize],"/",align_divisors[vectorSize],"+k];\n"
-            "   f[i] = vload", aligned ? "a" : "", "_half",vector_size_name_extensions[vectorSize],"( lid, hdata_p );\n"
+            "   for (k=0; k<VECTOR_LEN; k++)\n"
+            "     *(__local ",
+            "ALIGN_TYPE*)&(data[(lid*VECTOR_LEN+k)*sizeof(ALIGN_TYPE)]) = ",
+            "i_p[i*VECTOR_LEN+k];\n"
+            "   f[i] = vload",
+            aligned ? "a" : "",
+            "_half",
+            vector_size_name_extensions[vectorSize],
+            "( lid, hdata_p );\n"
             "}\n"
         };