Commit aabc723326a2c2f80b48f43e753bfca47e0cd156
2012-07-14T00:30:25
Make sure goffset is set for scrypt and drop padbuffer8 to something manageable for now.
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46
diff --git a/driver-opencl.c b/driver-opencl.c
index d50f1ff..5e79d2b 100644
--- a/driver-opencl.c
+++ b/driver-opencl.c
@@ -1470,7 +1470,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
status = clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, NULL,
globalThreads, localThreads, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
- applog(LOG_ERR, "Error: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)");
+ applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status);
return -1;
}
diff --git a/ocl.c b/ocl.c
index 48d78b4..95f2562 100644
--- a/ocl.c
+++ b/ocl.c
@@ -428,8 +428,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
gpus[gpu].vwidth = preferred_vwidth;
}
- if ((clState->chosen_kernel == KL_POCLBM || clState->chosen_kernel == KL_DIABLO || clState->chosen_kernel == KL_DIAKGCN) &&
- clState->vwidth == 1 && clState->hasOpenCL11plus)
+ if (((clState->chosen_kernel == KL_POCLBM || clState->chosen_kernel == KL_DIABLO || clState->chosen_kernel == KL_DIAKGCN) &&
+ clState->vwidth == 1 && clState->hasOpenCL11plus) || opt_scrypt)
clState->goffset = true;
if (gpus[gpu].work_size && gpus[gpu].work_size <= clState->max_work_size)
@@ -536,7 +536,7 @@ build:
char *CompilerOptions = calloc(1, 256);
if (opt_scrypt) {
- sprintf(CompilerOptions, "-D LOOKUP_GAP=2 -D CONCURRENT_THREADS=6144 -D WORKSIZE=%d",
+ sprintf(CompilerOptions, "-D LOOKUP_GAP=1 -D CONCURRENT_THREADS=1 -D WORKSIZE=%d",
(int)clState->wsize);
} else {
sprintf(CompilerOptions, "-D WORKSIZE=%d -D VECTORS%d -D WORKVEC=%d",
@@ -735,7 +735,7 @@ built:
#ifdef USE_SCRYPT
if (opt_scrypt) {
clState->CLbuffer0 = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, 128, NULL, &status);
- clState->padbuffer8 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, 402653184, NULL, &status);
+ clState->padbuffer8 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, 131072, NULL, &status);
}
#endif
clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_WRITE_ONLY, BUFFERSIZE, NULL, &status);