kernel: allow choosing Alexey Karimov's kernels.

Am I the only one who gets a segfault?..

configure.ac

@@ -317,10 +317,11 @@ fi
 
 AC_DEFINE_UNQUOTED([SGMINER_PREFIX], ["$prefix/bin"], [Path to sgminer install])
 
+AC_DEFINE_UNQUOTED([ALEXKARNEW_KERNNAME], ["alexkarnew"], [Filename for Alexey Karimov's optimised kernel for Catalyst >=13.4])
+AC_DEFINE_UNQUOTED([ALEXKAROLD_KERNNAME], ["alexkarold"], [Filename for Alexey Karimov's optimised kernel for Catalyst <13.4])
 AC_DEFINE_UNQUOTED([CKOLIVAS_KERNNAME], ["ckolivas"], [Filename for original scrypt kernel])
 AC_DEFINE_UNQUOTED([ZUIKKIS_KERNNAME], ["zuikkis"], [Filename for Zuikkis' optimised kernel])
 
-
 AC_SUBST(OPENCL_LIBS)
 AC_SUBST(OPENCL_FLAGS)
 AC_SUBST(JANSSON_LIBS)

driver-opencl.c

@@ -196,6 +196,10 @@ char *set_thread_concurrency(char *arg)
 
 static enum cl_kernels select_kernel(char *arg)
 {
+	if (!strcmp(arg, "alexkarnew"))
+		return KL_ALEXKARNEW;
+	if (!strcmp(arg, "alexkarold"))
+		return KL_ALEXKAROLD;
 	if (!strcmp(arg, "ckolivas"))
 		return KL_CKOLIVAS;
 	if (!strcmp(arg, "zuikkis"))
@@ -1212,6 +1216,12 @@ static bool opencl_thread_prepare(struct thr_info *thr)
 	if (!cgpu->kname)
 	{
 		switch (clStates[i]->chosen_kernel) {
+			case KL_ALEXKARNEW:
+				cgpu->kname = "alexkarnew";
+				break;
+			case KL_ALEXKAROLD:
+				cgpu->kname = "alexkarold";
+				break;
 			case KL_CKOLIVAS:
 				cgpu->kname = "ckolivas";
 				break;

kernel/alexkarnew.cl

@@ -813,8 +813,8 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 	unshittify(X);
 }
 
-#define SCRYPT_FOUND (0xFF)
+#define FOUND (0xFF)
-#define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
+#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void search(__global const uint4 * restrict input,

kernel/alexkarold.cl

@@ -810,8 +810,8 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 	unshittify(X);
 }
 
-#define SCRYPT_FOUND (0xFF)
+#define FOUND (0xFF)
-#define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
+#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

@@ -372,6 +372,8 @@ enum dev_enable {
 
 enum cl_kernels {
 	KL_NONE,
+	KL_ALEXKARNEW,
+	KL_ALEXKAROLD,
 	KL_CKOLIVAS,
 	KL_ZUIKKIS,
 };

ocl.c

@@ -416,16 +416,28 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 		preferred_vwidth = 2;
 
 	switch (clState->chosen_kernel) {
+		case KL_ALEXKARNEW:
+			strcpy(filename, ALEXKARNEW_KERNNAME".cl");
+			strcpy(binaryfilename, ALEXKARNEW_KERNNAME);
+			/* Kernel only supports vector 1 */
+			cgpu->vwidth = 1;
+			break;
+		case KL_ALEXKAROLD:
+			strcpy(filename, ALEXKAROLD_KERNNAME".cl");
+			strcpy(binaryfilename, ALEXKAROLD_KERNNAME);
+			/* Kernel only supports vector 1 */
+			cgpu->vwidth = 1;
+			break;
 		case KL_CKOLIVAS:
 			strcpy(filename, CKOLIVAS_KERNNAME".cl");
 			strcpy(binaryfilename, CKOLIVAS_KERNNAME);
-			/* Scrypt only supports vector 1 */
+			/* Kernel only supports vector 1 */
 			cgpu->vwidth = 1;
 			break;
 		case KL_ZUIKKIS:
 			strcpy(filename, ZUIKKIS_KERNNAME".cl");
 			strcpy(binaryfilename, ZUIKKIS_KERNNAME);
-			/* Scrypt only supports vector 1 */
+			/* Kernel only supports vector 1 */
 			cgpu->vwidth = 1;
 			break;
 		case KL_NONE: /* Shouldn't happen */

sgminer.c

@@ -4121,6 +4121,12 @@ void write_config(FILE *fcfg)
 			switch (gpus[i].kernel) {
 				case KL_NONE: // Shouldn't happen
 					break;
+				case KL_ALEXKARNEW:
+					fprintf(fcfg, "alexkarnew");
+					break;
+				case KL_ALEXKAROLD:
+					fprintf(fcfg, "alexkarold");
+					break;
 				case KL_CKOLIVAS:
 					fprintf(fcfg, "ckolivas");
 					break;