Commit b198badcf44d33a1f66fb68d23919cb655cc6f3d
2011-07-17T14:42:44
The poclbm kernel needs to be updated to work with the change to 4k sized output buffers.
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 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91
diff --git a/poclbm.cl b/poclbm.cl
index 709cf4f..c5358a4 100644
--- a/poclbm.cl
+++ b/poclbm.cl
@@ -80,7 +80,6 @@ __kernel void search( const uint state0, const uint state1, const uint state2, c
u W[24];
u Vals[8];
u nonce;
- uint it = get_local_id(0);
#ifdef VECTORS4
nonce = base + (get_global_id(0)<<2) + (uint4)(0, 1, 2, 3);
@@ -628,70 +627,32 @@ __kernel void search( const uint state0, const uint state1, const uint state2, c
Vals[7]+=0x5be0cd19U;
-#define MAXBUFFERS (4 * 512)
+#define MAXBUFFERS (4095)
+#define NFLAG (0xFFFUL)
#if defined(VECTORS4) || defined(VECTORS2)
if (Vals[7].x == 0)
{
- // Unlikely event there is something here already !
- if (output[it]) {
- for (it = 0; it < MAXBUFFERS; it++) {
- if (!output[it])
- break;
- }
- }
- output[it] = nonce.x;
- output[MAXBUFFERS] = 1;
+ output[MAXBUFFERS] = output[NFLAG & nonce.x] = nonce.x;
}
if (Vals[7].y == 0)
{
- it += 512;
- if (output[it]) {
- for (it = 0; it < MAXBUFFERS; it++) {
- if (!output[it])
- break;
- }
- }
- output[it] = nonce.y;
- output[MAXBUFFERS] = 1;
+ output[MAXBUFFERS] = output[NFLAG & nonce.y] = nonce.y;
}
#ifdef VECTORS4
if (Vals[7].z == 0)
{
- it += 1024;
- if (output[it]) {
- for (it = 0; it < MAXBUFFERS; it++) {
- if (!output[it])
- break;
- }
- }
- output[it] = nonce.z;
- output[MAXBUFFERS] = 1;
+ output[MAXBUFFERS] = output[NFLAG & nonce.z] = nonce.z;
}
if (Vals[7].w == 0)
{
- it += 1536;
- if (output[it]) {
- for (it = 0; it < MAXBUFFERS; it++) {
- if (!output[it])
- break;
- }
- }
- output[it] = nonce.w;
- output[MAXBUFFERS] = 1;
+ output[MAXBUFFERS] = output[NFLAG & nonce.w] = nonce.w;
}
#endif
#else
if (Vals[7] == 0)
{
- if (output[it]) {
- for (it = 0; it < MAXBUFFERS; it++) {
- if (!output[it])
- break;
- }
- }
- output[it] = nonce;
- output[MAXBUFFERS] = 1;
+ output[MAXBUFFERS] = output[NFLAG & nonce] = nonce;
}
#endif
}
\ No newline at end of file