Faster GPU-ECM with CGBN

[SethTro]

Originally Posted by bsquared

I have run 2^997-1 up to B1=10M so far with no problems. I will go up another order of magnitude; should take a few hours.

(2^499-1)/20959 has a factor that is found on my card in stage 1 at B1=2M (sigma=3890).

You might try (2^419-1)/4398294875195008479937661267298757530097217 which has a 38 digit factor or
(2^569-1)/160592976218334727184554268072735638438202191 which has a 42 digit factor

Originally Posted by chris2be8

I got that error too:

$ echo "((2^499-1)/20959)" | ./ecm -gpu -cgbn -sigma 3:1000 20000 0
GMP-ECM 7.0.5-dev [configured with GMP 6.1.2, --enable-asm-redc, --enable-gpu, --enable-assert] [ECM]
Input number is ((2^499-1)/20959) (146 digits)
Using B1=20000, B2=0, sigma=3:1000-3:1383 (384 curves)
CUDA error occurred: cannot set while device is active in this process
While running cudaSetDeviceFlags (cudaDeviceScheduleBlockingSync)   (file cgbn_stage1.cu, line 601)

Though the different line number makes me suspect I may have an old version of your code. What command should I run to download the latest version?

You can run `git pull` (tries to pull any changes I've made but fails if I do a bad thing that I like to do). If that fails you can run this command to reset to the state of my branch (this is destructive of any changes you made)
`git fetch`
`git reset --hard origin/gpu_integration`

Thanks for confirming you also see the cudaError, I'll investigate more now that multiple people see it.

[frmky]

That was a painful configure. I'm using nvhpc rather than the CUDA toolkit, so the directory structure is different and the cuda lib is just a stub so the check for cuInit() fails. I can point to the actual lib, but the version doesn't match the nvhpc version (which is fine, nvhpc doesn't require them to match) but that check fails. Then the check for CGBN ignores the --with-gmp directory so that check fails because it can't find GMP but reports that it can't find cgbn.h. After bypassing all of those checks, I get a working binary. Yay!

$ echo "(2^499-1)/20959" | ./ecm -gpu -cgbn -sigma 3:1000 20000 0
GMP-ECM 7.0.5-dev [configured with GMP 6.2.1, --enable-asm-redc, --enable-gpu, --enable-assert] [ECM]
Input number is (2^499-1)/20959 (146 digits)
Using B1=20000, B2=0, sigma=3:1000-3:6119 (5120 curves)
Computing 5120 Step 1 took 282ms of CPU time / 705ms of GPU time

$ echo "((2^499-1)/20959)" | ./ecm -gpu -cgbn -sigma 3:1000 2000000 0
GMP-ECM 7.0.5-dev [configured with GMP 6.2.1, --enable-asm-redc, --enable-gpu, --enable-assert] [ECM]
Input number is ((2^499-1)/20959) (146 digits)
Using B1=2000000, B2=0, sigma=3:1000-3:6119 (5120 curves)
GPU: factor 1998447222711143545931606352264121 found in Step 1 with curve 2890 (-sigma 3:3890)
Computing 5120 Step 1 took 47146ms of CPU time / 112480ms of GPU time
********** Factor found in step 1: 1998447222711143545931606352264121
Found prime factor of 34 digits: 1998447222711143545931606352264121
Prime cofactor (((2^499-1)/20959))/1998447222711143545931606352264121 has 113 digits

[SethTro]

Originally Posted by frmky

That was a painful configure. I'm using nvhpc rather than the CUDA toolkit, so the directory structure is different and the cuda lib is just a stub so the check for cuInit() fails. I can point to the actual lib, but the version doesn't match the nvhpc version (which is fine, nvhpc doesn't require them to match) but that check fails. Then the check for CGBN ignores the --with-gmp directory so that check fails because it can't find GMP but reports that it can't find cgbn.h. After bypassing all of those checks, I get a working binary. Yay!

$ echo "(2^499-1)/20959" | ./ecm -gpu -cgbn -sigma 3:1000 20000 0
GMP-ECM 7.0.5-dev [configured with GMP 6.2.1, --enable-asm-redc, --enable-gpu, --enable-assert] [ECM]
Input number is (2^499-1)/20959 (146 digits)
Using B1=20000, B2=0, sigma=3:1000-3:6119 (5120 curves)
Computing 5120 Step 1 took 282ms of CPU time / 705ms of GPU time

$ echo "((2^499-1)/20959)" | ./ecm -gpu -cgbn -sigma 3:1000 2000000 0
GMP-ECM 7.0.5-dev [configured with GMP 6.2.1, --enable-asm-redc, --enable-gpu, --enable-assert] [ECM]
Input number is ((2^499-1)/20959) (146 digits)
Using B1=2000000, B2=0, sigma=3:1000-3:6119 (5120 curves)
GPU: factor 1998447222711143545931606352264121 found in Step 1 with curve 2890 (-sigma 3:3890)
Computing 5120 Step 1 took 47146ms of CPU time / 112480ms of GPU time
********** Factor found in step 1: 1998447222711143545931606352264121
Found prime factor of 34 digits: 1998447222711143545931606352264121
Prime cofactor (((2^499-1)/20959))/1998447222711143545931606352264121 has 113 digits

Glad you got a working binary! Would you mind measuring the speedup of echo "2^997-1" with -gpu vs -cgbn?

I feel for you on setup. I didn't know any automake and how to configure it took me a whole day.

If you were willing you could try making this change and seeing if the cgbn.h checks respects the --with_gmp_lib path after it.

diff --git a/acinclude.m4 b/acinclude.m4
index fbbf94df..04694003 100644
--- a/acinclude.m4
+++ b/acinclude.m4
@@ -612,7 +612,7 @@ AS_IF([test "x$enable_gpu" = "xyes" ],
#include <gmp.h>
#include <cgbn.h>
],
- [-I$cgbn_include -lgmp],
+ [-I$cgbn_include $GMPLIB],
[AC_MSG_RESULT([yes])],
[
AC_MSG_RESULT([no])
@@ -620,7 +620,7 @@ AS_IF([test "x$enable_gpu" = "xyes" ],
]
)
AC_DEFINE([HAVE_CGBN_H], [1], [Define to 1 if cgbn.h exists])
- NVCCFLAGS="-I$with_cgbn_include -lgmp $NVCCFLAGS"
+ NVCCFLAGS="-I$with_cgbn_include $GMPLIB $NVCCFLAGS"
want_cgbn="yes"
])
])

[SethTro]

I made a fairly large change so that the GPU kernel is called in much smaller batches. This helps with system responsiveness also gives a progress status.

[chris2be8]

After further testing on the system I was trying it on I'm fairly sure it won't work because the GPU is sm_30 (Fermi architecture) and CGBN needs sm_35 (Kepler) or higher. I had to look in the Makefiles shipped with CGBN to find out what sm_.. corresponds to what marketing name Nvidia use.

I have a newer GPU which is sm_52. But that system has an old version of the CUDA Toolkit and runtime (7.5/7.5). And that doesn't support CGBN. So I'll have to upgrade CUDA on that system. Which will take a while.

I've just used git clone 'https://github.com/sethtroisi/gmp-ecm.git' to grab a clean copy of the latest code instead of downloading gmp-ecm-gpu_integration.zip from the web site. I assume that's the best way to get it.

PS. git clone 'https://github.com/sethtroisi/gmp-ecm.git' isn't right, it doesn't include cgbn_stage1.* So what should I use?

[EdH]

Originally Posted by chris2be8

After further testing on the system I was trying it on I'm fairly sure it won't work because the GPU is sm_30 (Fermi architecture) and CGBN needs sm_35 (Kepler) or higher. I had to look in the Makefiles shipped with CGBN to find out what sm_.. corresponds to what marketing name Nvidia use.
. ..

Thank you! I've been watching this thread, wondering if I should try it with my sm_30 card, that won't build GMP-ECM or Msieve with CUDA 10.x, which is supposed to support that architecture. CUDA 11.x does not support it.

[henryzz]

There still seems to be cudaDeviceSynchronize issues although I get it at a very different line number.

echo "(2^499-1)/20959" | ./ecm -gpu -cgbn -gpucurves 3584 -sigma 3:1000 20000
GMP-ECM 7.0.5-dev [configured with GMP 6.2.99, --enable-asm-redc, --enable-gpu, --enable-assert, --enable-openmp] [ECM]
Input number is (2^499-1)/20959 (146 digits)
Using B1=20000, B2=3804582, sigma=3:1000-3:4583 (3584 curves)
CUDA error occurred: unknown error
While running cudaDeviceSynchronize()   (file cgbn_stage1.cu, line 733)

I have managed to get the non-cgbn code working under WSL2 although this required updating windows to 21H2 to enable gpu support which is only available as a preview so far.

[chris2be8]

Originally Posted by EdH

Thank you! I've been watching this thread, wondering if I should try it with my sm_30 card, that won't build GMP-ECM or Msieve with CUDA 10.x, which is supposed to support that architecture. CUDA 11.x does not support it.

My system with that card has CUDA 9.1 on it. So if you can "upgrade" to that level ecm and msieve should build.

[Plutie]

Originally Posted by chris2be8

After further testing on the system I was trying it on I'm fairly sure it won't work because the GPU is sm_30 (Fermi architecture) and CGBN needs sm_35 (Kepler) or higher. I had to look in the Makefiles shipped with CGBN to find out what sm_.. corresponds to what marketing name Nvidia use.

I have a newer GPU which is sm_52. But that system has an old version of the CUDA Toolkit and runtime (7.5/7.5). And that doesn't support CGBN. So I'll have to upgrade CUDA on that system. Which will take a while.

I've just used git clone 'https://github.com/sethtroisi/gmp-ecm.git' to grab a clean copy of the latest code instead of downloading gmp-ecm-gpu_integration.zip from the web site. I assume that's the best way to get it.

PS. git clone 'https://github.com/sethtroisi/gmp-ecm.git' isn't right, it doesn't include cgbn_stage1.* So what should I use?

The correct command to download the CGBN branch should be git clone https://github.com/sethtroisi/gmp-ecm/ -b gpu_integration folder_name.


PS: Also receiving the same CUDA error as chris2be8,

$ echo "2^997-1" | ./ecm -gpu -cgbn -sigma 3:1000 20000 0
GMP-ECM 7.0.5-dev [configured with GMP 6.2.1, --enable-asm-redc, --enable-gpu, --enable-assert] [ECM]
Input number is 2^997-1 (301 digits)
Using B1=20000, B2=0, sigma=3:1000-3:1639 (640 curves)
CUDA error occurred: cannot set while device is active in this process
While running cudaSetDeviceFlags (cudaDeviceScheduleBlockingSync) (file cgbn_stage1.cu, line 600)

[SethTro]

WraithX proposed a patch which will hopefully resolve the `CUDA error occurred: cannot set while device is active in this process` error. Hopefully it will be checked in later today / tonight.

[henryzz]

With Wraith's patch(with the header fixed) I now get:

echo "(2^499-1)/20959" | ./ecm -gpu -cgbn -gpucurves 3584 -sigma 3:1000 2000
GMP-ECM 7.0.5-dev [configured with GMP 6.2.99, --enable-asm-redc, --enable-gpu, --enable-assert, --enable-openmp] [ECM]
Input number is (2^499-1)/20959 (146 digits)
Using B1=2000, B2=147396, sigma=3:1000-3:4583 (3584 curves)
Computing 3584 Step 1 took 1863ms of CPU time / 2291ms of GPU time
Computing 3584 Step 2 on CPU took 12404ms

echo "(2^499-1)/20959" | ./ecm -gpu -gpucurves 3584 -sigma 3:1000 2000
GMP-ECM 7.0.5-dev [configured with GMP 6.2.99, --enable-asm-redc, --enable-gpu, --enable-assert, --enable-openmp] [ECM]
Input number is (2^499-1)/20959 (146 digits)
Using B1=2000, B2=147396, sigma=3:1000-3:4583 (3584 curves)
GPU: Block: 32x32x1 Grid: 112x1x1 (3584 parallel curves)
Computing 3584 Step 1 took 3668ms of CPU time / 6199ms of GPU time
Computing 3584 Step 2 on CPU took 12445ms

Although with B1=20000 I still get:

echo "(2^499-1)/20959" | ./ecm -gpu -cgbn -gpucurves 3584 -sigma 3:1000 20000
GMP-ECM 7.0.5-dev [configured with GMP 6.2.99, --enable-asm-redc, --enable-gpu, --enable-assert, --enable-openmp] [ECM]
Input number is (2^499-1)/20959 (146 digits)
Using B1=20000, B2=3804582, sigma=3:1000-3:4583 (3584 curves)
CUDA error (702) occurred: the launch timed out and was terminated
While running cudaDeviceSynchronize()   (file cgbn_stage1.cu, line 731)