From 23eb242a5f74ea04c054bf730ea61f3eb83b5f2d Mon Sep 17 00:00:00 2001
From: Noel Maersk <veox@wemakethings.net>
Date: Sat, 1 Feb 2014 15:57:26 +0200
Subject: [PATCH] kernel: make psw usable, reorder cases alphabetically.

Closes https://github.com/veox/sgminer/issues/30
---
 AUTHORS.md      |  4 ++--
 doc/KERNEL.md   | 13 +++++++------
 driver-opencl.c |  2 +-
 kernel/psw.cl   |  4 ++--
 miner.h         |  2 +-
 5 files changed, 13 insertions(+), 12 deletions(-)

diff --git a/AUTHORS.md b/AUTHORS.md
index af0d7375..1e61bb6a 100644
--- a/AUTHORS.md
+++ b/AUTHORS.md
@@ -14,10 +14,10 @@
 All current kernels are based on `scrypt`, originally by Colin Percival,
 updated by many others.
 
-* ckolivas: Con Kolivas <kernel@kolivas.org> 15qSxP1SQcUX3o4nhkfdbgyoWEFMomJ4rZ
-* zuikkis: Zuikkis LeXck7EYgxyjw13zNDxZFmmgmWffFvhmSh
 * alexkarnew/alexkarold: Alexey Karimov LMqRcHdwnZtTMH6c2kWoxSoKM5KySfaP5C
+* ckolivas: Con Kolivas <kernel@kolivas.org> 15qSxP1SQcUX3o4nhkfdbgyoWEFMomJ4rZ
 * psw: Pavel Semjanov LP6GRFvgoMxKA6AW4TVF668cNezEGZvEtr
+* zuikkis: Zuikkis LeXck7EYgxyjw13zNDxZFmmgmWffFvhmSh
 
 
 ## Testing, bug fixes, improvements
diff --git a/doc/KERNEL.md b/doc/KERNEL.md
index c3c38d5c..dfcdd1c8 100644
--- a/doc/KERNEL.md
+++ b/doc/KERNEL.md
@@ -42,6 +42,13 @@ Con Kolivas in `cgminer` and renamed to reflect the fact.
 Only supports `vectors=1`.
 
 
+### psw
+
+Pavel Semjanov optimised kernel, SHA256 speedups.
+
+[Announcement](https://bitcointalk.org/index.php?topic=369858.0).
+
+
 ### zuikkis
 
 Zuikkis' optimised kernel, based on `ckolivas`.
@@ -50,12 +57,6 @@ Only supports `vectors=1`, `lookup-gap=2` and `worksize=256`.
 
 [Announcement](https://litecointalk.org/index.php?topic=6058.msg90873#msg90873).
 
-### psw
-
-Pavel Semjanov optimised kernel, SHA256 speedups.
-
-[Announcement](https://bitcointalk.org/index.php?topic=369858.0).
-
 
 ## Submitting new kernels
 
diff --git a/driver-opencl.c b/driver-opencl.c
index 9c64d048..d3275a1c 100644
--- a/driver-opencl.c
+++ b/driver-opencl.c
@@ -1345,8 +1345,8 @@ static bool opencl_thread_init(struct thr_info *thr)
 	case KL_ALEXKARNEW:
 	case KL_ALEXKAROLD:
 	case KL_CKOLIVAS:
-	case KL_ZUIKKIS:
 	case KL_PSW:
+	case KL_ZUIKKIS:
 		thrdata->queue_kernel_parameters = &queue_scrypt_kernel;
 		break;
 	default:
diff --git a/kernel/psw.cl b/kernel/psw.cl
index de084936..acd805e0 100644
--- a/kernel/psw.cl
+++ b/kernel/psw.cl
@@ -747,8 +747,8 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 	unshittify(X);
 }
 
-#define SCRYPT_FOUND (0xFF)
-#define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
+#define FOUND (0xFF)
+#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void search(__global const uint4 * restrict input,
diff --git a/miner.h b/miner.h
index 975461dd..107d4c5c 100644
--- a/miner.h
+++ b/miner.h
@@ -375,8 +375,8 @@ enum cl_kernels {
 	KL_ALEXKARNEW,
 	KL_ALEXKAROLD,
 	KL_CKOLIVAS,
-	KL_ZUIKKIS,
 	KL_PSW,
+	KL_ZUIKKIS,
 };
 
 enum dev_reason {