CUDA 9 and sm_20
 

The latest CUDA compiler no longer supports the cc2.x architecture. This has caused a number of problems and has now hit msieve's Makefile. :sad:

Anyone want to decide what to do and update msieve appropriately?

I'd do it but there are doubtless folks here more skilled than I am. For a start, I have no capability to build under Windoze.

FWIW, GMP-ECM has hit the same issue. I'm responsible for the CUDA build there but we explicitly don't support native Windoze builds and so the problem can be safely ignored for the time being.

Added in edit: I really don't want to do it. An attempt to kludge my way through led me in to a twisty little maze of nvcc options all different.

It gets worse:
[code]
pcl@anubis ~/nums/msieve-code/trunk $ ./msieve -i W852.n -g 0 -t 6 -np
integrator failed nan inf
[/code]
Oh well, lets see whether CADO-NFS still works.

Nope it didn't. :sad:

I can see the New Year's resolution is to get the old tool kits working again. Ho hum.

This message occured with "recent" CADO-NFS 2.3.0 with CUDA. For reference, since this was the only exact search result for that exact message I got on Google, I got it working by removing every temporary or intermediate file in the executable's directory. Before that, it crashed with the exact same message.

[QUOTE=kruoli;57010]... CADO-NFS 2.3.0 ...[/QUOTE]

Sleep more, think more... :davieddy:
I have no idea why I was thinking about CADO, this is of course totally wrong. Sorry for the misinformation. I was using msieve, SVN 998, and had the problem.

The last time I had to install msieve on a system with a CUDA capable GPU I had to update the Makefiles as follows:
[code]
chris@sirius:~/factordb$ diff -u /home/chris/msieve.1030/trunk/Makefile /home/chris/msieve.1030.cuda/trunk/Makefile
--- /home/chris/msieve.1030/trunk/Makefile 2019-12-07 19:15:56.039498635 +0000
+++ /home/chris/msieve.1030.cuda/trunk/Makefile 2019-12-07 20:01:18.819251935 +0000
@@ -196,8 +196,8 @@

#---------------------------------- GPU file lists -------------------------

+# stage1_core_sm20.ptx
GPU_OBJS += \
- stage1_core_sm20.ptx \
stage1_core_sm30.ptx \
stage1_core_sm35.ptx \
stage1_core_sm50.ptx \
@@ -334,8 +334,8 @@

# GPU build rules

-stage1_core_sm20.ptx: $(NFS_GPU_HDR)
- $(NVCC) -arch sm_20 -ptx -o $@ $<
+#stage1_core_sm20.ptx: $(NFS_GPU_HDR)
+# $(NVCC) -arch sm_20 -ptx -o $@ $<

stage1_core_sm30.ptx: $(NFS_GPU_HDR)
$(NVCC) -arch sm_30 -ptx -o $@ $<
@@ -347,4 +347,4 @@
$(NVCC) -arch sm_50 -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=300,350,520 && cd ..
[/code]

[code]
chris@sirius:~/factordb$ diff -u /home/chris/msieve.1030/trunk/cub/Makefile /home/chris/msieve.1030.cuda/trunk/cub/Makefile
--- /home/chris/msieve.1030/trunk/cub/Makefile 2019-12-07 19:16:01.763534223 +0000
+++ /home/chris/msieve.1030.cuda/trunk/cub/Makefile 2019-12-09 21:40:42.328453882 +0000
@@ -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

@@ -43,14 +43,6 @@
SM_TARGETS += -gencode=arch=compute_30,code=\"sm_30,compute_30\"
SM_DEF += -DSM300
endif
-ifeq (210, $(findstring 210, $(SM_ARCH)))
- SM_TARGETS += -gencode=arch=compute_20,code=\"sm_21,compute_20\"
- SM_DEF += -DSM210
-endif
-ifeq (200, $(findstring 200, $(SM_ARCH)))
- SM_TARGETS += -gencode=arch=compute_20,code=\"sm_20,compute_20\"
- SM_DEF += -DSM200
-endif

rwildcard=$(foreach d,$(wildcard $1*),$(call rwildcard,$d/,$2) $(filter $(subst *,%,$2),$d))
[/code]

After that it worked for me.

Chris

[QUOTE=xilman;475849]The latest CUDA compiler no longer supports the cc2.x architecture. This has caused a number of problems and has now hit msieve's Makefile. :sad:

Anyone want to decide what to do and update msieve appropriately?

I'd do it but there are doubtless folks here more skilled than I am. For a start, I have no capability to build under Windoze.

FWIW, GMP-ECM has hit the same issue. I'm responsible for the CUDA build there but we explicitly don't support native Windoze builds and so the problem can be safely ignored for the time being.

Added in edit: I really don't want to do it. An attempt to kludge my way through led me in to a twisty little maze of nvcc options all different.[/QUOTE]

GMP-ECM explicitly does support native Windows builds and has done so for many years (I maintain this part of the build which uses Microsoft Visual Studio). I also fully support the CUDA build of MSIEVE on Windows.

[QUOTE=Brian Gladman;570162]GMP-ECM explicitly does support native Windows builds and has done so for many years (I maintain this part of the build which uses Microsoft Visual Studio). I also fully support the CUDA build of MSIEVE on Windows.[/QUOTE]Excellent! :bow:

Do you have fully static builds available? As I can't get anything to work natively it might be worth trying to run a Windows executable under WINE.

[QUOTE=xilman;570164]Excellent! :bow:

Do you have fully static builds available? As I can't get anything to work natively it might be worth trying to run a Windows executable under WINE.[/QUOTE]

HI Paul,

The builds in the official repositories for GMP-ECM and MSIEVE are both static. If you don't know the repository locations I can let you have them (I am not sure about posting their locations here).

I would love a ptx compatible with a compute ability 7.5....

[QUOTE=firejuggler;570180]I would love a ptx compatible with a compute ability 7.5....[/QUOTE]

Seconded. Attempting to build msieve and GMP-ECM on my brand-new Kubuntu Focus (RTX 2060) was very frustrating, and I still don't think I did it right.

[QUOTE=Happy5214;570198]Seconded. Attempting to build msieve and GMP-ECM on my brand-new Kubuntu Focus (RTX 2060) was very frustrating, and I still don't think I did it right.[/QUOTE]
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.

[QUOTE=Brian Gladman;570224]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.[/QUOTE]
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?

[QUOTE=Happy5214;570302]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?[/QUOTE]
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[/CODE]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
{
[/CODE]Note that the msieve built this way only works with CUDA arch 7.5.

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];
[/CODE]I don't know for sure if this is correct, but it works.

Thank you, thank you, thank you! They both build now. yafu doesn't build, though. I'll report that to Ben.

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


edit : adding second screenshot

I think that the 1660 Ti needs a different shader model. But I don't know which.

a quick google search tell me it is shader model 6.5... so I should modify the sm_arch to 65?

[QUOTE=firejuggler;570323]a quick google search tell me it is shader model 6.5... so I should modify the sm_arch to 65?[/QUOTE]
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[/CODE]

[QUOTE=Brian Gladman;570224]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.[/QUOTE]

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[/CODE]

I'm not sure what file it's not finding or where it's not finding it.

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
{
[/CODE]

Working now, thanks!

[QUOTE=bsquared;580364]Working now, thanks![/QUOTE]Does that mean you have a working GPU version (Windows) you can share? The latest [url=https://download.mersenne.ca/msieve]version[/url] I have is from 2016...

[QUOTE=James Heinrich;580367]Does that mean you have a working GPU version (Windows) you can share? The latest [url=https://download.mersenne.ca/msieve]version[/url] I have is from 2016...[/QUOTE]

No, doesn't look like it. It builds ok and now it doesn't complain about missing files. But it seems to get stuck somewhere when I run it... uses 0% gpu and hangs.