–ptxas-options=-v里面很多参数看不懂，跪求高人解释。
同样看不懂。跪求高人解释。
#include <stdio.h>
#include <cuda_runtime.h>
#include <assert.h>
#define N 4
void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}
__constant__ float c_a=1.3;
__global__ void hey(float *d_a1)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
__shared__ float s_a[N];
// float r_a=1.1(此句加与不加都显示用了6个register)
s_a[tid]=1.0;
d_a1[tid] = d_a1[tid] + s_a[tid] + c_a;
}
int main()
{
int i;
float h_a[N];
float *d_a;
size_t msize=N*sizeof(float);
for (i=0; i<N; i++){
h_a[i]=i*2.0;
}
cudaMalloc((void**)&d_a, msize);
cudaMemcpy(d_a, h_a, msize, cudaMemcpyHostToDevice);
dim3 dimBlock( 1 );
dim3 dimThread( N );
hey <<< dimBlock, dimThread >>> (d_a);
checkCUDAError("kernel invocation");
cudaMemcpy (h_a, d_a, msize, cudaMemcpyDeviceToHost);
printf("h_a=%f,%f,%f,%f\n", h_a[0],h_a[1],h_a[2],h_a[3]);
cudaFree(d_a);
return 0;
}
编译时加–ptxas-options=-v，可得以下信息
ptxas info : 0 bytes gmem, 4 bytes cmem[2]
ptxas info : Compiling entry function ‘_Z3heyPf’ for ‘sm_20’
ptxas info : Function properties for _Z3heyPf
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 6 registers, 16 bytes smem, 40 bytes cmem[0]
可以看懂的是 4 bytes cmem[2]对应程序里面的__constant__ float c_a，以及16 bytes smem对应程序里面的__shared__ float s_a[4]
问题：(1)6个registers如何算得？
(2)40 bytes cmem[0]是什么？
(3)如程序中所示，如果在global函数中(22行)定义并赋初值float r_a=1.1，再次按–ptxas-options=-v编译，显示的也是used 6 registers，而且其他各种内存使用都不变，为何？是r_a未定义在register上吗？
楼主您好，无公开资料。无法评价ptxas。抱歉。
您可以等待NVIDIA原厂支持为您提供资料以及详细信息。
我用nvidia visual profiler拿来实验，在global函数中多定义一个float(就是上面代码的22行)，结果显示register并没有增加，不知为何
因为ptxas的相关资料不能被公开，所以您问的cmem[n]的问题无法直接回答。
您可以让NVIDIA原厂支持回答。
关于寄存器不增加，这个很正常，当且仅当：
（1）1个常量不能作为立即数
（2）1个常量不能放入constant cache
（3）1个常量只能放入寄存器
（4）当前已经分配的寄存器数目中，没有空闲的。
才可能会增加寄存器数量。
感谢来访。
嗯，又有了更多的理解，十分感谢！
服务您是我的荣幸。
期待您的再次来访。
那这就产生一个问题：我们在编程时无法准确估计寄存器占用的数量？
楼主您好，您在编码的时候无法直接得到这个。
但您编译的时候会得到used xxxx registers的信息的。
感谢来访。
那也只能这样了，但是寄存器的数量对代码的编写和算法产生限制，无法在编码的时候准确估计寄存器使用的话，是件挺麻烦的事情。
我建议您专心写代码即可，编译器的安排是编译器的事。
操心这个时间，不如多泡泡妞了，您觉得呢？
道理是没错，不过我是泡妞之余搞这玩意儿的
好吧。如果你真的需要一边写代码一些能计算出寄存器的使用，
你需要：
（1）统计每个变量的分配位置（例如是否真的需要在寄存器中，还是作为立即数还是作为常数等)
（2）统计到目前未知已经分配的寄存器数目，统计前面的那些寄存器变量是否已经超过有效生存期，是否可以被标记为空闲了，从而决定你是否需要分配一个新的寄存器。
（3）评估为了性能的影响，一个长期不用的寄存器是将它的值交换出去呢？让出新空间给新变量用呢？还是让它长期留着，节省交换代价，直接分配一个新寄存器。
（4）需要知道硬件的分配粒度，例如某些卡如果你知道分配13个和分配16个寄存器实际运行资源占用无区别，你可以大胆的多分配3个。还是不能。
（5）…
（6）…
正因为这些以及还有更多更多的原因，你大脑进行编译、以及后续的寄存器分配非常非常折腾。我还是建议您编译下看看编译器给出的结果。
真抱歉，我之前没有继续追问的意思，，，不过感谢你详细的回复，又可以让我学习学习。