[GPU计算]深入浅出谈CUDA技术_42pp.pdf

深入浅出谈深入浅出谈CUDA发表时间：2008-11-21“CUDA是NVIDIA的GPGPU(非图像计算)模型，它使用C语言为基础，可以直接以大多数人熟悉的C语言，写出在显示芯片上执行的程序，而不需要去学习特定的显示芯片的指令或是特殊的结构。”GPU就是早期的图形运算处理器，但是随着发展，拥有较高带宽的GPU被人们开发出来进行除开图形运算的其它运算，提高程序的执行效率即GPGPU；CUDA是什么？是什么？编者注：NVIDIA的GeFoce8800GTX发布后，它的通用计算架构CUDA经过一年多的推广后，现在已经在有相当多的论文发表，在商业应用软件等方面也初步出现了视频编解码、金融、地质勘探、科学计算等领域的产品，是时候让我们对其作更深一步的了解。为了让大家更容易了解CUDA，我们征得Hotball的本人同意，发表他最近亲自撰写的本文。这篇文章的特点是深入浅出，也包含了hotball本人编写一些简单CUDA程序的亲身体验，对于希望了解CUDA的读者来说是非常不错的入门文章，PCINLIFE对本文的发表没有作任何的删减，主要是把一些台湾的词汇转换成大陆的词汇以及作了若干编者注的注释。现代的显示芯片已经具有高度的可程序化能力，由于显示芯片通常具有相当高的内存带宽，以及大量的执行单元，因此开始有利用显示芯片来帮助进行一些计算工作的想法，即GPGPU。CUDA即是NVIDIA的GPGPU模型。NVIDIA的新一代显示芯片，包括GeForce8系列及更新的显示芯片都支持CUDA。NVIDIA免费提供CUDA的开发工具（包括Windows版本和Linux版本）、程序范例、文件等等，可以在CUDAZone下载。GPGPU的优缺点的优缺点使用显示芯片来进行运算工作，和使用CPU相比，主要有几个好处：1.显示芯片通常具有更大的内存带宽（前端总线）。例如，NVIDIA的GeForce8800GTX具有超过50GBs的内存带宽，而目前高阶CPU的内存带宽则在10GBs左右。对于带宽的解释对于带宽的解释详细见文档什么是带宽。详细见文档什么是带宽。2.显示芯片具有更大量的执行单元。例如GeForce8800GTX具有128个streamprocessors，频率为1.35GHz。CPU频率通常较高，但是执行单元的数目则要少得多。3.和高阶CPU相比，显卡的价格较为低廉。例如目前一张GeForce8800GT包括512MB内存的价格，和一颗2.4GHz四核心CPU的价格相若。当然，使用显示芯片也有它的一些缺点：1.显示芯片的运算单元数量很多，因此对于不能高度并行化的工作，所能带来的帮助就不大。2.显示芯片目前通常只支持32bits浮点数，且多半不能完全支持IEEE754规格，有些运算的精确度可能较低。目前许多显示芯片并没有分开的整数运算单元，因此整数运算的效率较差。3.显示芯片通常不具有分支预测等复杂的流程控制单元，因此对于具有高度分支的程序，效率会比较差。4.目前GPGPU的程序模型仍不成熟，也还没有公认的标准。例如NVIDIA和AMDATI就有各自不同的程序模型。整体来说，显示芯片的性质类似streamprocessor，适合一次进行大量相同的工作。CPU则比较有弹性，能同时进行变化较多的工作。CUDA架构架构CUDA是NVIDIA的GPGPU模型，它使用C语言为基础，可以直接以大多数人熟悉的C语言，写出在显示芯片上执行的程序，而不需要去学习特定的显示芯片的指令或是特殊的结构。在CUDA的架构下，一个程序分为两个部份：host端和device端。Host端是指在CPU上执行的部份，而device端则是在显示芯片上执行的部份。Device端的程序又称为kernel（核心）。通常host端程序会将数据准备好后，复制到显卡的内存中，再由显示芯片执行device端程序，完成后再由host端程序将结果从显卡的内存中取回。由于CPU存取显卡内存时只能透过PCIExpress接口，因此速度较慢（PCIExpressx16的理论带宽是双向各4GBs），因此不能太常进行这类动作，以免降低效率。在CUDA架构下，显示芯片执行时的最小单位是thread(线程线程)。数个thread可以组成一个block(块块)。一个block中的thread能存取同一块共享的内存，而且可以快速进行同步的动作。每一个block所能包含的thread数目是有限的。不过，执行相同程序的block，可以组成grid(格格子子)。不同block中的thread无法存取同一个共享的内存，因此无法直接互通或进行同步。因此，不同block中的thread能合作的程度是比较低的。不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的thread数目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个block中的thread顺序执行，而非同时执行。不同的grid则可以执行不同的程序（即kernel）。Grid、block和thread的关系，如下图所示：每个thread都有自己的一份register和localmemory的空间。同一个block中的每个thread则有共享的一份sharememory。此外，所有的thread（包括不同block的thread）都共享一份globalmemory、constantmemory、和texturememory。不同的grid则有各自的globalmemory、constantmemory和texturememory。这些不同的内存的差别，会在之后讨论。执行模式执行模式由于显示芯片大量并行计算的特性，它处理一些问题的方式，和一般CPU是不同的。主要的特点包括：1.内存存取latency的问题：CPU通常使用cache来减少存取主内存的次数，以避免内存latency影响到执行效率。显示芯片则多半没有cache（或很小），而利用并行化执行的方式来隐藏内存的latency（即，当第一个thread需要等待内存读取结果时，则开始执行第二个thread，依此类推）。2.分支指令的问题：CPU通常利用分支预测等方式来减少分支指令造成的pipelinebubble。显示芯片则多半使用类似处理内存latency的方式。不过，通常显示芯片处理分支的效率会比较差。因此，最适合利用CUDA处理的问题，是可以大量并行化的问题，才能有效隐藏内存的latency，并有效利用显示芯片上的大量执行单元。使用CUDA时，同时有上千个thread在执行是很正常的。因此，如果不能大量并行化的问题，使用CUDA就没办法达到最好的效率了。CUDAToolkit的安装的安装目前NVIDIA提供的CUDAToolkit（可从这里下载）支持Windows（32bits及64bits版本）及许多不同的Linux版本。CUDAToolkit需要配合CC+compiler。在Windows下，目前只支持VisualStudio7.x及VisualStudio8（包括免费的VisualStudioC+2005Express）。VisualStudio6和gcc在Windows下是不支援的。在Linux下则只支援gcc。这里简单介绍一下在Windows下设定并使用CUDA的方式。下载及安装下载及安装在Windows下，CUDAToolkit和CUDASDK都是由安装程序的形式安装的。CUDAToolkit包括CUDA的基本工具，而CUDASDK则包括许多范例程序以及链接库。基本上要写CUDA的程序，只需要安装CUDAToolkit即可。不过CUDASDK仍值得安装，因为里面的许多范例程序和链接库都相当有用。CUDAToolkit安装完后，预设会安装在C:CUDA目录里。其中包括几个目录：bin-工具程序及动态链接库doc-文件include-header檔lib-链接库档案open64-基于Open64的CUDAcompilersrc-一些原始码安装程序也会设定一些环境变量，包括：CUDA_BIN_PATH-工具程序的目录，默认为C:CUDAbinCUDA_INC_PATH-header文件的目录，默认为C:CUDAincCUDA_LIB_PATH-链接库文件的目录，默认为C:CUDAlib在在VisualStudio中使用中使用CUDACUDA的主要工具是nvcc，它会执行所需要的程序，将CUDA程序代码编译成执行档(或object檔)。在VisualStudio下，我们透过设定custombuildtool的方式，让VisualStudio会自动执行nvcc。这里以VisualStudio2005为例：1.首先，建立一个Win32Console模式的project（在ApplicationSettings中记得勾选Emptyproject），并新增一个档案，例如main.cu。2.在main.cu上右键单击，并选择Properties。点选General，确定Tool的部份是选择CustomBuildTool。3.选择CustomBuildStep，在CommandLine使用以下设定：oRelease模式模式：$(CUDA_BIN_PATH)nvcc.-ccbin$(VCInstallDir)bin-c-DWIN32-D_CONSOLE-D_MBCS-XcompilerEHscW3nologoWp64O2ZiMT-I$(CUDA_INC_PATH)-o$(ConfigurationName)$(Name).obj$(FileName)oDebug模式模式：$(CUDA_BIN_PATH)nvcc.-ccbin$(VCInstallDir)bin-c-D_DEBUG-DWIN32-D_CONSOLE-D_MBCS-XcompilerEHscW3nologoWp64OdZiRTC1MTd-I$(CUDA_INC_PATH)-o$(ConfigurationName)$(Name).obj$(FileName)4.如果想要使用软件仿真的模式，可以新增两个额外的设定：oEmuRelease模式模式：$(CUDA_BIN_PATH)nvcc.-ccbin$(VCInstallDir)bin-deviceemu-c-DWIN32-D_CONSOLE-D_MBCS-XcompilerEHscW3nologoWp64O2ZiMT-I$(CUDA_INC_PATH)-o$(ConfigurationName)$(Name).obj$(FileName)oEmuDebug模式模式：$(CUDA_BIN_PATH)nvcc.-ccbin$(VCInstallDir)bin-deviceemu-c-D_DEBUG-DWIN32-D_CONSOLE-D_MBCS-XcompilerEHscW3nologoWp64OdZiRTC1MTd-I$(CUDA_INC_PATH)-o$(ConfigurationName)$(Name).obj$(FileName)5.对所有的配置文件，在CustomBuildStep的Outputs中加入$(ConfigurationName)$(Name).obj。6.选择project，右键单击选择Properties，再点选Linker。对所有的配置文件修改以下设定：oGeneralEnableIncrementalLinking：NooGeneralAdditionalLibraryDirectories：$(CUDA_LIB_PATH)oAdditionalDependencies：cudart.lib这样应该就可以直接在VisualStudio的IDE中，编辑CUDA程序后，直接build以及执行程序了。第一个第一个CUDA程序程序CUDA目前有两种不同的API：RuntimeAPI和DriverAPI，两种API各有其适用的范围。由于runtimeAPI较容易使用，一开始我们会以runetimeAPI为主。CUDA的初始化的初始化首先，先建立一个档案first_cuda.cu。如果是使用VisualStudio的话，则请先按照这里的设定方式设定project。要使用runtimeAPI的时候，需要includecuda_runtime.h。所以，在程序的最前面，加上#include#include接下来是一个InitCUDA函式，会呼叫runtimeAPI中，有关初始化CUDA的功能：boolInitCUDA()intcountcudaGetDeviceCount(count)if(count=0)fprintf(stderrThereisnodevice.n)returnfalseintifor(i=0i=1)breakif(i=count)fprintf(stderrThereisnodevicesupportingCUDA1.x.n)returnfalsecudaSetDevice(i)returntrue这个函式会先呼叫cudaGetDeviceCount函式，取得支持CUDA的装置的数目。如果系统上没有支持CUDA的装置，则它会传回1，而device0会是一个仿真的装置，但不支持CUDA1.0以上的功能。所以，要确定系统上是否有支持CUDA的装置，需要对每个device呼叫cudaGetDeviceProperties函式，取得装置的各项数据，并判断装置支持的CUDA版本（prop.major和prop.minor分别代表装置支持的版本号码，例如1.0则prop.major为1而prop.minor为0）。透过cudaGetDeviceProperties函式可以取得许多数据，除了装置支持的CUDA版本之外，还有装置的名称、内存的大小、最大的thread数目、执行单元的频率等等。详情可参考NVIDIA的CUDAProgrammingGuide。在找到支持CUDA1.0以上的装置之后，就可以呼叫cudaSetDevice函式，把它设为目前要使用的装置。最后是main函式。在main函式中我们直接呼叫刚才的InitCUDA函式，并显示适当的讯息：intmain()if(!InitCUDA()return0printf(CUDAinitialized.n)return0这样就可以利用nvcc来compile这个程序了。使用VisualStudio的话，若按照先前的设定方式，可以直接BuildProject并执行。nvcc是CUDA的compile工具，它会将.cu檔拆解出在GPU上执行的部份，及在host上执行的部份，并呼叫适当的程序进行compile动作。在GPU执行的部份会透过NVIDIA提供的compiler编译成中介码，而host执行的部份则会透过系统上的C+compiler编译（在Windows上使用VisualC+而在Linux上使用gcc）。编译后的程序，执行时如果系统上有支持CUDA的装置，应该会显示CUDAinitialized.的讯息，否则会显示相关的错误讯息。利用利用CUDA进行运算进行运算到目前为止，我们的程序并没有做什么有用的工作。所以，现在我们加入一个简单的动作，就是把一大堆数字，计算出它的平方和。首先，把程序最前面的include部份改成：#include#include#include#defineDATA_SIZE1048576intdataDATA_SIZE并加入一个新函式GenerateNumbers：voidGenerateNumbers(intnumberintsize)for(inti=0i(参数.)呼叫完后，还要把结果从显示芯片复制回主内存上。在main函式中加入以下的程序：sumOfSquares(gpudataresult)intsumcudaMemcpy(sumresultsizeof(int)cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)printf(sum:%dnsum)因为这个程序只使用一个thread，所以block数目、thread数目都是1。我们也没有使用到任何sharedmemory，所以设为0。编译后执行，应该可以看到执行的结果。为了确定执行的结果正确，我们可以加上一段以CPU执行的程序代码，来验证结果：sum=0for(inti=0i(gpudataresulttime)intsumclock_ttime_usedcudaMemcpy(sumresultsizeof(int)cudaMemcpyDeviceToHost)cudaMemcpy(time_usedtimesizeof(clock_t)cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)printf(sum:%dtime:%dnsumtime_used)编译后执行，就可以看到执行所花费的时间了。如果计算实际运行时间的话，可能会注意到它的执行效率并不好。这是因为我们的程序并没有利用到CUDA的主要的优势，即并行化执行。在下一段文章中，会讨论如何进行优化的动作。改良第一个改良第一个CUDA程序程序在上一篇文章中，我们做了一个计算一大堆数字的平方和的程序。不过，我们也提到这个程序的执行效率并不理想。当然，实际上来说，如果只是要做计算平方和的动作，用CPU做会比用GPU快得多。这是因为平方和的计算并不需要太多运算能力，所以几乎都是被内存带宽所限制。因此，光是把数据复制到显卡内存上的这个动作，所需要的时间，可能已经和直接在CPU上进行计算差不多了。不过，如果进行平方和的计算，只是一个更复杂的计算过程的一部份的话，那么当然在GPU上计算还是有它的好处的。而且，如果数据已经在显卡内存上（例如在GPU上透过某种算法产生），那么，使用GPU进行这样的运算，还是会比较快的。刚才也提到了，由于这个计算的主要瓶颈是内存带宽，所以，理论上显卡的内存带宽是相当大的。这里我们就来看看，倒底我们的第一个程序，能利用到多少内存带宽。程序的并行化程序的并行化我们的第一个程序，并没有利用到任何并行化的功能。整个程序只有一个thread。在GeForce8800GT上面，在GPU上执行的部份（称为kernel）大约花费640M个频率。GeForce8800GT的执行单元的频率是1.5GHz，因此这表示它花费了约0.43秒的间。1M个32bits数字的数据量是4MB，因此，这个程序实际上使用的内存带宽，只有9.3MBs左右！这是非常糟糕的表现。为什么会有这样差的表现呢？这是因为GPU的架构特性所造成的。在CUDA中，一般的数据复制到的显卡内存的部份，称为globalmemory。这些内存是没有cache的，而且，存取globalmemory所需要的时间（即latency）是非常长的，通常是数百个cycles。由于我们的程序只有一个thread，所以每次它读取globalmemory的内容，就要等到实际读取到数据、累加到sum之后，才能进行下一步。这就是为什么它的表现会这么的差。由于globalmemory并没有cache，所以要避开巨大的latency的方法，就是要利用大量的threads。假设现在有大量的threads在同时执行，那么当一个thread读取内存，开始等待结果的时候，GPU就可以立刻切换到下一个thread，并读取下一个内存位置。因此，理想上当thread的数目够多的时候，就可以完全把globalmemory的巨大latency隐藏起来了。要怎么把计算平方和的程序并行化呢？最简单的方法，似乎就是把数字分成若干组，把各组数字分别计算平方和后，最后再把每组的和加总起来就可以了。一开始，我们可以把最后加总的动作，由CPU来进行。首先，在first_cuda.cu中，在#defineDATA_SIZE的后面增加一个#define，设定thread的数目：#defineDATA_SIZE1048576#defineTHREAD_NUM256接着，把kernel程序改成：_global_staticvoidsumOfSquares(intnumintresultclock_ttime)constinttid=threadIdx.xconstintsize=DATA_SIZETHREAD_NUMintsum=0inticlock_tstartif(tid=0)start=clock()for(i=tidsizei(gpudataresulttime)intsumTHREAD_NUMclock_ttime_usedcudaMemcpy(sumresultsizeof(int)THREAD_NUMcudaMemcpyDeviceToHost)cudaMemcpy(time_usedtimesizeof(clock_t)cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)cudaFree(time)可以注意到我们在呼叫sumOfSquares函式时，指定THREAD_NUM为thread的数目。最后，在CPU端把计算好的各组数据的平方和进行加总：intfinal_sum=0for(inti=0ithread2-.因此，在同一个thread中连续存取内存，在实际执行时反而不是连续了。要让实际执行结果是连续的存取，我们应该要让thread0读取第一个数字，thread1读取第二个数字依此类推。所以，我们可以把kernel程序改成如下：_global_staticvoidsumOfSquares(intnumintresultclock_ttime)constinttid=threadIdx.xintsum=0inticlock_tstartif(tid=0)start=clock()for(i=tidi(gpudataresulttime)intsumTHREAD_NUMBLOCK_NUMclock_ttime_usedBLOCK_NUM2cudaMemcpy(sumresultsizeof(int)THREAD_NUMBLOCK_NUMcudaMemcpyDeviceToHost)cudaMemcpy(time_usedtimesizeof(clock_t)BLOCK_NUM2cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)cudaFree(time)intfinal_sum=0for(inti=0i(gpudataresulttime)intsumBLOCK_NUMclock_ttime_usedBLOCK_NUM2cudaMemcpy(sumresultsizeof(int)BLOCK_NUMcudaMemcpyDeviceToHost)cudaMemcpy(time_usedtimesizeof(clock_t)BLOCK_NUM2cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)cudaFree(time)intfinal_sum=0for(inti=0i=1_syncthreads()这样同时也省去了mask变数。因此，这个版本的执行的效率就可以再提高一些。在GeForce8800GT上，这个版本执行的时间是约137Kcycles。当然，这时差别已经很小了。如果还要再提高效率，可以把树状加法整个展开：if(tid(acnbcnccnn)cudaMemcpy2D(csizeof(float)ldcccsizeof(float)nsizeof(float)nncudaMemcpyDeviceToHost)cudaFree(ac)cudaFree(bc)cudaFree(cc)end=clock()returnend-start这个函式相当单纯，就是在显卡内存中配置存放矩阵的内存，然后把主内存中的矩阵数据复制到显卡内存上。不过，因为我们的矩阵乘法函式可以指定pitch（即lda、ldb、和ldc），所以如果用一般的cudaMemcpy函式来复制内存的话，会需要每个row都分开复制，那会需要呼叫很多次cudaMemcpy函式，会使效率变得很差。因此，在这里我们用了一个新的cudaMemcpy2D函式，它是用来复制二维数组，可以指定数组的pitch。这样就可以透过一次函数调用就可以了。进行计算的kernel如下：_global_staticvoidmatMultCUDA(constfloatasize_tldaconstfloatbsize_tldbfloatcsize_tldcintn)constinttid=threadIdx.xconstintbid=blockIdx.xconstintidx=bidblockDim.x+tidconstintrow=idxnconstintcolumn=idx%nintiif(row(acnbcnccnn)kernel的部份则改成：_global_staticvoidmatMultCUDA(constfloatasize_tldaconstfloatbsize_tldbfloatcsize_tldcintn)extern_shared_floatdataconstinttid=threadIdx.xconstintrow=blockIdx.xintijfor(i=tidi(acpitch_asizeof(float)bcpitch_bsizeof(float)ccpitch_csizeof(float)n)同样的，把计算结果复制回到主内存时，也要使用传回的宽度值：cudaMemcpy2D(csizeof(float)ldcccpitch_csizeof(float)nncudaMemcpyDeviceToHost)这样就完成了。Kernel部份则不需要修改。这样的修改有多大的效果呢？在GeForce8800GT上执行，结果如下：Maxerror:1.19209e-007Averageerror:4.22751e-008Timeused:0.1250(16.00GFLOPS)可以看到，执行速度又再大幅提高了三倍多！而这只是把内存的配置方式稍微修改一下而已。虽然执行速度提高了很多，但是，和前面提到的理论值相比，其实还是有相当的差距。这是因为，前面也提到过，这样的做法需要n3+n2次的内存读取，和n2次的内存写入动作。由于n=1000，每个数字的大小是32bits，所以总共的内存存取数据量约为4GB。除以实际执行的时间0.125秒，得到的带宽数值是约32GBs，这已经接近GeForce8800GT显卡内存的带宽了。由于我们计算时间的时候，把配置内存、以及数据的复制动作也计算进去，因此实际上花费在kernel的时间是更短的（约0.09秒）。因此，可以很明显的看出，这个程序的效率，是受限于内存带宽的。进一步的改良进一步的改良上一节的结论显示出，矩阵乘法的程序，效率是受限于内存带宽的。那有没有办法降低内存的存取次数呢？答案当然是有的，不然就不会有这一节了:)要进一步降低内存带宽的使用，可以注意到，在上一节的方法中，虽然A矩阵的存取次数被减至最低，但是B矩阵的存取次数并没有减少。这是因为我们只将A矩阵的row加载到sharedmemory中，但是B矩阵的column也是有被重复使用的。理想上应该也可以避免重复加载才对。不过，由于B矩阵的column使用的时机，和A矩阵的row是不同的，所以并不能直接这样做。解决方法是blocking。也就是把整个矩阵乘法的动作，切割成很多小矩阵的乘法。例如，要计算C矩阵的(00)(1515)的值，可以把它想成：A(015015)B(015015)+A(0151631)B(1631015)+A(0153247)B(3247015)+.这样一来，我们就可以把两个小矩阵加载到sharedmemory，则小矩阵本身的乘法就不需要再存取任何外部的内存了！这样一来，假设小矩阵的大小是k，则实际上需要的内存存取次数就会变成约2k2(nk)3=2n3k。由于目前CUDA每个block的thread数目最多是512，因此k=16似乎是一个相当理想的数字（共256个threads）。因此，对于一个n=1000的矩阵来说，我们可以把内存存取的量减少到约500MB，也就是上一节的存取量的18。理论上，这样应该可以让效率提高八倍（假设没有遇到别的瓶颈）。为了方便进行区块的计算，我们让每个block有16x16个threads，再建立(n16)x(n16)个blocks。把呼叫kernel的地方改成：intbx=(n+BLOCK_SIZE-1)BLOCK_SIZEdim3blocks(bxbx)dim3threads(BLOCK_SIZEBLOCK_SIZE)matMultCUDA(acpitch_asizeof(float)bcpitch_bsizeof(float)ccpitch_csizeof(float)n)BLOCK_SIZE则是定义成16。dim3是CUDA的一种数据型态，表示一个3D的向量。在这里，我们透过dim3来建立16x16个threads的block，和(n16)x(n16)个blocks。Kernel程序的部份，则改成：_global_staticvoidmatMultCUDA(constfloatasize_tldaconstfloatbsize_tldbfloatcsize_tldcintn)_shared_floatmatABLOCK_SIZEBLOCK_SIZE_shared_floatmatBBLOCK_SIZEBLOCK_SIZEconstinttidc=threadIdx.xconstinttidr=threadIdx.yconstintbidc=blockIdx.xBLOCK_SIZEconstintbidr=blockIdx.yBLOCK_SIZEintijfloatresults=0floatcomp=0for(j=0j正文CUDA计算初步计算初步发表时间：2008-08-29CUDA计算体系是一个新鲜玩艺俺也来凑个热闹刚好硬件能对付DELL1420自带的显卡是8400MGS属于比较低端的支持CUDA的显卡。要查看你的显卡是否支持CUDA请访问官网。CUDA计算体系是一个新鲜玩艺俺也来凑个热闹刚好硬件能对付DELL1420自带的显卡是8400MGS属于比较低端的支持CUDA的显卡。要查看你的显卡是否支持CUDA请访问官网。DELL1420官方驱动是不带CUDA的可能一般的OEM驱动都不会默认安装CUDA。NVIDIA的驱动更新太快装了几个版本的驱动都不能打开8400MGS这个显卡的CUDA支持无奈选用177.79版本的驱动配合修改inf来强制打开.终于可以用了.运行devicequery例子可以得到如下输出代表CUDA正常工作了(不容易啊，用了几个小时):Thereis1devicesupportingCUDADevice0:GeForce8400MGSMajorrevisionnumber:1Minorrevisionnumb