From d5fbf8f987bcf2127e034d4b9ee68bf627560875 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Sun, 24 Jul 2011 01:09:30 +1000
Subject: [PATCH] Some hardware might benefit from the less OPS so there's no
 harm in leaving it there apart from readability of the code.

Revert "Simplify the output part of the kernel. There's no demonstrable advantage from more complexity."

This reverts commit 53b53cb666ff09dfd9034fc6f6987516d0f6fcc7.
---
 phatk110722.cl | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/phatk110722.cl b/phatk110722.cl
index 77379d1c..43cdacf7 100644
--- a/phatk110722.cl
+++ b/phatk110722.cl
@@ -394,26 +394,26 @@ __kernel void search(	const uint state0, const uint state1, const uint state2, c
 #if defined(VECTORS4) || defined(VECTORS2)
 	if (Vals[7].x == -H[7])
 	{
-		output[MAXBUFFERS] = output[NFLAG & W_3.x] =  W_3.x;
+		output[MAXBUFFERS] = output[NFLAG & (W[3].x >> 2)] =  W_3.x;
 	}
 	if (Vals[7].y == -H[7])
 	{
-		output[MAXBUFFERS] = output[NFLAG & W_3.y] =  W_3.y;
+		output[MAXBUFFERS] = output[NFLAG & (W[3].y >> 2)] =  W_3.y;
 	}
 #ifdef VECTORS4
 	if (Vals[7].z == -H[7])
 	{
-		output[MAXBUFFERS] = output[NFLAG & W_3.z] =  W_3.z;
+		output[MAXBUFFERS] = output[NFLAG & (W[3].z >> 2)] =  W_3.z;
 	}
 	if (Vals[7].w == -H[7])
 	{
-		output[MAXBUFFERS] = output[NFLAG & W_3.w] =  W_3.w;
+		output[MAXBUFFERS] = output[NFLAG & (W[3].w >> 2)] =  W_3.w;
 	}
 #endif
 #else
 	if (Vals[7] == -H[7])
 	{
-		output[MAXBUFFERS] = output[NFLAG & W_3] =  W_3;
+		output[MAXBUFFERS] = output[NFLAG & (W[3] >> 2)] =  W_3;
 	}
 #endif
 }