Comment by dragontamer | Hacker News Reader

Comment by dragontamer | Hacker News Reader

R49081 means this was a ~49081 digit number, or a number that could be represented in maybe 163043-ish bits or ~20kB or so?

I'm kinda curious if this is a data-size that might be better on CPU rather than GPU. GPUs of course are better compute, but GPU-registers tend to cap out at ~1024 bytes or so, and a number of this size probably would have to be stored in RAM (plus all the associated data structures in ... whatever algorithm happened here).

If this were small enough, I can kinda imagine CPUs being better due to caching effects (L1 and L2 are incredibly fast on CPU, and GPUs kind of miss those tiers in some sense since GPUs focus on register space).

It could just be that this eliptical-curve prime-testing program is currently a CPU kinda function. But I'm curious if this algorithm would be (or wouldn't be) suitable for GPU compute, with more TFLOPS and more VRAM bandwidth for cheaper.

EDIT: A 20kB number would fit inside of __shared__ memory on GPU, so perhaps various operations (multiplication and/or division??) could be parallelized onto a 256-wide/1024-wide workgroup (~1024 to 4096 bytes per clock tick, or ~5 to ~20 clock ticks to perform most kinds of operations with the 20kB number on one compute unit). It really depends on the kind of algorithm going on here, as well as what needs to be searched for. I did a brief read into ECPP testing, and I honestly don't get anything they're saying in the paper (A. O. L. ATKIN AND F. MORAIN, 1993. "ELLIPTIC CURVES AND PRIMALITY PROVING").

For the mersenne project (searching for prime numbers that are 110 million bits in size) GPUs are quite efficient.

One of the algorithms used, called PRP (probable prime), which uses the reciprocal of Fermat's little theorem (a^(p-1)==1 mod p if p is prime) involves computing 3^p mod p (where p is a 110million bits number). This is called "modular exponentiation", and is done efficiently with FFTs (Fast Fourrier Transform, used to implement multiplication, thus squaring).

The FFT of 110Mbits can be implemented efficiently on GPUs. The key for keeping the hot data in GPU registers or caches is "locality of data access". Although the FFT algorithm is non-local by excellence (has a tendency to access all the data all the time), it can be split into "blocks" which have a smaller hot-data size. And this allows efficient GPU implementations.

The problem with R49081 may be that it's a small number, thus hard to fill all the processing units of the GPU with the amount of parallel work the algo offers at this size. Another problem may be that the algorithm involving eliptic-curves is more complex, thus more work is needed to express it in GPU terms.

http://mersenne.org/

> The problem with R49081 may be that it's a small number, thus hard to fill all the processing units of the GPU with the amount of parallel work the algo offers at this size. Another problem may be that the algorithm involving eliptic-curves is more complex, thus more work is needed to express it in GPU terms.

Well, 20kB is small, but also large. Too large to fit inside of 32-bit registers on the GPU (or even in 200+ such registers), too small to really take advantage of the GPU's much faster VRAM.

Its an interesting size. Larger primes almost certainly would benefit from a GPU no doubt. Smaller primes also would benefit from GPU (see cryptocurrency miners, where the entire program fits inside of a singular shader).

20kB though? Its... a weird size. Very interesting. If it were much bigger, or smaller, I'd be confident about porting the algorithm over to the GPU. But its actually more intimidating to be at that size (especially since 64kB L1 caches of CPUs is so darn good).

If I were forced to do this in GPU space, I'd have one-workgroup per number. I'd hope that the majority of the compute-time were spent on multiplication/division routines that would benefit from all 1024-threads (maxed size workgroup). But it'd be a more difficult program to write than a larger or smaller number.

These are "repunit" primes, meaning that it's the digit "1" repeated N times, in this case 49081 times. So yes 20KB.

47.9KiB in fact, assuming we are using 8-bits to the byte.

Where are you getting that?

log2((10^49081 - 1) / 9) / 8 / 1024 ≈ 19.9

About 20 KiB

Ah. For some reason my head was locked in binary, and even then I was forgetting to convert Kbit->KiB. My only excuse is a bad night's sleep…

Only if you're representing it as an ASCII or UTF-8 string with one byte per character, or if you're using binary-coded decimal with 8 bits instead of the typical 4 for each decimal digit.

256 (3 digits) in base 2 is one byte, 65,536 (5 digits) is two bytes, 16,777,216 (8 digits) is three bytes, 4,294,967,296 (10 digits) is four bytes, etc. Those aren't 3, 5, 8, and 10 bytes, they're 1, 2, 3, and 4.

> R49081 means this was a ~49081 digit number, or a number that could be represented in maybe 163043-ish bits or ~20kB or so?

A single (very large) number, that I could not represent on any computer I owned until Christmas morning 1983.

Mad.

Intel is offering customer silicon on their latest architecture. So you could be ambitious and design your own block with as much SRAM as you can sacrifice other accelerator units for, depending on routing latency.

Edit: customer not custom, small difference but important if you're asking..

Custom / semi-custom is probably an old Altera FPGA kinda thing, which means verilog (or similar hardware-level programming languages), which is much more difficult to use in my experience.

I generally see the "compute" hierarchy as:

1. General Purpose CPUs -- easiest.

2. CPU acceleration -- AVX512, NEON, SVE, etc. etc. Specialized CPU instructions that only exist on certain versions of hardware.

3. General Purpose GPU -- OpenCL, DirectCompute, ROCm, CUDA. General purpose GPU instructions that work on a wide variety of hardware.

4. GPU accelerated units -- Raytracing, matrix-multiplication, bfloat16 support, wave intrinsics. I guess DirectX12 Ultimate has these but you definitely need to be checking to see if your hardware supports this before using them.

5. FPGAs -- Over here baby, but much much harder than #4 in practice.

-------

I know Intel has FPGA + Intel core chips, and that Microsoft has used them before for Bing search and other such projects. But no one ever told me it was easy.

Reply on news.ycombinator.com