kernel: make psw usable, reorder cases alphabetically.

Closes https://github.com/veox/sgminer/issues/30

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

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
 

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:

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,

miner.h

@@ -375,8 +375,8 @@ enum cl_kernels {
 	KL_ALEXKARNEW,
 	KL_ALEXKAROLD,
 	KL_CKOLIVAS,
-	KL_ZUIKKIS,
 	KL_PSW,
+	KL_ZUIKKIS,
 };
 
 enum dev_reason {