mersenneforum.org  

Go Back   mersenneforum.org > Factoring Projects > Msieve

Reply
 
Thread Tools
Old 2021-01-28, 02:34   #12
Happy5214
 
Happy5214's Avatar
 
"Alexander"
Nov 2008
The Alamo City

2×3×5×23 Posts
Default

Quote:
Originally Posted by Brian Gladman View Post
Anyone who is having problems in building Windows x64 CUDA versions of MSIEVE or GMP-ECM using Visual Studio 2019 is welcome to report their issues here and I will do what I can to help.
FWIW I think xilman's post is a more appropriate quote than mine, since I was describing a Kubuntu (Linux) computer. Do you have any advice for us Linux users?
Happy5214 is offline   Reply With Quote
Old 2021-01-28, 06:31   #13
Gimarel
 
Apr 2010

2·83 Posts
Default

Quote:
Originally Posted by Happy5214 View Post
FWIW I think xilman's post is a more appropriate quote than mine, since I was describing a Kubuntu (Linux) computer. Do you have any advice for us Linux users?
I'm using debian and a GTX 2060 Super.
First you have to get rid of the cub that comes with msieve and use the cub that comes with cuda. On debian I did these commands in the msieve directory:
Code:
cd cub
rm -rf cub
ln -s /usr/include/cub
Then you need the following patch:
Code:
Index: Makefile
===================================================================
--- Makefile    (Revision 1037)
+++ Makefile    (Arbeitskopie)
@@ -34,6 +34,9 @@
          -DMSIEVE_SVN_VERSION="\"$(SVN_VERSION)\"" \
         -I. -Iaprcl -Iinclude -Ignfs -Ignfs/poly -Ignfs/poly/stage1
 
+CUDA = 1
+NO_ZLIB = 1
+
 # tweak the compile flags
 
 ifeq ($(ECM),1)
@@ -197,10 +200,7 @@
 #---------------------------------- GPU file lists -------------------------
 
 GPU_OBJS += \
-    stage1_core_sm20.ptx \
-    stage1_core_sm30.ptx \
-    stage1_core_sm35.ptx \
-    stage1_core_sm50.ptx \
+    stage1_core_sm75.ptx \
     cub/built
 
 #---------------------------------- NFS file lists -------------------------
@@ -346,5 +346,8 @@
 stage1_core_sm50.ptx: $(NFS_GPU_HDR)
     $(NVCC) -arch sm_50 -ptx -o $@ $<
 
+stage1_core_sm75.ptx: $(NFS_GPU_HDR)
+    $(NVCC) -arch sm_75 -ptx -o $@ $<
+
 cub/built:
-    cd cub && make WIN=$(WIN) WIN64=$(WIN64) sm=200,300,350,520 && cd ..
+    cd cub && make WIN=$(WIN) WIN64=$(WIN64) sm=750 && cd ..
Index: cub/Makefile
===================================================================
--- cub/Makefile    (Revision 1037)
+++ cub/Makefile    (Arbeitskopie)
@@ -16,7 +16,7 @@
     NVCC = "$(shell which nvcc)"
     CUDA_ROOT = $(shell dirname $(NVCC))/../
     EXT = so
-    NVCCFLAGS += -Xptxas -v -Xcudafe -\# -shared -Xptxas -abi=no \
+    NVCCFLAGS += -Xptxas -v -Xcudafe -\# -shared \
             -Xcompiler -fPIC -Xcompiler -fvisibility=hidden
 endif
 
@@ -27,6 +27,10 @@
     SM_ARCH = 200
 endif
 
+ifeq (750, $(findstring 750, $(SM_ARCH)))
+    SM_TARGETS     += -gencode=arch=compute_75,code=\"sm_75,compute_75\" 
+    SM_DEF         += -DSM750
+endif
 ifeq (520, $(findstring 520, $(SM_ARCH)))
     SM_TARGETS     += -gencode=arch=compute_52,code=\"sm_52,compute_52\" 
     SM_DEF         += -DSM520
Index: gnfs/poly/stage1/stage1_sieve_gpu.c
===================================================================
--- gnfs/poly/stage1/stage1_sieve_gpu.c    (Revision 1037)
+++ gnfs/poly/stage1/stage1_sieve_gpu.c    (Arbeitskopie)
@@ -1113,7 +1113,7 @@
             CUDA_TRY(cuModuleLoad(&t->gpu_module, "stage1_core_sm35.ptx"))
     }
     else if (d->gpu_info->compute_version_major >= 5) {
-        CUDA_TRY(cuModuleLoad(&t->gpu_module, "stage1_core_sm50.ptx"))
+        CUDA_TRY(cuModuleLoad(&t->gpu_module, "stage1_core_sm75.ptx"))
     }
     else 
     {
Note that the msieve built this way only works with CUDA arch 7.5.
Gimarel is offline   Reply With Quote
Old 2021-01-28, 06:47   #14
Gimarel
 
Apr 2010

2×83 Posts
Default

For ecm to compile with current CUDA I had to apply this patch:
Code:
Index: cudakernel_default.cu
===================================================================
--- cudakernel_default.cu    (Revision 3092)
+++ cudakernel_default.cu    (Arbeitskopie)
@@ -7,7 +7,7 @@
   carry_t cytemp;
   unsigned int thm1;
 
-  while(__any(cy[threadIdx.x])!=0)
+  while(__any_sync(__activemask(), cy[threadIdx.x])!=0)
   {
     thm1 = (threadIdx.x - 1) % ECM_GPU_NB_DIGITS;
     cytemp = cy[thm1];
I don't know for sure if this is correct, but it works.
Gimarel is offline   Reply With Quote
Old 2021-01-28, 09:28   #15
Happy5214
 
Happy5214's Avatar
 
"Alexander"
Nov 2008
The Alamo City

12628 Posts
Default

Thank you, thank you, thank you! They both build now. yafu doesn't build, though. I'll report that to Ben.
Happy5214 is offline   Reply With Quote
Old 2021-01-28, 10:06   #16
firejuggler
 
firejuggler's Avatar
 
Apr 2010
Over the rainbow

23·52·13 Posts
Default

I do have the PTX but I still have trouble with the sort_engine...


edit : adding second screenshot
Attached Thumbnails
Click image for larger version

Name:	Screenshot_7.jpg
Views:	78
Size:	42.6 KB
ID:	24240   Click image for larger version

Name:	Screenshot_8.jpg
Views:	84
Size:	167.2 KB
ID:	24241  
Attached Files
File Type: txt report.txt (27.9 KB, 72 views)

Last fiddled with by firejuggler on 2021-01-28 at 10:24 Reason: adding another screenshot and a report
firejuggler is offline   Reply With Quote
Old 2021-01-28, 10:25   #17
Gimarel
 
Apr 2010

2×83 Posts
Default

I think that the 1660 Ti needs a different shader model. But I don't know which.
Gimarel is offline   Reply With Quote
Old 2021-01-28, 10:37   #18
firejuggler
 
firejuggler's Avatar
 
Apr 2010
Over the rainbow

50508 Posts
Default

a quick google search tell me it is shader model 6.5... so I should modify the sm_arch to 65?
firejuggler is offline   Reply With Quote
Old 2021-01-28, 15:38   #19
Gimarel
 
Apr 2010

2468 Posts
Default

Quote:
Originally Posted by firejuggler View Post
a quick google search tell me it is shader model 6.5... so I should modify the sm_arch to 65?
The CUDA Samples contain a deviceQuery progam. The compile contains these options for CUDA 11.1:
Code:
-gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_75,code=sm_75 -gencode arch=compute_80,code=sm_80 -gencode arch=compute_86,code=sm_86 -gencode arch=compute_86,code=compute_86
Gimarel is offline   Reply With Quote
Old 2021-06-08, 16:28   #20
bsquared
 
bsquared's Avatar
 
"Ben"
Feb 2007

22×3×293 Posts
Default

Quote:
Originally Posted by Brian Gladman View Post
Anyone who is having problems in building Windows x64 CUDA versions of MSIEVE or GMP-ECM using Visual Studio 2019 is welcome to report their issues here and I will do what I can to help.
After adjusting the project files to point to my nvidia cuda toolkit location, I just built msieve with compute_75. I have an up-to-date sort_engine.dll and stage1_core_sm75.ptx file. But when trying to run msieve-gpu.exe it gives me this error:

Code:
Msieve v. 1.54 (SVN 998)
Tue Jun  8 11:24:02 2021
random seeds: 2f8aea7c 74608973
factoring 138924029959401366454963864059579437250850355925904953363654825080008713183159095653855715163496880698665441863162263 (117 digits)
searching for 15-digit factors
commencing number field sieve (117-digit input)
commencing number field sieve polynomial selection
polynomial degree: 5
max stage 1 norm: 2.24e+18
max stage 2 norm: 7.93e+14
min E-value: 4.22e-10
poly select deadline: 6326
time limit set to 1.76 CPU-hours
expecting poly E from 5.27e-10 to > 6.06e-10
searching leading coefficients from 1 to 213804
using GPU 0 (TITAN RTX)
selected card has CUDA arch 7.5
deadline: 6326 CPU-seconds per coefficient
error (line 1116): CUDA_ERROR_FILE_NOT_FOUND
I'm not sure what file it's not finding or where it's not finding it.
bsquared is offline   Reply With Quote
Old 2021-06-08, 16:41   #21
Gimarel
 
Apr 2010

2468 Posts
Default

My SVN revision is a bit newer, but the solution is the same:
Code:
Index: gnfs/poly/stage1/stage1_sieve_gpu.c
===================================================================
--- gnfs/poly/stage1/stage1_sieve_gpu.c    (Revision 1030)
+++ gnfs/poly/stage1/stage1_sieve_gpu.c    (Arbeitskopie)
@@ -1113,7 +1113,7 @@
             CUDA_TRY(cuModuleLoad(&t->gpu_module, "stage1_core_sm35.ptx"))
     }
     else if (d->gpu_info->compute_version_major >= 5) {
-        CUDA_TRY(cuModuleLoad(&t->gpu_module, "stage1_core_sm50.ptx"))
+        CUDA_TRY(cuModuleLoad(&t->gpu_module, "stage1_core_sm75.ptx"))
     }
     else 
     {
Gimarel is offline   Reply With Quote
Old 2021-06-08, 16:59   #22
bsquared
 
bsquared's Avatar
 
"Ben"
Feb 2007

DBC16 Posts
Default

Working now, thanks!
bsquared is offline   Reply With Quote
Reply

Thread Tools


Similar Threads
Thread Thread Starter Forum Replies Last Post
CUDA 5.5 ET_ GPU Computing 2 2013-06-13 15:50
AVX CPU LL vs CUDA LL nucleon GPU Computing 11 2012-01-04 17:52
Best CUDA GPU for the $$ Christenson GPU Computing 24 2011-05-01 00:06
CUDA P-1? nucleon GPU Computing 2 2010-11-17 17:52
CUDA? Xentar Conjectures 'R Us 6 2010-03-31 07:43

All times are UTC. The time now is 07:43.


Sat Jul 24 07:43:42 UTC 2021 up 1 day, 2:12, 1 user, load averages: 2.10, 1.59, 1.40

Powered by vBulletin® Version 3.8.11
Copyright ©2000 - 2021, Jelsoft Enterprises Ltd.

This forum has received and complied with 0 (zero) government requests for information.

Permission is granted to copy, distribute and/or modify this document under the terms of the GNU Free Documentation License, Version 1.2 or any later version published by the Free Software Foundation.
A copy of the license is included in the FAQ.