从GTX 680升级到titan出错

加速计算 CUDA
1

GTX 680下程序运行正常。 其中block数为5， thread per block为1024。
升级到titan后，显示系统不识别titan，更新驱动程序至NVIDIA-Linux-x86_64-319.23后，可以识别系统。
但是程序无法运行。
将thread per block改为256后，程序又可以运行了。

问题：为什么titan支持的thread数量反而少了?

2

LZ您好，TITAN可支持的thread per block不比680少的，您可能是其他方面的问题。

请您再检查下。

祝您好运~

3

可是软件根本没动（ubuntu 11.10），换一块卡就不行了。实在想不出来是为什么。难道驱动对titan支持不好？

4

LZ您好，我无法凭空判断的。
geforce 680的计算能力是3.0，TITAN是3.5。
根据CUDA C Programming Guide Appendix.F，两者的Maximum number of threads per block均为1024。此数据由NVIDIA官方保证可靠性。

所以唯有建议您仔细检查，并升级至最新版的驱动和CUDA Toolkit了。
如果驱动有BUG，欢迎您及时反馈，论坛有常驻NV原厂支持，可以向原厂反馈BUG的。

祝您好运~

5

特别说一下，如果您使用的CUDA Toolkit是5.0之前的版本，请升级至5.0。
CUDA 4.2版本使用计算能力3.0的卡也许无问题，但是3.5的卡并无法保证。以及，即便是3.0的卡，也建议使用CUDA 5.0的。

6

[

我忘了说了, 我现在用的是cuda tookit 5.0. 编译选项中所有的计算能力都选上了。明天试试只选计算能力3.5 或者只选3.0，看看结果怎么样。

7

编译选项多选一些，一般是不影响的，会内部生成多分可执行代码/PTX代码的，根据实际硬件信息执行相应版本的。

如果没有对应硬件的二进制代码，也会从PTX中jit编译一份出来的。

不过您还是多加尝试和寻找吧。

8

谢了。我发现使用计算能力3.5编译才会出错，如果只使用3.0不会出错。换用cuda 5.5也是一样的。

9

发代码，发你遇到什么错误。

光说“换上Titan程序无法运行”不能说明什么。
你启动kernel后，然后执行一次cudaDeviceSynchronize(); 看看返回值是什么。

10

我用getlasterror返回的错误是too many resource requested for launch!
用cudaDeviceSynchronize返回值是0
下面是我使用的Makefile，我担心是我的Makefile使用的不对。因为我使用Nsight编译没有任何问题，但使用我自己的Makefile就会出问题。
我坚持使用自己的Makefile的原因是Nsight编译出的代码速度很慢，即使改成Release，把-G -g去掉也没什么效果。
这个是我使用的Makefile，不知道有没有什么问题

OS Name (Linux or Darwin)

OSUPPER = $(shell uname -s 2>/dev/null | tr [:lower:] [:upper:])
OSLOWER = $(shell uname -s 2>/dev/null | tr [:upper:] [:lower:])

Flags to detect 32-bit or 64-bit OS platform

OS_SIZE = $(shell uname -m | sed -e “s/i.86/32/” -e “s/x86_64/64/”)
OS_ARCH = $(shell uname -m | sed -e “s/i386/i686/”)

These flags will override any settings

ifeq ($(i386),1)
OS_SIZE = 32
OS_ARCH = i686
endif

ifeq ($(x86_64),1)
OS_SIZE = 64
OS_ARCH = x86_64
endif

Flags to detect either a Linux system (linux) or Mac OSX (darwin)

DARWIN = $(strip $(findstring DARWIN, $(OSUPPER)))

Location of the CUDA Toolkit binaries and libraries

CUDA_PATH ?= /usr/local/cuda-5.5
CUDA_INC_PATH ?= $(CUDA_PATH)/include
CUDA_BIN_PATH ?= $(CUDA_PATH)/bin
ifneq ($(DARWIN),)
CUDA_LIB_PATH ?= $(CUDA_PATH)/lib
else
ifeq ($(OS_SIZE),32)
CUDA_LIB_PATH ?= $(CUDA_PATH)/lib
else
CUDA_LIB_PATH ?= $(CUDA_PATH)/lib64
endif
endif

Common binaries

NVCC ?= $(CUDA_BIN_PATH)/nvcc
GCC ?= gcc

Extra user flags

EXTRA_NVCCFLAGS ?=
EXTRA_LDFLAGS ?=
EXTRA_CCFLAGS ?=

CUDA code generation flags

GENCODE_SM10 := -gencode arch=compute_10,code=sm_10
GENCODE_SM20 := -gencode arch=compute_20,code=sm_20
GENCODE_SM30 := -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35
#-gencode arch=compute_30,code=sm_30

-gencode arch=compute_35,code=sm_35

#GENCODE_FLAGS := $(GENCODE_SM10) $(GENCODE_SM20) $(GENCODE_SM30)
GENCODE_FLAGS := $(GENCODE_SM30)

OS-specific build flags

ifneq ($(DARWIN),)
LDFLAGS := -Xlinker -rpath $(CUDA_LIB_PATH) -L$(CUDA_LIB_PATH) -lcudart
CCFLAGS := -arch $(OS_ARCH)
else
ifeq ($(OS_SIZE),32)
LDFLAGS := -L$(CUDA_LIB_PATH) -lcudart
CCFLAGS := -m32
else
LDFLAGS := -L$(CUDA_LIB_PATH) -lcudart
CCFLAGS := -m64
endif
endif

OS-architecture specific flags

ifeq ($(OS_SIZE),32)
NVCCFLAGS := -m32
else
NVCCFLAGS := -m64
endif

Common includes and paths for CUDA

INCLUDES := -I$(CUDA_INC_PATH) -I. -I… -I…/…/common/inc
LDFLAGS += $(LIBPATH_OPENGL)

Target rules

all: build

build: main

main: GPU.o filetest.o selftest.o softtest.o tool.o speedtest.o EllipticCurve.o Mpi.o sm3hash.o sm4.o sm2.o sm2cpp.o
g++ $(CCFLAGS) -o $@ $+ $(LDFLAGS) $(EXTRA_LDFLAGS) -pthread -march=i686

softtest.o : test/softtest.cpp
g++ -c test/softtest.cpp -o softtest.o

tool.o : test/tool.cpp
g++ -c test/tool.cpp -o tool.o

speedtest.o : test/speedtest.cu
$(CUDA_BIN_PATH)/nvcc -c test/speedtest.cu -o speedtest.o $(NVCCFLAGS) $(EXTRA_NVCCFLAGS) $(GENCODE_FLAGS) $(INCLUDES) --ptxas-options=-v

selftest.o : test/selftest.cpp
g++ -c test/selftest.cpp -o selftest.o

filetest.o : test/filetest.cpp
g++ -c test/filetest.cpp -o filetest.o

GPU.o : test/GPU.cpp
g++ -c test/GPU.cpp -o GPU.o
sm2.o: GPU/sm2.cu
$(CUDA_BIN_PATH)/nvcc -c GPU/sm2.cu -o sm2.o $(NVCCFLAGS) $(EXTRA_NVCCFLAGS) $(GENCODE_FLAGS) $(INCLUDES) --ptxas-options=-v
sm4.o: GPU/sm4.cu
$(CUDA_BIN_PATH)/nvcc -c GPU/sm4.cu -o sm4.o $(NVCCFLAGS) $(EXTRA_NVCCFLAGS) $(GENCODE_FLAGS) $(INCLUDES) --ptxas-options=-v
sm2cpp.o: GPU/sm2cpp.cu
$(CUDA_BIN_PATH)/nvcc -c GPU/sm2cpp.cu -o sm2cpp.o $(NVCCFLAGS) $(EXTRA_NVCCFLAGS) $(GENCODE_FLAGS) $(INCLUDES) --ptxas-options=-v
EllipticCurve.o : algorithm/EllipticCurve.cpp
g++ -c algorithm/EllipticCurve.cpp -o EllipticCurve.o
sm3hash.o : algorithm/sm3hash.cpp
g++ -c algorithm/sm3hash.cpp -o sm3hash.o
Mpi.o : algorithm/Mpi.cpp
g++ -c algorithm/Mpi.cpp -o Mpi.o
OBJS = GPU.o sm2.o sm2cpp.o

clean :
rm -f *.o

run: build
./GPUsign

11

Nsight和这个Makefile编译出的代码，性能差距大约有十倍，这也不知道为什么。我只会修改release debug -G -g 等选项，不知道是否还需要修改别的？

12

LZ您好，我认为您的这个现象应该有其他原因，而不是说您的代码无法以3.5的计算能力标准编译。

13

LZ您好，我来简要说一下返回值查错的问题。

论坛曾经有帖子详细讨论了cudaGetLastError()和通过返回值检查问题的方法，大致结论为：

1：cudaGetLastError()自己维护了一个缓冲，可以缓冲一次错误，以及可以通过调用该函数将缓冲中的内容释放出来。因为该缓冲，所以有时得到的错误报告会对人产生误导。

2：对于同步完成的函数，如常规的cudaMemcpy等，可以直接检查其返回值，判断执行成功与否。

3：对于kernel调用等异步函数，可以在紧跟该函数之后，使用cudaDeviceSynchronize()，并检查cudaDeviceSynchronize()的返回值，这个返回值表示了cudaDeviceSynchronize()之前紧挨着这个异步函数的执行情况。

4：上述2：，3：都无缓冲机制，需要您每步都做检查。

需要指出的是，无论是使用IDE配合相关的模板，还是您自己写Makefile，其原理都是一致的，都是使用NVCC编译device端代码，原host的编译器编译host端代码，最后再合在一起。在参数对等的情况下，并不会出现一个比另外一个慢10倍的情况。

同时您指出使用IDE配合模板的方法编译是无问题的，使用您的Makefile会出问题，这多少也说明了您前面表示的“nsight编译结果慢”（IDE+模板编译结果慢）并没有做出合适的对比。

大致如此，我不用linux也不用Makefile，无法就您书写细节给出建议了，请其他人补充。

祝您好运~

修改若干错别字。

14

自发自回。我已经找到怎么设置release模式了，除了上面说的，还有在project视图里，右键Release下面的可执行文件 run-> 1 local c/c++ Application，以后release模式下的可执行文件才加入到run中。之前一直是在debug模式下。
而且我的Makefile文件确实是有问题的，表现为：
使用我的Makefile, 启动10000个线程时，程序计算结果有误
使用Nsight Release模式，程序计算结果无误
目前还不知道是哪里写的有问题