search

cbuchner1 / CudaMiner

a CUDA accelerated litecoin mining application based on pooler's CPU miner
Other
687 stars 304 forks source link

ptxas error while compile BUT fix included! #163

Open LouisCyfer opened 6 years ago

LouisCyfer commented 6 years ago

Hiyas!

using Ubuntu 16.04, cuda 9.1, nVidia GTX 750ti

I'm getting this errors while compiling:

ptxas error   : Invalid value 'no' for option -abi.
ptxas fatal   : Ptx assembly aborted due to errors

I'm getting curious (for cuda 9.1 there is no "-abi" option listed in the official documentation)

Nevertheless, I did change all compute_10 to compute_50 also exchanged ALL sm_* with sm_50 in makefile.am

BOOM! it compiled!

then when I run ./cudaminer --no-autotune --benchmark it thows this errors:

[2017-12-31 04:39:36] NVML GPU monitoring is not available.
[2017-12-31 04:39:36] 1 miner threads started, using 'scrypt' algorithm.
[2017-12-31 04:39:39] GPU #0: GeForce GTX 750 Ti with compute capability 5.0
[2017-12-31 04:39:39] GPU #0: interactive: 1, tex-cache: 0 , single-alloc: 0
[2017-12-31 04:39:39] GPU #0: 32 hashes / 4.0 MB per warp.
[2017-12-31 04:39:39] GPU #0: using launch configuration T120x2
[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(sha256_h, host_sha256_h, sizeof(host_sha256_h), 0, cudaMemcpyHostToDevice)' (sha256.cu line 415)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(sha256_k, host_sha256_k, sizeof(host_sha256_k), 0, cudaMemcpyHostToDevice)' (sha256.cu line 416)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(keypad, host_keypad, sizeof(host_keypad), 0, cudaMemcpyHostToDevice)' (sha256.cu line 417)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(innerpad, host_innerpad, sizeof(host_innerpad), 0, cudaMemcpyHostToDevice)' (sha256.cu line 418)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(outerpad, host_outerpad, sizeof(host_outerpad), 0, cudaMemcpyHostToDevice)' (sha256.cu line 419)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(finalblk, host_finalblk, sizeof(host_finalblk), 0, cudaMemcpyHostToDevice)' (sha256.cu line 420)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(pdata, host_pdata, 20*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)' (sha256.cu line 423)

[2017-12-31 04:39:39] GPU #0: cudaError 13 (invalid device symbol) calling 'cudaMemcpyToSymbol(midstate, host_midstate, 8*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)' (sha256.cu line 424)

note: ./cudaminer -H 0 --no-autotune --benchmark gives me

[2017-12-31 04:49:34] NVML GPU monitoring is not available.
[2017-12-31 04:49:34] 1 miner threads started, using 'scrypt' algorithm.
[2017-12-31 04:49:34] Binding thread 0 to cpu 0
[2017-12-31 04:49:37] GPU #0: GeForce GTX 750 Ti with compute capability 5.0
[2017-12-31 04:49:37] GPU #0: interactive: 1, tex-cache: 0 , single-alloc: 0
[2017-12-31 04:49:37] GPU #0: 32 hashes / 4.0 MB per warp.
[2017-12-31 04:49:37] GPU #0: using launch configuration T120x2
[2017-12-31 04:49:38] GPU #0: GeForce GTX 750 Ti, 396.57 khash/s
[2017-12-31 04:49:38] Total: 396.57 khash/s

is this appropiate?

roman3017 commented 6 years ago

@LouisCyfer: I have tried your changes but getting error: ptxas fatal : redefinition of argument 'gpu-name'

Can you please share your Makefile.am?

LouisCyfer commented 6 years ago

I figured exchanging all compute_* and sm_* to sm_50 did the trick better with my GPU (GTX 750ti) and cuda version (9.1)

here's the Makefile.am:


if WANT_JANSSON
JANSSON_INCLUDES= -I$(top_srcdir)/compat/jansson
else
JANSSON_INCLUDES=
endif

EXTRA_DIST      = autogen.sh README.txt LICENSE.txt \
              cudaminer.sln cudaminer.vcxproj cudaminer.vcxproj.filters \
              compat/gettimeofday.c compat/getopt/getopt_long.c cpuminer-config.h.in

SUBDIRS     = compat

bin_PROGRAMS    = cudaminer

cudaminer_SOURCES   = elist.h miner.h compat.h \
              compat/inttypes.h compat/stdbool.h compat/unistd.h \
              compat/sys/time.h compat/getopt/getopt.h \
              cpu-miner.c util.c \
              wrapnvml.cu wrapnvml.h \
              sha2.c \
              scrypt.cpp \
              maxcoin.cpp \
              blakecoin.cpp \
              sha3.cpp sha3.h \
              scrypt-jane.cpp scrypt-jane.h \
              salsa_kernel.cu salsa_kernel.h \
              sha256.cu sha256.h \
              keccak.cu keccak.h \
              blake.cu blake.h blake.c sph_blake.h sph_types.h \
              fermi_kernel.cu fermi_kernel.h \
              kepler_kernel.cu kepler_kernel.h \
              test_kernel.cu test_kernel.h \
              nv_kernel.cu nv_kernel.h \
              nv_kernel2.cu nv_kernel2.h \
              titan_kernel.cu titan_kernel.h

cudaminer_LDFLAGS   = $(PTHREAD_FLAGS) @CUDA_LDFLAGS@
cudaminer_LDADD     = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @WS2_LIBS@ @CUDA_LIBS@ @OPENMP_CFLAGS@ @LIBS@ -ldl
cudaminer_CPPFLAGS  = -msse2 @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME

.cu.o:
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=64 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<

sha256.o: sha256.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=64 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<

keccak.o: keccak.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=64 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<

blake.o: blake.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=64 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<

# NOTE: now compiling for compute 1.0 again, as it's using less power and runs way faster on Linux
fermi_kernel.o: fermi_kernel.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=64 $(JANSSON_INCLUDES) -o $@ -c $<

kepler_kernel.o: kepler_kernel.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=32 $(JANSSON_INCLUDES) -o $@ -c $<

titan_kernel.o: titan_kernel.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=32 $(JANSSON_INCLUDES) -o $@ -c $<

# NOTE: now compiling for compute 1.0 again, as it's using less power and runs way faster on Linux
test_kernel.o: test_kernel.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=32 $(JANSSON_INCLUDES) -o $@ -c $<

nv_kernel.o: nv_kernel.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=63 $(JANSSON_INCLUDES) -o $@ -c $<

nv_kernel2.o: nv_kernel2.cu
    $(NVCC) @CFLAGS@ -Xptxas "-abi=no -v" -arch=sm_50 --maxrregcount=80 $(JANSSON_INCLUDES) -o $@ -c $<
roman3017 commented 6 years ago

@LouisCyfer: Thank you very much for sharing. With your Makefile.am I am getting the original issue:

ptxas error   : Invalid value 'no' for option -abi.
ptxas fatal   : Ptx assembly aborted due to errors

I also have cuda version 9.1. My card is GeForce GT 750M.

I have got my card working with changes here: https://github.com/roman3017/CudaMiner/commit/4f4a100f5faf3d818221fcfde6695beced3667cf

I wonder, would removing all:

"-abi=no -v" -arch=xxx

work for you? It somehow works for me.

LouisCyfer commented 6 years ago

if you remove these two params cuda might not know what exact routine to support, have a look on the -arch parameter description:

--gpu-architecture <arch>                  (-arch)                           
        Specify the name of the class of NVIDIA 'virtual' GPU architecture for which
        the CUDA input files must be compiled.
        With the exception as described for the shorthand below, the architecture
        specified with this option must be a 'virtual' architecture (such as compute_20).
        Normally, this option alone does not trigger assembly of the generated PTX
        for a 'real' architecture (that is the role of nvcc option '--gpu-code',
        see below); rather, its purpose is to control preprocessing and compilation
        of the input to PTX.
        For convenience, in case of simple nvcc compilations, the following shorthand
        is supported.  If no value for option '--gpu-code' is specified, then the
        value of this option defaults to the value of '--gpu-architecture'.  In this
        situation, as only exception to the description above, the value specified
        for '--gpu-architecture' may be a 'real' architecture (such as a sm_20),
        in which case nvcc uses the specified 'real' architecture and its closest
        'virtual' architecture as effective architecture values.  For example, 'nvcc
        --gpu-architecture=sm_20' is equivalent to 'nvcc --gpu-architecture=compute_20
        --gpu-code=sm_20,compute_20'.
        Allowed values for this option:  'compute_20','compute_30','compute_32',
        'compute_35','compute_37','compute_50','compute_52','compute_53','sm_20',
        'sm_21','sm_30','sm_32','sm_35','sm_37','sm_50','sm_52','sm_53'.

note:

nvcc --gpu-architecture=sm_50 or -arch=sm_50
is equivalent to
'nvcc --gpu-architecture=compute_50 --gpu-code=sm_50,compute_50

You might wanna try replacing all compute_* with sm_50 and then just run this command ./autogen.sh && ./configure.sh && make

Removing these parameters might just not give the miner engine less specific workspace or rather said, the taxi driver doesn't know about the fuel or electricity in a car before driving, means cuda might just not know what to do with it if compiled, and you might have issues with it if the miner tries to access the libs.

I might forgot to mention that I'm running 64bit but it should work on 32bit too.

roman3017 commented 6 years ago

Unfortunately when I use -arch=sm_50 I am still getting the original issue:

. autogen.sh
. configure.sh
make
...
ptxas error   : Invalid value 'no' for option -abi.
ptxas fatal   : Ptx assembly aborted due to errors
...

You do not need to remove at least this part?

"-abi=no -v"

I do not have build.sh but hope it is similar to the three commands above.

LouisCyfer commented 6 years ago

I edited my previous post because Idid mistake it with another miner, the command line is like you commented.

well you might also try a fresh git clone which will grab the most recent/latest release (is how I made it), no zip download.

probably there is some other option, did you also tried 30 or 32 instead of 50?

roman3017 commented 6 years ago

Yes, I have tried. It seems I cannot use -arch option:

$ /usr/local/cuda/bin/nvcc -O3 -Xptxas -arch=sm_30 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o wrapnvml.o -c wrapnvml.cu
ptxas fatal   : redefinition of argument 'gpu-name'
$ /usr/local/cuda/bin/nvcc -O3 -Xptxas -gpu-architecture=sm_30 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o wrapnvml.o -c wrapnvml.cu
ptxas fatal   : Unknown option 'gpu-architecture'
$ /usr/local/cuda/bin/nvcc -O3 -Xptxas --gpu-architecture=sm_30 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o wrapnvml.o -c wrapnvml.cu
ptxas fatal   : Unknown option '-gpu-architecture'

Anyway, I have got it working by removing -arch and -abi options.

LouisCyfer commented 6 years ago

-gpu-architecture should be --gpu-architecture with two minus characters OR -arch with one minus character

roman3017 commented 6 years ago

Ok, finally got it working with -arch option:

-Xptxas "-v" -arch=sm_30

@LouisCyfer: Thank you very much for your help, I really apreciate it.

LouisCyfer commented 6 years ago

YAY! very welcome!

joshuamsalazar commented 3 years ago

I looked almost everywhere, and in the CUDA documentation there is no detail abut the -abi flag. I am compiling a CUDA app in Linux. I installed all the required deppendencies, toolkits and libraries required. However, I cannot get rid of that error.

I followed @roman3017 suggestion of reomving "-abi=no -v" and just setting "-v" into the instructions:

 nvcc -g -02 -Xptas "-abi=no -v" -arch=sm_35

would do the trick. However I get this big set of errors:

make  all-recursive
make[1]: Entering directory '/data/Install/UNIX/CudaMiner-master'
Making all in compat
make[2]: Entering directory '/data/Install/UNIX/CudaMiner-master/compat'
Making all in jansson
make[3]: Entering directory '/data/Install/UNIX/CudaMiner-master/compat/jansson'
make[3]: Nothing to be done for 'all'.
make[3]: Leaving directory '/data/Install/UNIX/CudaMiner-master/compat/jansson'
make[3]: Entering directory '/data/Install/UNIX/CudaMiner-master/compat'
make[3]: Nothing to be done for 'all-am'.
make[3]: Leaving directory '/data/Install/UNIX/CudaMiner-master/compat'
make[2]: Leaving directory '/data/Install/UNIX/CudaMiner-master/compat'
make[2]: Entering directory '/data/Install/UNIX/CudaMiner-master'
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o wrapnvml.o -c wrapnvml.cu
ptxas info    : 0 bytes gmem
gcc -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-sha2.o -MD -MP -MF .deps/cudaminer-sha2.Tpo -c -o cudaminer-sha2.o `test -f 'sha2.c' || echo './'`sha2.c
mv -f .deps/cudaminer-sha2.Tpo .deps/cudaminer-sha2.Po
g++ -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-scrypt.o -MD -MP -MF .deps/cudaminer-scrypt.Tpo -c -o cudaminer-scrypt.o `test -f 'scrypt.cpp' || echo './'`scrypt.cpp
scrypt.cpp:60:38: warning: dynamic exception specifications are deprecated in C++11 [-Wdeprecated]
   60 |     void * operator new(size_t size) throw(std::bad_alloc) { void *p; if (posix_memalign(&p, 16, size) < 0) { static const std::bad_alloc nomem; throw nomem; } return (p); }
      |                                      ^~~~~
scrypt.cpp:62:40: warning: dynamic exception specifications are deprecated in C++11 [-Wdeprecated]
   62 |     void * operator new[](size_t size) throw(std::bad_alloc) { void *p; if (posix_memalign(&p, 16, size) < 0) { static const std::bad_alloc nomem; throw nomem; } return (p); }
      |                                        ^~~~~
mv -f .deps/cudaminer-scrypt.Tpo .deps/cudaminer-scrypt.Po
g++ -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-maxcoin.o -MD -MP -MF .deps/cudaminer-maxcoin.Tpo -c -o cudaminer-maxcoin.o `test -f 'maxcoin.cpp' || echo './'`maxcoin.cpp
mv -f .deps/cudaminer-maxcoin.Tpo .deps/cudaminer-maxcoin.Po
g++ -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-blakecoin.o -MD -MP -MF .deps/cudaminer-blakecoin.Tpo -c -o cudaminer-blakecoin.o `test -f 'blakecoin.cpp' || echo './'`blakecoin.cpp
mv -f .deps/cudaminer-blakecoin.Tpo .deps/cudaminer-blakecoin.Po
g++ -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-sha3.o -MD -MP -MF .deps/cudaminer-sha3.Tpo -c -o cudaminer-sha3.o `test -f 'sha3.cpp' || echo './'`sha3.cpp
mv -f .deps/cudaminer-sha3.Tpo .deps/cudaminer-sha3.Po
g++ -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-scrypt-jane.o -MD -MP -MF .deps/cudaminer-scrypt-jane.Tpo -c -o cudaminer-scrypt-jane.o `test -f 'scrypt-jane.cpp' || echo './'`scrypt-jane.cpp
mv -f .deps/cudaminer-scrypt-jane.Tpo .deps/cudaminer-scrypt-jane.Po
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o salsa_kernel.o -c salsa_kernel.cu
salsa_kernel.cu(718): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu(719): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu(720): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu(721): warning: conversion from a string literal to "char *" is deprecated

ptxas info    : 0 bytes gmem
salsa_kernel.cu(718): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu(719): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu(720): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu(721): warning: conversion from a string literal to "char *" is deprecated

salsa_kernel.cu: In function ‘void cuda_shutdown(int)’:
salsa_kernel.cu:225:16: warning: ‘cudaError_t cudaThreadExit()’ is deprecated [-Wdeprecated-declarations]
  225 |     cudaThreadExit();
      |                ^
/usr/include/cuda_runtime_api.h:941:46: note: declared here
  941 | extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadExit(void);
      |                                              ^~~~~~~~~~~~~~
salsa_kernel.cu:225:16: warning: ‘cudaError_t cudaThreadExit()’ is deprecated [-Wdeprecated-declarations]
  225 |     cudaThreadExit();
      |                ^
/usr/include/cuda_runtime_api.h:941:46: note: declared here
  941 | extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadExit(void);
      |                                              ^~~~~~~~~~~~~~
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o sha256.o -c sha256.cu
ptxas info    : 0 bytes gmem, 588 bytes cmem[3]
ptxas info    : Compiling entry function '_Z16cuda_post_sha256PjS_S_S_' for 'sm_80'
ptxas info    : Function properties for _Z16cuda_post_sha256PjS_S_S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 55 registers, 384 bytes cmem[0]
ptxas info    : Compiling entry function '_Z15cuda_pre_sha256PjS_S_j' for 'sm_80'
ptxas info    : Function properties for _Z15cuda_pre_sha256PjS_S_j
    192 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads
ptxas info    : Used 64 registers, 380 bytes cmem[0]
ptxas info    : Function properties for _Z21cuda_sha256_transformPjPKj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o keccak.o -c keccak.cu
ptxas info    : 0 bytes gmem, 576 bytes cmem[3]
ptxas info    : Compiling entry function '_Z11crypto_hashPmjPjb' for 'sm_80'
ptxas info    : Function properties for _Z11crypto_hashPmjPjb
    16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads
ptxas info    : Used 64 registers, 377 bytes cmem[0]
ptxas info    : Compiling entry function '_Z19cuda_post_keccak512PjS_j' for 'sm_80'
ptxas info    : Function properties for _Z19cuda_post_keccak512PjS_j
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 124 registers, 372 bytes cmem[0]
ptxas info    : Compiling entry function '_Z18cuda_pre_keccak512Pjj' for 'sm_80'
ptxas info    : Function properties for _Z18cuda_pre_keccak512Pjj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 168 registers, 364 bytes cmem[0]
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=64 --ptxas-options=-v -I./compat/jansson -o blake.o -c blake.cu
ptxas info    : 0 bytes gmem, 112 bytes cmem[3]
ptxas info    : Compiling entry function '_Z18cuda_blake256_hashPmjPjb' for 'sm_80'
ptxas info    : Function properties for _Z18cuda_blake256_hashPmjPjb
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 48 registers, 377 bytes cmem[0]
gcc -DHAVE_CONFIG_H -I.  -msse2  -fopenmp -pthread -fno-strict-aliasing -I./compat/jansson -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME   -g -O2 -MT cudaminer-blake.o -MD -MP -MF .deps/cudaminer-blake.Tpo -c -o cudaminer-blake.o `test -f 'blake.c' || echo './'`blake.c
mv -f .deps/cudaminer-blake.Tpo .deps/cudaminer-blake.Po
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=64 -I./compat/jansson -o fermi_kernel.o -c fermi_kernel.cu
fermi_kernel.cu(112): warning: variable "sleeptime" was declared but never referenced

ptxas info    : 0 bytes gmem, 32768 bytes cmem[3]
ptxas info    : Compiling entry function '_Z28fermi_scrypt_core_kernelB_LGILi1EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z28fermi_scrypt_core_kernelB_LGILi1EEvPjjj
    176 bytes stack frame, 428 bytes spill stores, 432 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0]
ptxas info    : Compiling entry function '_Z28fermi_scrypt_core_kernelB_LGILi0EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z28fermi_scrypt_core_kernelB_LGILi0EEvPjjj
    176 bytes stack frame, 452 bytes spill stores, 456 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0]
ptxas info    : Compiling entry function '_Z32fermi_scrypt_core_kernelB_LG_texILi1ELi2EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z32fermi_scrypt_core_kernelB_LG_texILi1ELi2EEvPjjj
    160 bytes stack frame, 384 bytes spill stores, 384 bytes spill loads
ptxas info    : Used 64 registers, 372 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z32fermi_scrypt_core_kernelB_LG_texILi0ELi2EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z32fermi_scrypt_core_kernelB_LG_texILi0ELi2EEvPjjj
    160 bytes stack frame, 400 bytes spill stores, 400 bytes spill loads
ptxas info    : Used 64 registers, 372 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z32fermi_scrypt_core_kernelB_LG_texILi1ELi1EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z32fermi_scrypt_core_kernelB_LG_texILi1ELi1EEvPjjj
    160 bytes stack frame, 384 bytes spill stores, 384 bytes spill loads
ptxas info    : Used 64 registers, 372 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z32fermi_scrypt_core_kernelB_LG_texILi0ELi1EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z32fermi_scrypt_core_kernelB_LG_texILi0ELi1EEvPjjj
    160 bytes stack frame, 400 bytes spill stores, 400 bytes spill loads
ptxas info    : Used 64 registers, 372 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z25fermi_scrypt_core_kernelBILi1EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z25fermi_scrypt_core_kernelBILi1EEvPjj
    40 bytes stack frame, 68 bytes spill stores, 68 bytes spill loads
ptxas info    : Used 64 registers, 364 bytes cmem[0]
ptxas info    : Compiling entry function '_Z25fermi_scrypt_core_kernelBILi0EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z25fermi_scrypt_core_kernelBILi0EEvPjj
    40 bytes stack frame, 68 bytes spill stores, 68 bytes spill loads
ptxas info    : Used 64 registers, 364 bytes cmem[0]
ptxas info    : Compiling entry function '_Z29fermi_scrypt_core_kernelB_texILi1ELi2EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z29fermi_scrypt_core_kernelB_texILi1ELi2EEvPjj
    56 bytes stack frame, 72 bytes spill stores, 72 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z29fermi_scrypt_core_kernelB_texILi0ELi2EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z29fermi_scrypt_core_kernelB_texILi0ELi2EEvPjj
    56 bytes stack frame, 68 bytes spill stores, 68 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z29fermi_scrypt_core_kernelB_texILi1ELi1EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z29fermi_scrypt_core_kernelB_texILi1ELi1EEvPjj
    72 bytes stack frame, 100 bytes spill stores, 100 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z29fermi_scrypt_core_kernelB_texILi0ELi1EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z29fermi_scrypt_core_kernelB_texILi0ELi1EEvPjj
    72 bytes stack frame, 100 bytes spill stores, 100 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0], 1 textures
ptxas info    : Compiling entry function '_Z28fermi_scrypt_core_kernelA_LGILi1EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z28fermi_scrypt_core_kernelA_LGILi1EEvPjjj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0]
ptxas info    : Compiling entry function '_Z28fermi_scrypt_core_kernelA_LGILi0EEvPjjj' for 'sm_80'
ptxas info    : Function properties for _Z28fermi_scrypt_core_kernelA_LGILi0EEvPjjj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 64 registers, 368 bytes cmem[0]
ptxas info    : Compiling entry function '_Z25fermi_scrypt_core_kernelAILi1EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z25fermi_scrypt_core_kernelAILi1EEvPjj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 64 registers, 364 bytes cmem[0]
ptxas info    : Compiling entry function '_Z25fermi_scrypt_core_kernelAILi0EEvPjj' for 'sm_80'
ptxas info    : Function properties for _Z25fermi_scrypt_core_kernelAILi0EEvPjj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 64 registers, 364 bytes cmem[0]
fermi_kernel.cu: In member function ‘virtual bool FermiKernel::bindtexture_1D(uint32_t*, size_t)’:
fermi_kernel.cu:69:64: warning: ‘cudaError_t cudaBindTexture(size_t*, const textureReference*, const void*, const cudaChannelFormatDesc*, size_t)’ is deprecated [-Wdeprecated-declarations]
   69 |     checkCudaErrors(cudaBindTexture(NULL, &texRef1D_4_V, d_V, &channelDesc4, size));
      |                                                                ^
/usr/include/cuda_runtime_api.h:7714:46: note: declared here
 7714 | extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t size __dv(UINT_MAX));
      |                                              ^~~~~~~~~~~~~~~
fermi_kernel.cu:69:64: warning: ‘cudaError_t cudaBindTexture(size_t*, const textureReference*, const void*, const cudaChannelFormatDesc*, size_t)’ is deprecated [-Wdeprecated-declarations]
   69 |     checkCudaErrors(cudaBindTexture(NULL, &texRef1D_4_V, d_V, &channelDesc4, size));
      |                                                                ^
/usr/include/cuda_runtime_api.h:7714:46: note: declared here
 7714 | extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t size __dv(UINT_MAX));
      |                                              ^~~~~~~~~~~~~~~
fermi_kernel.cu: In member function ‘virtual bool FermiKernel::bindtexture_2D(uint32_t*, int, int, size_t)’:
fermi_kernel.cu:85:82: warning: ‘cudaError_t cudaBindTexture2D(size_t*, const textureReference*, const void*, const cudaChannelFormatDesc*, size_t, size_t, size_t)’ is deprecated [-Wdeprecated-declarations]
   85 |     checkCudaErrors(cudaBindTexture2D(NULL, &texRef2D_4_V, d_V, &channelDesc4, width, height, pitch));
      |                                                                                  ^
/usr/include/cuda_runtime_api.h:7773:46: note: declared here
 7773 | extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaBindTexture2D(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t width, size_t height, size_t pitch);
      |                                              ^~~~~~~~~~~~~~~~~
fermi_kernel.cu:85:82: warning: ‘cudaError_t cudaBindTexture2D(size_t*, const textureReference*, const void*, const cudaChannelFormatDesc*, size_t, size_t, size_t)’ is deprecated [-Wdeprecated-declarations]
   85 |     checkCudaErrors(cudaBindTexture2D(NULL, &texRef2D_4_V, d_V, &channelDesc4, width, height, pitch));
      |                                                                                  ^
/usr/include/cuda_runtime_api.h:7773:46: note: declared here
 7773 | extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaBindTexture2D(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t width, size_t height, size_t pitch);
      |                                              ^~~~~~~~~~~~~~~~~
fermi_kernel.cu: In member function ‘virtual bool FermiKernel::unbindtexture_1D()’:
fermi_kernel.cu:91:31: warning: ‘cudaError_t cudaUnbindTexture(const texture<T, dim, readMode>&) [with T = uint4; int dim = 1; cudaTextureReadMode readMode = cudaReadModeElementType; cudaError_t = cudaError]’ is deprecated [-Wdeprecated-declarations]
   91 |     checkCudaErrors(cudaUnbindTexture(texRef1D_4_V));
      |                               ^
/usr/include/cuda_runtime.h:1304:53: note: declared here
 1304 | static __CUDA_DEPRECATED __inline__ __host__ cudaError_t cudaUnbindTexture(
      |                                                     ^~~~~~~~~~~~~~~~~
fermi_kernel.cu:91:31: warning: ‘cudaError_t cudaUnbindTexture(const texture<T, dim, readMode>&) [with T = uint4; int dim = 1; cudaTextureReadMode readMode = cudaReadModeElementType; cudaError_t = cudaError]’ is deprecated [-Wdeprecated-declarations]
   91 |     checkCudaErrors(cudaUnbindTexture(texRef1D_4_V));
      |                               ^
/usr/include/cuda_runtime.h:1304:53: note: declared here
 1304 | static __CUDA_DEPRECATED __inline__ __host__ cudaError_t cudaUnbindTexture(
      |                                                     ^~~~~~~~~~~~~~~~~
fermi_kernel.cu: In member function ‘virtual bool FermiKernel::unbindtexture_2D()’:
fermi_kernel.cu:97:31: warning: ‘cudaError_t cudaUnbindTexture(const texture<T, dim, readMode>&) [with T = uint4; int dim = 2; cudaTextureReadMode readMode = cudaReadModeElementType; cudaError_t = cudaError]’ is deprecated [-Wdeprecated-declarations]
   97 |     checkCudaErrors(cudaUnbindTexture(texRef2D_4_V));
      |                               ^
/usr/include/cuda_runtime.h:1304:53: note: declared here
 1304 | static __CUDA_DEPRECATED __inline__ __host__ cudaError_t cudaUnbindTexture(
      |                                                     ^~~~~~~~~~~~~~~~~
fermi_kernel.cu:97:31: warning: ‘cudaError_t cudaUnbindTexture(const texture<T, dim, readMode>&) [with T = uint4; int dim = 2; cudaTextureReadMode readMode = cudaReadModeElementType; cudaError_t = cudaError]’ is deprecated [-Wdeprecated-declarations]
   97 |     checkCudaErrors(cudaUnbindTexture(texRef2D_4_V));
      |                               ^
/usr/include/cuda_runtime.h:1304:53: note: declared here
 1304 | static __CUDA_DEPRECATED __inline__ __host__ cudaError_t cudaUnbindTexture(
      |                                                     ^~~~~~~~~~~~~~~~~
nvcc -g -O2 -Xptxas "-v" -arch=sm_80 --maxrregcount=32 -I./compat/jansson -o kepler_kernel.o -c kepler_kernel.cu
kepler_kernel.cu(69): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(69): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(69): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(69): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(69): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(69): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(69): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(69): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(149): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(149): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(150): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(150): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(151): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(151): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(154): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(154): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(155): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(155): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(156): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(156): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(309): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(309): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(310): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(310): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(311): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(311): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(324): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(324): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(325): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(325): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(326): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(326): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(343): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(343): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(344): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(344): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(345): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(345): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(352): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(352): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(353): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(353): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(354): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(354): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(399): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(399): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(400): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(400): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(401): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(401): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(409): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(409): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(410): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(410): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(411): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(411): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(428): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(428): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(429): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(429): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(430): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(430): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(438): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(438): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(439): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(439): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(440): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(440): error: no suitable conversion function from "uint4" to "unsigned int" exists

kepler_kernel.cu(715): warning: variable "num_sleeps" was declared but never referenced

kepler_kernel.cu(716): warning: variable "sleeptime" was declared but never referenced

kepler_kernel.cu(102): error: no suitable constructor exists to convert from "int" to "uint4"
          detected during instantiation of "void kepler_scrypt_core_kernelA<ALGO,SCHEME>(const uint32_t *, int, int) [with ALGO=0, SCHEME=ANDERSEN]" 
(722): here

kepler_kernel.cu(102): error: no operator "+" matches these operands
            operand types are: uint4 + int
          detected during instantiation of "void kepler_scrypt_core_kernelA<ALGO,SCHEME>(const uint32_t *, int, int) [with ALGO=0, SCHEME=ANDERSEN]" 
(722): here

kepler_kernel.cu(117): error: no suitable constructor exists to convert from "int" to "uint4"
          detected during instantiation of "void kepler_scrypt_core_kernelA<ALGO,SCHEME>(const uint32_t *, int, int) [with ALGO=0, SCHEME=ANDERSEN]" 
(722): here

kepler_kernel.cu(117): error: no operator "+" matches these operands
            operand types are: uint4 + int
          detected during instantiation of "void kepler_scrypt_core_kernelA<ALGO,SCHEME>(const uint32_t *, int, int) [with ALGO=0, SCHEME=ANDERSEN]" 
(722): here

kepler_kernel.cu(560): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(560): error: no operator "&" matches these operands
            operand types are: uint4 & uint32_t

kepler_kernel.cu(591): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(591): error: no operator "&" matches these operands
            operand types are: uint4 & uint32_t

kepler_kernel.cu(600): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(600): error: no operator "&" matches these operands
            operand types are: uint4 & uint32_t

kepler_kernel.cu(624): error: no suitable constructor exists to convert from "int" to "uint4"

kepler_kernel.cu(624): error: no operator "&" matches these operands
            operand types are: uint4 & uint32_t

80 errors detected in the compilation of "kepler_kernel.cu".
make[2]: *** [Makefile:1163: kepler_kernel.o] Error 1
make[2]: Leaving directory '/data/Install/UNIX/CudaMiner-master'
make[1]: *** [Makefile:701: all-recursive] Error 1
make[1]: Leaving directory '/data/Install/UNIX/CudaMiner-master'
make: *** [Makefile:419: all] Error 2

What shall I do?