1.CUDA5.5 SDK中利用共享存储器编写的矩阵乘法分为runtime api版本和driver api版本，理论上driver api版本的运行速度是快于runtime api 版本的运行速度，在debug配置下编译是符合理论情况的，但是在release配置下runtime api编写的矩阵乘法性能大幅上升，但driver api编写的矩阵乘法性能没改变，这是什么问题，难道release配置下编译器不对driver api编写的程序进行优化吗？求版主解惑？
2.顺便问一下内建变量占用的存储资源是什么（local memory还是别的什么）？
楼主您需要了解driver api和runtime api的区别，
这些区别里面最重要的一点就是driver api是要单独编译kernel的，你的kernel的编译质量，在你分开编译kernel的过程中就已经固定了（其实这么说不确切，不过楼主可以先这么看着），而您的host code的项目代码，使用debug还是release配置，均不会影响您的kernel, 因此性能不变。
关于您咨询的第二个问题，一般情况下，一些是特殊的寄存器(special register), 例如threadIdx.x这种，一些是在constant memory里，例如blockDim.x这种。但他们不是从local memory中寻找的。
感谢来访。
您所说的分开编译的kernel的过程中已经固定是不是指显示加载地ptx已经是固定？
我安照这种思路将项目文件中.ptx删除，然后在vs中对kernel进行release配置下的编译，然后再在.cpp进行release配置下的编译。还是没改变。是我的理解哪里还不到位吗？求您为我解惑。
if (module_path.rfind(“ptx”) != string::npos)
{
// in this branch we use compilation with parameters
const unsigned int jitNumOptions = 3;
CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
void *jitOptVals = new void[jitNumOptions];
// set up size of compilation log buffer
jitOptions[0] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
int jitLogBufferSize = 1024;
jitOptVals[0] = (void *)(size_t)jitLogBufferSize;
// set up pointer to the compilation log buffer
jitOptions[1] = CU_JIT_INFO_LOG_BUFFER;
char *jitLogBuffer = new char[jitLogBufferSize];
jitOptVals[1] = jitLogBuffer;
// set up pointer to set the Maximum # of registers for a particular kernel
jitOptions[2] = CU_JIT_MAX_REGISTERS;
int jitRegCount = 32;
jitOptVals[2] = (void *)(size_t)jitRegCount;
status = cuModuleLoadDataEx(&cuModule, ptx_source.c_str(), jitNumOptions, jitOptions, (void **)jitOptVals);
printf("> PTX JIT log:\n%s\n", jitLogBuffer);
}
本来sdk是浮点型数组相乘，我将内核函数的参数变为双精度，主机端代码分配也改为分配双精度大小的地址，运行以后结果检查是正确的，但是仔细查看上面的信息.exe其中有一行显示ptxas application ptx input,line 120;
Double is not supported.Demoting to float.(对应红色代码），这个代码怎么改才能不让它有waring
楼主您需要1.3/2.x/3.x的计算能力才能用double的。
请您检查您的ptx是否使用nvcc -ptx -arch sm_30之类的参数从.cu生成。
或者如果ptx是您手写的，请您检查开头是否有.target sm_30之类的这样
（这里用假设您的计算能力是3.0的，这是目前最常见的情况）
感谢您的来访。
对了，如果您是手写的，请您同时也别忘记去掉map_f64_to_f32的选项。
谢谢。
（如果是nvcc -ptx生成的，请无视此条。直接修改-arch sm_xx即可）
我在vs上通过.cu的属性在CUDAC/C++的device中的code generation改为compute_30,sm_30
编译运行就没有waring了。
您说的nvcc -ptx -arch 参数在哪查看，是在命令窗口下查看吗？我一般不怎么用命令窗口，如果是在命令窗口，请告诉我具体操作步骤，麻烦玫瑰版主告知。(我的ptx是CUDA C语言编写的程序生成的）
其次虽然release配置下编译的Waring没了，但性能上还是和debug配置下性能一样，没有提升，横扫版主告诉我“在你分开编译kernel的过程中就已经固定了”，但我单独在release配置下编译kernel怎么还没提升？
上文的单独构建是指您将编译产生一个单独的ptx或者cubin文件，然后手工载入。而绝非你单独一个人在release配置下进行编译。
您当前的项目配置可能在构建的时候自动将您的kernel进行了构建（从.cu到.ptx或者.cubin)。
您这就不是命令行了，您需要自己看下您的.cu属性里的构建参数有无-G（一般debug配置下有的），和-O2之类的配置参数（一般release下有的）。一般-G的代码质量非常差，所以很慢。
如果您的debug配置也是-O2（例如您使用的某个例子的配置已经配好了），则您无需操心。
以及上，起作用的是是否有-G是否有-O2，则不是配置名称叫“debug”还是“release”。这个无关的。
您应该看下配置的参数，这些参数将在编译的时候应用到nvcc的命令行上（自动）。
感谢来访。
在.cu的属性里CUDAC/C++的device下generate GPU debug information有两个选择否和-G,没有-O2，
我自己输入-O2，会提示：error MSB4030: “-O2”对于“CudaCompile”任务的“GPUDebugInfo”参数是无效值。“GPUDebugInfo”参数的类型为“System.Boolean”。求玫瑰版主能为我解惑-O2在VS .cu中哪里设置
那个是是否启用调试信息，是否-G才是影响性能的关键。
至于-O2, 有没有无所谓其实，默认会启用的（-O2或者-O3）（当没有-G的情况下，有-G会禁用优化）。
您只需要关心是否有-G即可。
也就是说，
如果您的debug和release配置下的-G都是一样配置的（都有或者都没有），
那么您得不到性能提升是正常的。
果然如此，generate GPU debug information选择“否”的情况（约80ms)下果然要比-g（约170ms）的性能好，可是速度依然比runtime api 在release配置下（约70ms），也就是generate GPU debug information选择“否”的情况下慢10ms。那这个问题出错在哪呢，劳烦玫瑰版主解惑。
您还有更多选择可以调节的，
请您查看use_fast_math选项是否为真。
如果启用该选项，将降低您的计算精度，但可以进一步的提高您的代码速度。
请您查看这个。
没什么提升（是不是use_fast_math是主机端的选项，我测得是内核函数的时间，所以没什么提升）。。。试了好多属性配置都没什么提升。。。求玫瑰版主救我于苦海！！！
use_fast_math是如假包换的device端选项，请勿怀疑。
那是我理解错了，但是use_fast_math不是在CUDAC/C++里Host里的选项，这里的HOST是什么意思呢？
LZ您好：
use_fast_math是NVCC的一个参数，您可以在NVCC的手册中看到。
以及您可以在CUDA C Programming Guide中的Mathematical Functions章节看到相关的介绍（使用之后一些数学函数的精度变化等）。
至于为什么会放置在VS里面的CUDA设置的host选项下，这个我没有研究过，或许NV认为这个参数是和host端有关的或者是由host端设置并影响device端全局的。
大致如此，祝您好运~
那不可能。完全一样的代码，一样的编译参数，无论用driver api还是runtime api将得到完全一样的运行时间。
除非您前面的2个因素之一有区别。请您检查这点。
（例如上文说的编译参数，以及计算能力等，都是有关参数，如果完全一样，不可能时间不同）
我想起横扫版主和ICE版主曾经对我说内核函数初次运行会花部分时间初始化，我的驱动程序没有避开第一次，更没有多次循环求平均值，我也懒得改程序，于是我想那就用profiler测测，他是自动进行多次测时间并求平均值，于是我就用profiler测driver api程序，但是没有timeline，我就查看console选项卡下的显示。
以前我认为console下的信息应该是和.exe显示的信息一模一样，但是令我诧异的是有个信息不一样，那就是我打印出的时间，在.exe中显示是约86ms，但是在profiler的console下显示的是约76ms（和理论相当），我虽然和高兴，但是有一点不解，console下的信息难道不应该是按照我程序执行吗，应该和.exe时间一样并且内核函数执行一次，难道console这里的测得时间已经在我不知情的情况下幕后经过多次测量求平衡值了吗？
没看懂你在说什么和你想说什么。
以及，你居然没用profiler测试时间，就简单的说我的时间是多少多少，真心很自信。我拒绝对此进行任何评价。
以及，driver api也一样可以用profiler的，请您在您的程序结束前加上：
cuCtxSynchronize();
cuCtxDestroy(你的Context);
感谢一大早来访。
好吧，我错了，加过以后，为什么profiler没有timeline呢？