diff --git a/cnf.h b/cnf.h index 1c63bd3..bd85b57 100644 --- a/cnf.h +++ b/cnf.h @@ -3,6 +3,7 @@ #include #include #include +/* #define CHECK(X, Y) if (X == NULL) { \ @@ -31,3 +32,4 @@ void sortlastnum(cnf* c, u64 N); void freecnf(cnf* c); +*/ \ No newline at end of file diff --git a/gpusolver.c b/gpusolver.c index 45a4e7a..1709c14 100644 --- a/gpusolver.c +++ b/gpusolver.c @@ -1,9 +1,11 @@ #include "gpusolver.h" #include #include "time.h" +#include "gmp.h" -#define GLOBAL_SIZE (256) -#define LOCAL_SIZE (GLOBAL_SIZE) + +#define LOCAL_SIZE (128) +#define GLOBAL_SIZE (1024) #define CHECKASGN (true) @@ -28,7 +30,7 @@ i32 gpusolve(cnf* c) { source_size = fread( source_str, 1, 0x100000, fp); fclose( fp ); - u32 wordcnt = 1 + ((c->varcnt) >> 5U); + u32 wordcnt = 1 + ((c->cnts[0]) >> 5U); u32* solution = calloc((wordcnt + 1), sizeof(u32)); if (solution == NULL) { @@ -36,6 +38,13 @@ i32 gpusolve(cnf* c) { exit(1); } + mpz_t gmpmax; + mpz_init(gmpmax); + mpz_ui_pow_ui(gmpmax, 2, c->cnts[0]); + mpz_div_ui(gmpmax, gmpmax, GLOBAL_SIZE); + mpz_export(solution + 1, NULL, -1, sizeof(u32), 0, 0, gmpmax); + mpz_clear(gmpmax); + cl_int res = clGetPlatformIDs(1, &platformid, &numplatforms); if (res != CL_SUCCESS) { printf("Failed to retrieve OpenCL platform IDs\n"); @@ -74,38 +83,27 @@ i32 gpusolve(cnf* c) { */ // TODO: Look into DMA, maybe? Could do clause learning CPU-side and just update the GPU buffer - cl_mem gpuheader = clCreateBuffer(context, CL_MEM_READ_ONLY, 3 * sizeof(cl_uint), NULL, &res); + cl_mem gpuheader = clCreateBuffer(context, CL_MEM_READ_ONLY, 2 * sizeof(cl_uint), NULL, &res); if (res != CL_SUCCESS) { printf("Failed to create CNF header buffer\n"); exit(1); } - cl_mem gpulvars = clCreateBuffer(context, CL_MEM_READ_ONLY, c->clausecnt * sizeof(cl_uint), NULL, &res); + cl_mem gpulvars = clCreateBuffer(context, CL_MEM_READ_ONLY, 3 * c->cnts[1] * sizeof(cl_uint), NULL, &res); if (res != CL_SUCCESS) { printf("Failed to create CNF lvar buffer\n"); exit(1); } - cl_mem gpuvariables = clCreateBuffer(context, CL_MEM_READ_ONLY, c->litcnt * sizeof(cl_uint), NULL, &res); + cl_mem gpuvariables = clCreateBuffer(context, CL_MEM_READ_ONLY, c->cnts[2] * sizeof(cl_uint), NULL, &res); if (res != CL_SUCCESS) { printf("Failed to create CNF variable buffer\n"); exit(1); } - cl_mem gpuclauses = clCreateBuffer(context, CL_MEM_READ_ONLY, c->litcnt * sizeof(cl_uint), NULL, &res); - if (res != CL_SUCCESS) { - printf("Failed to create CNF clause buffer\n"); - exit(1); - } - cl_mem gpuparities = clCreateBuffer(context, CL_MEM_READ_ONLY, c->litcnt * sizeof(cl_uchar), NULL, &res); + cl_mem gpuparities = clCreateBuffer(context, CL_MEM_READ_ONLY, c->cnts[2] * sizeof(cl_uchar), NULL, &res); if (res != CL_SUCCESS) { printf("Failed to create CNF parity buffer\n"); exit(1); } - // Allocate scratchpad memory - cl_mem gpuscratchpad = clCreateBuffer(context, CL_MEM_READ_WRITE, c->clausecnt * sizeof(cl_uchar), NULL, &res); - if (res != CL_SUCCESS) { - printf("Failed to create CNF subsumption scratchpad buffer\n"); - exit(1); - } /* cl_mem gpumaxvals = clCreateBuffer(context, CL_MEM_READ_WRITE, GLOBAL_SIZE * sizeof(cl_uint), NULL, &res); if (res != CL_SUCCESS) { @@ -120,30 +118,29 @@ i32 gpusolve(cnf* c) { exit(1); } - u32 cnfheader[3] = { c->litcnt, c->varcnt, c->clausecnt }; - // Load buffers to GPU - res = clEnqueueWriteBuffer(commqueue, gpuheader, CL_TRUE, 0, 3 * sizeof(cl_uint), cnfheader, 0, NULL, NULL); + res = clEnqueueWriteBuffer(commqueue, gpuheader, CL_TRUE, 0, 2 * sizeof(cl_uint), c->cnts, 0, NULL, NULL); if (res != CL_SUCCESS) { printf("Failed to queue CNF header write\n"); exit(1); } - res = clEnqueueWriteBuffer(commqueue, gpulvars, CL_TRUE, 0, c->clausecnt * sizeof(cl_uint), c->lastvars, 0, NULL, NULL); + res = clEnqueueWriteBuffer(commqueue, gpulvars, CL_TRUE, 0, 3 * c->cnts[1] * sizeof(cl_uint), c->clausedat, 0, NULL, NULL); if (res != CL_SUCCESS) { printf("Failed to queue CNF lvar write\n"); exit(1); } - res = clEnqueueWriteBuffer(commqueue, gpuvariables, CL_TRUE, 0, c->litcnt * sizeof(cl_uint), c->variables, 0, NULL, NULL); + res = clEnqueueWriteBuffer(commqueue, gpuvariables, CL_TRUE, 0, c->cnts[2] * sizeof(cl_uint), c->variables, 0, NULL, NULL); if (res != CL_SUCCESS) { printf("Failed to queue CNF variable write\n"); exit(1); } - res = clEnqueueWriteBuffer(commqueue, gpuclauses, CL_TRUE, 0, c->litcnt * sizeof(cl_uint), c->clauses, 0, NULL, NULL); + res = clEnqueueWriteBuffer(commqueue, gpuparities, CL_TRUE, 0, c->cnts[2] * sizeof(cl_uchar), c->parities, 0, NULL, NULL); if (res != CL_SUCCESS) { - printf("Failed to queue CNF clause write\n"); + printf("Failed to queue CNF parity write\n"); exit(1); } - res = clEnqueueWriteBuffer(commqueue, gpuparities, CL_TRUE, 0, c->litcnt * sizeof(cl_uchar), c->pars, 0, NULL, NULL); + + res = clEnqueueWriteBuffer(commqueue, gpuoutput, CL_TRUE, 0, (wordcnt + 1) * sizeof(cl_uint), solution, 0, NULL, NULL); if (res != CL_SUCCESS) { printf("Failed to queue CNF parity write\n"); exit(1); @@ -179,18 +176,16 @@ i32 gpusolve(cnf* c) { size_t maxworkgrpu = 0; res = clGetKernelWorkGroupInfo(kernel, deviceid, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &maxworkgrpu, NULL); - printf("Max work group size: %lu\n", maxworkgrpu); + // printf("Max work group size: %lu\n", maxworkgrpu); res = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &gpuheader); res = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &gpulvars); res = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &gpuvariables); - res = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void*) &gpuclauses); - res = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void*) &gpuparities); + res = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void*) &gpuparities); - res = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void*) &gpuoutput); + res = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void*) &gpuoutput); - res = clSetKernelArg(kernel, 6, sizeof(cl_mem), (void*) &gpuscratchpad); - res = clSetKernelArg(kernel, 7, LOCAL_SIZE * sizeof(cl_uint), NULL); + res = clSetKernelArg(kernel, 5, 2 * wordcnt * sizeof(cl_uint) * LOCAL_SIZE, NULL); // u64 starttime = utime(); size_t itemsize[2] = {GLOBAL_SIZE, LOCAL_SIZE }; @@ -207,34 +202,39 @@ i32 gpusolve(cnf* c) { } // u64 endtime = utime(); - if (solution[0] == 1) { + if (solution[0] == 0) { printf("UNSAT\n"); - } else if (solution[0] == 0) { - printf("SAT\n"); - for (u32 k = 0; k < c->varcnt; ++k) { - u32 vind = (c->varcnt - 1) - k; + } else if (solution[0] == 1) { + printf("SAT: "); + for (u32 k = 0; k < c->cnts[0]; ++k) { + u32 vind = (c->cnts[0] - 1) - k; u32 iind = vind >> 5U; u32 bind = vind & 0b11111U; u8 par = (solution[iind + 1] >> bind) & 1U; printf("%u", par); } - printf("\n"); if (CHECKASGN) { - u8* assigncheck = calloc(c->clausecnt, sizeof(u8)); - for (u32 i = 0; i < c->litcnt; ++i) { - u32 g = ((c->varcnt - 1) - c->variables[i]) >> 5U; - u32 h = ((c->varcnt - 1) - c->variables[i]) & 0b11111U; - u8 paract = (solution[g + 1] >> h) & 1U; - if (c->pars[i] == paract) assigncheck[c->clauses[i]] = true; - } - for (u32 i = 0; i < c->clausecnt; ++i) { - if (!assigncheck[i]) { - printf("Failed assignment check\n"); - solution[0] = 4; + u8 checkres = 0; + for (u32 i = 0; i < c->cnts[1]; ++i) { + checkres = 0; + for (u32 j = 0; j < c->clausedat[3 * i + 1]; ++j) { + u32 v = c->variables[c->clausedat[3 * i] + j]; + u32 vv = c->cnts[0] - 1; + u32 g = (vv - v) >> 5U; + u32 h = (vv - v) & 0b11111U; + u8 paract = (solution[g + 1] >> h) & 1U; + if (c->parities[c->clausedat[3 * i] + j] == paract) { + checkres = 1; + break; + } } + if (!checkres) break; + } + if (checkres) { + printf(" \xE2\x9C\x93\n"); + } else { + printf(" -\n"); } - free(assigncheck); - printf("Passed assignment check\n"); } } else { printf("What the fuck???\n"); @@ -253,16 +253,14 @@ i32 gpusolve(cnf* c) { res = clReleaseMemObject(gpuheader); res = clReleaseMemObject(gpulvars); res = clReleaseMemObject(gpuvariables); - res = clReleaseMemObject(gpuclauses); res = clReleaseMemObject(gpuparities); res = clReleaseMemObject(gpuoutput); - res = clReleaseMemObject(gpuscratchpad); res = clReleaseDevice(deviceid); - i32 retval = solution[0]; + i32 retval = (i32) solution[0]; free(solution); free(source_str); return retval; -} \ No newline at end of file +} diff --git a/gpusolver.h b/gpusolver.h index 65e3728..a1d91a1 100644 --- a/gpusolver.h +++ b/gpusolver.h @@ -1,4 +1,4 @@ #pragma once -#include "cnf.h" +#include "ncnf.h" i32 gpusolve(cnf* c); \ No newline at end of file diff --git a/main.c b/main.c index 455956b..21663ee 100644 --- a/main.c +++ b/main.c @@ -1,7 +1,7 @@ #include -// #include "gpusolver.h" +#include "gpusolver.h" #include "time.h" -// #include "tests/masterTest.h" +#include "tests/masterTest.h" #include "gmp.h" #include "rng.h" #include "ncnf.h" @@ -10,8 +10,9 @@ #define CMP (1) #define CHK (2) + void ctrthings2(cnf* c, u32* state, u32* ctr, u32* max) { - u32 wcnt = 1U + (c->cnts[0] / 32U); + u32 wcnt = 1U + (c->cnts[0] >> 5U); u32* mode = state; u32* index = state + 1; u32* addval = state + 2; @@ -27,101 +28,40 @@ void ctrthings2(cnf* c, u32* state, u32* ctr, u32* max) { u8 corpar = (ctr[vword] >> vbit) & 1U; u8 isvalid = (par == corpar); u8 islvar = ((*addval + 1) == c->clausedat[3 * chkcls + 1]); - if (*mode == CHK) { - // printf("> %u %u\n", *index, *addval); - // printf("%u %u %u\n", var, par, corpar); - // printf("%u %u\n", chkcls, chkind); - // printf("%u %u %u\n", islvar, isvalid, c->clausedat[3 * chkcls + 1]); - - /* - * if last var - * if valid, add - * if invalid, iterate clause - * else - * if valid, iterate addval up to len - * if invalid, iterate claus - * If current var is valid: - * if last var: - * - */ - if (isvalid) { - if (islvar) { - u32 j = c->clausedat[3 * chkcls + 2]; - *mode -= 2; - printf("j: %u\n", j); - *index = j >> 5U; - *addval = 1U << (j & 0b11111U); - } else { - *addval += 1U; - } - } else { - *addval = 0; - *index += 1U; - if (*index == c->cnts[1]) { - printf("SAT\n"); - *mode = 4; - return; - } - } - } /* else { - - //printf("YEET\n"); - if (*index >= wcnt) printf("FUCK\n"); - u32 nval = ctr[*index] + *addval; // Find the result of the current step if it was addition - *addval = (nval < ctr[*index]) * (*mode == ADD); // If in add mode, set addval to carry, else set 0 - ctr[*index] = nval * (*mode == ADD) + ctr[*index] * (*mode != ADD); // If in add mode, set new ctr val, otherwise leave unchanged - *addval -= (ctr[*index] < max[*index]) * (*mode == CMP); // If in comparison mode, decrement addval if less than - *addval += (ctr[*index] > max[*index]) * (*mode == CMP); // If in comparison mode, increment addval if greater than - bool addcond = (*addval == 0) | (*index == (wcnt - 1)); // Exit condition for the ADD state: If addval is zero (no carry) or we're at the last word - bool cmpcond = (*addval != 0) | (*index == 0); // Exit condition for the CMP state: if addval is nonzero (lt or gt) or we're at the least significant word - if (*mode == CMP && cmpcond && *addval != -1) { // If in cmpmode and the comparison result is not less than, unsat - printf("UNSAT\n"); - *mode = 4; - return; - } - bool cmpdone = cmpcond * (*mode == CMP); // if comparison completion conditions are satisfied and in CMP mode - u32 addindex = (*index + 1) * !addcond + (wcnt - 1) * addcond; // if add completion is satisfied, set index to most significant word, else increment by 1 - *index = addindex * (*mode == ADD) + (*index - 1) * (*mode == CMP); // If in add mode, use addindex; if in cmp mode, decrement index by 1 - *index *= !cmpdone; - // Leave adval alone if: - // not in add mode - // add mode isn't done - // not in cmp mode - // cmp mode isn't done - // - *addval *= !(((addcond) & (*mode != ADD)) & cmpdone); // If add is complete, zero addval, else leave unchanged - *mode += addcond * (*mode == ADD) + cmpdone; // If in add mode and add completion is reached, increment mode. If in cmp mode and cmp completion reached, increment mode. - - } - */ + u8 isbchk0 = (*mode == CHK); + u8 isbchk1 = isbchk0 & isvalid; + u8 isbchk2 = isbchk1 & islvar; + u32 j = c->clausedat[3 * chkcls + 2]; + *mode -= 2 * isbchk2; + *index = (j >> 5U) * isbchk2 + *index * (!isbchk2); + *addval = (1U << (j & 0b11111U)) * isbchk2 + *addval * (!isbchk2); + *addval += ((isbchk1) & (!islvar)); + u8 isbchk3 = (isbchk0 & (!isvalid)); + *addval *= (!isbchk3); + *index += (isbchk3); + u8 issat = (*index == c->cnts[1]) * (isbchk3); u32 cmpaddind = *index * (*mode != CHK); - if (cmpaddind >= wcnt) printf("FUCK\n"); u32 nval = ctr[cmpaddind] + *addval; // Find the result of the current step if it was addition *addval = (nval < ctr[cmpaddind]) * (*mode == ADD) + (*addval) * (*mode == CHK); // If in add mode, set addval to carry. If in cmp mode, set to 0. If in check mode, leave alone. - ctr[cmpaddind] = nval * (*mode == ADD) + ctr[cmpaddind] * (*mode != ADD); // If in add mode, set new ctr val, otherwise leave unchanged + ctr[cmpaddind] = nval * ((*mode == ADD) & !issat) + ctr[cmpaddind] * ((*mode != ADD) | issat); // If in add mode, set new ctr val, otherwise leave unchanged *addval -= (ctr[cmpaddind] < max[cmpaddind]) * (*mode == CMP); // If in comparison mode, decrement addval if less than *addval += (ctr[cmpaddind] > max[cmpaddind]) * (*mode == CMP); // If in comparison mode, increment addval if greater than - bool addcond = (*addval == 0) | (cmpaddind == (wcnt - 1)); // Exit condition for the ADD state: If addval is zero (no carry) or we're at the last word - bool cmpcond = (*addval != 0) | (cmpaddind == 0); // Exit condition for the CMP state: if addval is nonzero (lt or gt) or we're at the least significant word - if (*mode == CMP && cmpcond && *addval != -1) { // If in cmpmode and the comparison result is not less than, unsat - printf("UNSAT\n"); + u8 addcond = (*addval == 0) | (cmpaddind == (wcnt - 1)); // Exit condition for the ADD state: If addval is zero (no carry) or we're at the last word + u8 cmpcond = (*addval != 0) | (cmpaddind == 0); // Exit condition for the CMP state: if addval is nonzero (lt or gt) or we're at the least significant word + u8 exittime = (*mode == CMP) & cmpcond & (*addval != -1); + exittime |= issat; + if (exittime) { // If in cmpmode and the comparison result is not less than, unsat + printf("Result: %u\n", issat); *mode = 4; return; } - bool cmpdone = cmpcond & (*mode == CMP); // if comparison completion conditions are satisfied and in CMP mode + u8 cmpdone = cmpcond & (*mode == CMP); // if comparison completion conditions are satisfied and in CMP mode u32 addindex = (cmpaddind + 1) * !addcond + (wcnt - 1) * addcond; // if add completion is satisfied, set index to most significant word, else increment by 1 *index = addindex * (*mode == ADD) + (*index - (*mode == CMP)) * (*mode != ADD); // If in add mode, use addindex; if in cmp mode, decrement index by 1 *index *= !cmpdone; - // Leave adval alone if: - // not in add mode - // add mode isn't done - // not in cmp mode - // cmp mode isn't done - // *addval *= !(((addcond) & (*mode == ADD)) | cmpdone); // If add is complete, or cmp is complete, zero. Else leave unchanged. *mode += addcond * (*mode == ADD) + cmpdone; // If in add mode and add completion is reached, increment mode. If in cmp mode and cmp completion reached, increment mode. - } void printbits(unsigned a) { @@ -134,6 +74,38 @@ void printbits(unsigned a) { #define TESTS (274877906944LU >> 10U) #define CSZE (83LU) #define eqprob (0.01f) + + +void mul(u32* c, u32 len, u32* a, u32 b) { + u32 carry = 0; + for (u32 i = 0; i < len; ++i) { + u32 ncarry; + u32 blo = a[i] & 0xFFFFU; + u32 bhi = a[i] >> 16U; + u32 ilo = b & 0xFFFFU; + u32 ihi = b >> 16U; + + *(c + i) = ilo * blo; + u32 b1 = ilo * bhi; + u32 c1 = ihi * blo; + ncarry = ihi * bhi; + + b1 += c1; + ncarry += (b1 < c1) << 16U; + u32 bblo = b1 & 0xFFFFU; + u32 bbhi = b1 >> 16U; + bblo <<= 16U; + *(c + i) += bblo; + u8 acarry = *(c + i) < bblo; + + ncarry += bbhi + acarry; + c[i] += carry; + ncarry += c[i] < carry; + carry = ncarry; + } +} + + int main() { /* printf("Tests: %lu\n", TESTS); @@ -276,38 +248,20 @@ int main() { /* Expects a path to a DIMACS file */ - cnf* c = readDIMACS("/home/lev/Downloads/uf20/uf20-022.cnf"); + + + /* + cnf* c = readDIMACS("/home/lev/Downloads/logistics/logistics.d.cnf"); sortlastnum(c); - // printcnf(c); - u32 wcnt = 1U + (c->cnts[0] / 32U); + printf("%u\n", c->cnts[0]); + gpusolve(c); - u32* ctr = calloc(wcnt, sizeof(u32)); - u32* max = calloc(wcnt, sizeof(u32)); - - max[c->cnts[0] >> 5U] = 1U << (c->cnts[0] & 0b11111U); - - u32 state[3]; - - state[0] = 2; - state[1] = state[2] = 0; - - - u32 mtr = 0; - while (state[0] < 3U) { - u32 cmd = state[0]; - ctrthings2(c, state, ctr, max); - if (state[0] != 2 && cmd == 2) { - //printf("\n"); - //for (unsigned i = wcnt - 1; i < wcnt; --i) printbits(ctr[i]); - //printf("\n"); - //mtr++; - //if (mtr == 10) exit(15); - // printf("%u %u %u\n", state[0], state[1], state[2]); - } - } - + freecnf(c); + return 0; + */ + runTests(); return 0; /* diff --git a/ncnf.c b/ncnf.c index 37b88b4..e9e8f05 100644 --- a/ncnf.c +++ b/ncnf.c @@ -127,6 +127,7 @@ cnf* readDIMACS(char* path) { // Realloc the arrays to exactly match the number of literals c->variables = realloc(c->variables, sizeof(u32) * cnt); c->parities = realloc(c->parities, sizeof(u8) * cnt); + c->cnts[2] = cnt; free(buf); if (fclose(f)) { printf("Failed to close file\n"); diff --git a/ncnf.h b/ncnf.h index c61d80f..6e84f26 100644 --- a/ncnf.h +++ b/ncnf.h @@ -10,7 +10,7 @@ } typedef struct { - u32 cnts[2]; // { varcnt, clausecnt } + u32 cnts[3]; // { varcnt, clausecnt } u32* clausedat; // { ind, len, jval } u32* variables; u8* parities; diff --git a/psat.cl b/psat.cl index c29dcff..13fd003 100644 --- a/psat.cl +++ b/psat.cl @@ -32,120 +32,122 @@ static inline void stateaddpow(uint wcnt, uint* state, uint pow) { } } -__kernel void vectorSAT(__global const uint* cnfheader, __global const uint* lvars, __global const uint* vars, __global const uint* clauses, __global const uchar* pars, __global uint* output, __global uchar* scratchpad, __local uint* maxvals) { - output[0] = 2; +void mul(uint* c, uint len, uint* a, uint b) { + uint carry = 0; + for (uint i = 0; i < len; ++i) { + uint ncarry; + uint blo = a[i] & 0xFFFFU; + uint bhi = a[i] >> 16U; + uint ilo = b & 0xFFFFU; + uint ihi = b >> 16U; - __local uint setmax; + *(c + i) = ilo * blo; + uint b1 = ilo * bhi; + uint c1 = ihi * blo; + ncarry = ihi * bhi; - uint cnt = cnfheader[0]; - uint vcnt = cnfheader[1]; - uint ccnt = cnfheader[2]; - - uint wcnt = 1 + (vcnt >> 5U); - - uint maxctr = 1U << (vcnt & 0b11111U); - - //uint glbid = get_global_id(0); - //uint glbsz = get_global_size(0); + b1 += c1; + ncarry += (b1 < c1) << 16U; + uint bblo = b1 & 0xFFFFU; + uint bbhi = b1 >> 16U; + bblo <<= 16U; + *(c + i) += bblo; + uchar acarry = *(c + i) < bblo; + ncarry += bbhi + acarry; + c[i] += carry; + ncarry += c[i] < carry; + carry = ncarry; + } +} +__kernel void vectorSAT(__global const uint* cnfhdr, __global const uint* clausedat, __global const uint* vars, __global const uchar* pars, __global uint* output, __local uint* lctrs) { uint locid = get_local_id(0); uint locsz = get_local_size(0); // uint grpid = get_group_id(0); // uint grpcn = get_num_groups(0); + uint globid = get_global_id(0); + uint globsz = get_global_size(0); - // Zero out the counter - for (uint i = 0; i < wcnt; ++i) output[i + 1] = 0; + uint wcnt = 1U + (cnfhdr[0] >> 5U); - // Set all scratchpad clauses to true - for (uint j = 0; j < ccnt; j += locsz) { - uchar cond = (j + locid) < ccnt; - j = j * cond + (!cond) * (ccnt - locid - 1); - scratchpad[j + locid] = 1; + uint mode = 2; + uint index = 0; + uint addval = 0; + + + if (globid == 0) output[0] = 0; + + uint* ctr = lctrs + wcnt * 2 * locid; + uint* max = lctrs + wcnt * (2 * locid + 1); + + for (uint i = 0; i < wcnt; ++i) { + ctr[i] = max[i] = 0; } - __local uint firstind[1]; - while (output[0] == 2) { - firstind[0] = ccnt; - - setmax = 0; - uint maxnumx = 0; - - for (uint j = 0; j < cnt; j += locsz) { - - uchar cond = (j + locid) < cnt; - // Last element cap - j = j * cond + (!cond) * (cnt - locid - 1); - uint varind = vars[j + locid]; - varind = (vcnt - 1) - varind; - uint iind = varind >> 5U; - uint bind = varind & 0b11111U; - uchar cpar = (output[iind + 1] >> bind) & 1U; - - if (cpar != pars[j + locid]) { - scratchpad[clauses[j + locid]] = 0; - } - } - - barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); + mul(ctr, wcnt, output + 1, globid); + if (globid == globsz) { + stateaddpow(wcnt, max, cnfhdr[0]); + } else { + mul(max, wcnt, output + 1, globid + 1); + } - for (uint j = 0; j < ccnt; j += locsz) { - if (scratchpad[j + locid] == 1 && (j + locid) < ccnt) { - setmax = 1; - } - } + barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); - barrier(CLK_LOCAL_MEM_FENCE); + uint varcnt = cnfhdr[0] - 1; + while (output[0] == 0) { - if (setmax) { - // Set maxval array to zero - maxvals[locid] = 0; + uint chkmsk = 0xFFFFFFFFU * (mode == 2U); + uint chkcls = index & chkmsk; + uint chkind = clausedat[3 * chkcls] + (addval & chkmsk); + uint var = vars[chkind]; + uchar par = pars[chkind]; + uint vword = (varcnt - var) >> 5U; + uint vbit = (varcnt - var) & 0b11111U; + uchar corpar = (ctr[vword] >> vbit) & 1U; + uchar isvalid = (par == corpar); + uchar islvar = ((addval + 1) == clausedat[3 * chkcls + 1]); + uchar isbchk0 = (mode == 2U); + uchar isbchk1 = isbchk0 & isvalid; + uchar isbchk2 = isbchk1 & islvar; + uint j = clausedat[3 * chkcls + 2]; + mode -= 2 * isbchk2; + // if (isbchk2) printf("j: %u\n", j); + index = (j >> 5U) * isbchk2 + index * (!isbchk2); + addval = (1U << (j & 0b11111U)) * isbchk2 + addval * (!isbchk2); + addval += ((isbchk1) & (!islvar)); + uchar isbchk3 = (isbchk0 & (!isvalid)); + addval *= (!isbchk3); + index += (isbchk3); + uchar issat = (index == cnfhdr[1]) * (isbchk3); - // Accumulate and reduce the maximums - for (uint j = 0; j < ccnt; j += locsz) { - //uint a = maxvals[locid]; - //uint b = lvars[j + locid]; - // uint c = max(a, b); - if ((j + locid) < ccnt && scratchpad[j + locid] == 1) { - //maxvals[locid] = c; - atomic_min(firstind, (j + locid)); + uint cmpaddind = index * (mode != 2U); + uint nval = ctr[cmpaddind] + addval; // Find the result of the current step if it was addition + addval = (nval < ctr[cmpaddind]) * (mode == 0) + (addval) * (mode == 2U); // If in add mode, set addval to carry. If in cmp mode, set to 0. If in check mode, leave alone. + ctr[cmpaddind] = nval * ((mode == 0) & !issat) + ctr[cmpaddind] * ((mode != 0) | issat); // If in add mode, set new ctr val, otherwise leave unchanged + addval -= (ctr[cmpaddind] < max[cmpaddind]) * (mode == 1); // If in comparison mode, decrement addval if less than + addval += (ctr[cmpaddind] > max[cmpaddind]) * (mode == 1); // If in comparison mode, increment addval if greater than + uchar addcond = (addval == 0) | (cmpaddind == (wcnt - 1)); // Exit condition for the ADD state: If addval is zero (no carry) or we're at the last word + uchar cmpcond = (addval != 0) | (cmpaddind == 0); // Exit condition for the CMP state: if addval is nonzero (lt or gt) or we're at the least significant word + uchar exittime = (mode == 1) & cmpcond & (addval != -1); + exittime |= issat; + if (exittime) { // If in cmpmode and the comparison result is not less than, unsat + if (issat) { + if (atomic_cmpxchg(output, 0, 1) == 0) { + for (uint i = 0; i < wcnt; ++i) { + output[i + 1] = ~ctr[i]; + } + output[0] = 1; } } - - barrier(CLK_LOCAL_MEM_FENCE); - - uint maxj = lvars[firstind[0]]; - - // Set all scratchpad clauses to true - for (uint j = 0; j < ccnt; j += locsz) { - uchar cond = (j + locid) < ccnt; - j = j * cond + (!cond) * (ccnt - locid - 1); - scratchpad[j + locid] = 1; - } - - // Final reduction pass - /* - uint maxj = maxvals[0]; - for (uint j = 1; j < locsz; ++j) { - maxj = max(maxj, maxvals[j]); - } - */ - - // Add to the counter - if (locid == 0) { - stateaddpow(wcnt, output + 1, maxj); - } - - if (output[wcnt] >= maxctr) { - output[0] = 1; - } - } else { - output[0] = 0; - if (locid == 0) { - for (uint i = 0; i < wcnt; ++i) output[i + 1] = ~output[i + 1]; - } + return; } - barrier(CLK_LOCAL_MEM_FENCE); + uchar cmpdone = cmpcond & (mode == 1); // if comparison completion conditions are satisfied and in CMP mode + uint addindex = (cmpaddind + 1) * !addcond + (wcnt - 1) * addcond; // if add completion is satisfied, set index to most significant word, else increment by 1 + index = addindex * (mode == 0) + (index - (mode == 1)) * (mode != 0); // If in add mode, use addindex; if in cmp mode, decrement index by 1 + index *= !cmpdone; + addval *= !(((addcond) & (mode == 0)) | cmpdone); // If add is complete, or cmp is complete, zero. Else leave unchanged. + mode += addcond * (mode == 0) + cmpdone; // If in add mode and add completion is reached, increment mode. If in cmp mode and cmp completion reached, increment mode. } } \ No newline at end of file diff --git a/tests/masterTest.c b/tests/masterTest.c index 471d369..db87a35 100644 --- a/tests/masterTest.c +++ b/tests/masterTest.c @@ -24,7 +24,7 @@ i32 runTests() { i32 runuf20() { - printf("Running against uf20\n"); + // printf("Running against uf20\n"); u32 passed = 0; u64 tottime = 0; for (u32 i = 0; i < 1000; ++i) { @@ -33,62 +33,7 @@ i32 runuf20() { cnf* c = readDIMACS(buf); - // TODO: Uncomment - // sortlastnum(c, c->litcnt); - - u64 start = utime(); - i32 res = gpusolve(c); - u64 stop = utime(); - tottime += (stop - start); - - freecnf(c); - if (res == 0) passed++; - } - printf("Passed %u / 1000 tests\n", passed); - printf("Took %lf s total, %lf s on avg\n", ((f64) tottime) / 1000000.0, ((f64) tottime) / 1000000000.0); - if (passed == 1000) return 0; - return 1; -} - -i32 runuf50() { - printf("Running against uf50\n"); - u32 passed = 0; - u64 tottime = 0; - for (u32 i = 0; i < 1000; ++i) { - char buf[128]; - i32 len = sprintf(buf, "/home/lev/Downloads/uf50/uf50-0%u.cnf", i + 1); - - cnf* c = readDIMACS(buf); - - // TODO: Uncomment - // sortlastnum(c, c->litcnt); - - u64 start = utime(); - i32 res = gpusolve(c); - u64 stop = utime(); - tottime += (stop - start); - - freecnf(c); - if (res == 0) passed++; - } - printf("Passed %u / 1000 tests\n", passed); - printf("Took %lf s total, %lf s on avg\n", ((f64) tottime) / 1000000.0, ((f64) tottime) / 1000000000.0); - if (passed == 1000) return 0; - return 1; -} - -i32 runuuf50() { - printf("Running against uuf50\n"); - u32 passed = 0; - u64 tottime = 0; - for (u32 i = 0; i < 1000; ++i) { - char buf[128]; - i32 len = sprintf(buf, "/home/lev/Downloads/uuf50/uuf50-0%u.cnf", i + 1); - - cnf* c = readDIMACS(buf); - - // TODO: Uncomment - // sortlastnum(c, c->litcnt); + sortlastnum(c); u64 start = utime(); i32 res = gpusolve(c); @@ -98,8 +43,60 @@ i32 runuuf50() { freecnf(c); if (res == 1) passed++; } - printf("Passed %u / 1000 tests\n", passed); - printf("Took %lf s total, %lf s on avg\n", ((f64) tottime) / 1000000.0, ((f64) tottime) / 1000000000.0); + // printf("Passed %u / 1000 tests\n", passed); + // printf("Took %lf s total, %lf s on avg\n", ((f64) tottime) / 1000000.0, ((f64) tottime) / 1000000000.0); + if (passed == 1000) return 0; + return 1; +} + +i32 runuf50() { + // printf("Running against uf50\n"); + u32 passed = 0; + u64 tottime = 0; + for (u32 i = 0; i < 1000; ++i) { + char buf[128]; + i32 len = sprintf(buf, "/home/lev/Downloads/uf50/uf50-0%u.cnf", i + 1); + + cnf* c = readDIMACS(buf); + + sortlastnum(c); + + u64 start = utime(); + i32 res = gpusolve(c); + u64 stop = utime(); + tottime += (stop - start); + + freecnf(c); + if (res == 1) passed++; + } + // printf("Passed %u / 1000 tests\n", passed); + // printf("Took %lf s total, %lf s on avg\n", ((f64) tottime) / 1000000.0, ((f64) tottime) / 1000000000.0); + if (passed == 1000) return 0; + return 1; +} + +i32 runuuf50() { + // printf("Running against uuf50\n"); + u32 passed = 0; + u64 tottime = 0; + for (u32 i = 0; i < 1000; ++i) { + char buf[128]; + i32 len = sprintf(buf, "/home/lev/Downloads/uuf50/uuf50-0%u.cnf", i + 1); + + cnf* c = readDIMACS(buf); + + sortlastnum(c); + + u64 start = utime(); + i32 res = gpusolve(c); + u64 stop = utime(); + tottime += (stop - start); + + freecnf(c); + if (res == 0) passed++; + } + // printf("Passed %u / 1000 tests\n", passed); + // printf("Took %lf s total, %lf s on avg\n", ((f64) tottime) / 1000000.0, ((f64) tottime) / 1000000000.0); if (passed == 1000) return 0; return 1; } \ No newline at end of file