From fad6a005c99932346e60b5add36fb83ac583597d Mon Sep 17 00:00:00 2001
From: Stephen Clarke <stephen.clarke@imgtec.com>
Date: Wed, 3 Feb 2021 14:34:18 +0000
Subject: [PATCH] Fix test_vector_swizzle possible overwrite of padding of
 3-element vectors (#1124)

Use vstore3 to ensure padding is not overwritten in destination buffer.

Fixes #1120
---
 .../basic/test_vector_swizzle.cpp             | 48 +++++++++++++------
 1 file changed, 33 insertions(+), 15 deletions(-)

diff --git a/test_conformance/basic/test_vector_swizzle.cpp b/test_conformance/basic/test_vector_swizzle.cpp
index 67bf7537..5ab3ea4f 100644
--- a/test_conformance/basic/test_vector_swizzle.cpp
+++ b/test_conformance/basic/test_vector_swizzle.cpp
@@ -94,11 +94,17 @@ __kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
     int index = 0;
 
     // lvalue swizzles
-    dst[index++].x = value.x;
-    dst[index++].y = value.x;
-    dst[index++].z = value.x;
-    dst[index++].xyz = value;
-    dst[index++].zyx = value;
+    TYPE t;
+    t = dst[index]; t.x = value.x;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.y = value.x;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.z = value.x;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.xyz = value;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.zyx = value;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
 
     // rvalue swizzles
     vstore3(value.x, 0, (__global BASETYPE*)(dst + index++));
@@ -114,11 +120,17 @@ __kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
     int index = 0;
 
     // lvalue swizzles
-    dst[index++].r = value.r;
-    dst[index++].g = value.r;
-    dst[index++].b = value.r;
-    dst[index++].rgb = value;
-    dst[index++].bgr = value;
+    TYPE t;
+    t = dst[index]; t.r = value.r;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.g = value.r;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.b = value.r;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.rgb = value;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.bgr = value;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
 
     // rvalue swizzles
     vstore3(value.r, 0, (__global BASETYPE*)(dst + index++));
@@ -134,11 +146,17 @@ __kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
     int index = 0;
 
     // lvalue swizzles
-    dst[index++].s0 = value.s0;
-    dst[index++].s1 = value.s0;
-    dst[index++].s2 = value.s0;
-    dst[index++].s012 = value;
-    dst[index++].s210 = value;
+    TYPE t;
+    t = dst[index]; t.s0 = value.s0;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.s1 = value.s0;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.s2 = value.s0;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.s012 = value;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
+    t = dst[index]; t.s210 = value;
+    vstore3(t, 0, (__global BASETYPE*)(dst + index++));
 
     // rvalue swizzles
     vstore3(value.s0, 0, (__global BASETYPE*)(dst + index++));