有时候我在分配数据到寄存器时，希望通过一些专业软件查看这些数据是否分配到我希望的寄存器或者事与愿违分配到局部存储器从而导致性能下降，请问visual profiler和nsight monitor哪个软件能提供这样的功能，并且如何操作？
你无需操心这个：
（1）如果有可能，编译器总是安排到寄存器的。
（2）如果真的无法安排到寄存器，你也不能强迫它安排到寄存器。
（3）请ICE补充。
感谢来访。
LZ您好：
我来补充一下2#玫瑰斑竹。
如果您没有在kernel里面使用非常多的变量，以至于必须同时使用的变量超过了一个线程所能使用的寄存器数量，以及您没有在kernel里面定义下标访问规律在编译时无法得知的数组，那么一般都会将变量放置在寄存器中的，以及如玫瑰斑竹所说，真的放不进去，你也没其他办法，只能修改您的实现。
您在nvvp的profile结果中可以看到local memory传输的统计，如果该项统计不为零，说明您的代码用到了local memory。
大致如此，祝您好运~
“如果您没有在kernel里面使用非常多的变量，以至于必须同时使用的变量超过了一个线程所能使用的寄存器数量，以及您没有在kernel里面定义下标访问规律在编译时无法得知的数组”还是不太懂你的意思。
1.您第一句是不是应该这样写道““如果您（没有）在kernel里面使用非常多的变量，以至于必须同时使用的变量超过了一个线程所能使用的寄存器数量”，多写了两个字“没有”。
2.其次如何定义下标访问规律才能使得数组分配到寄存器。
3.local memory传输统计在visual profiler哪里能看到？我只看到了timeline。
LZ您好：
我来把第一句改成短句重新说一下：
1：如果您的kernel里面使用了非常多的变量，经过编译器尽可能安排以后，同时需要使用的变量依然超过一个线程所能使用的最大寄存器的数量，那么会有部分变量被放置到local memory。
2：如果您在kernel里面定义了数组，并且这个数组里面的元素的下标访问规律在编译的时候无法得知（即在运行的时候才能知道），那么该数组将被放置到local memory（反之会根据已知的下标访问规律安排使用编译器）。
如果不是上述的1：，2：两种情况，那么一般编译器会安排使用寄存器。
关于您的问题2：
这个只需要在编译的时候就知道各个元素的具体行为。
比如
a[0]=…;
a[1]=…;
这种在编译的时候可以直接知道下标的访问规律，那么虽然逻辑上是一个数组，实际编译器会安排多个寄存器使用，和直接使用多个变量是一样的。
而如果a=…;且x是运行时才知道的数（比如传入参数或者依赖于运行时得到的数据才能确定的数），那么在编译时编译器无法确定这究竟是哪个元素的行为，同时N卡的寄存器不支持寻址，所以只能放置在lobal memory中。
关于您的第三个问题：
您可以选择Analyze All（在界面的左下角），分析完成以后，看下面的Details选项卡即可看到local memory相关的项。
或者在timeline中选定一个kernel，然后在右侧的Properties中看。
如果不为0%，那说明使用了local memory。
祝您好运~
您好，您的意思是说对数组元素分别赋值，编译器就知道了数组元素的访问规律，那么为什么不能有for循环赋值呢，如果我申请了100个浮点型数组，分别赋值太麻烦了。
分别有两个关于local memory的选项，一个是local load(store) transactions,另一个是local load（store),分别指的是什么？
LZ您好：
我没说不能用for的，上面那个只是一个简单举例而已。
以及，在kernel内部包括初始化赋值在内的所有的对数组元素的访问，如果在编译的时候就知道具体是那个元素执行哪些确定的操作，那么会安排使用寄存器。
这都是编译器判断之后决定的。
float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};但是这样赋值为什么不行呢？
float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};为什么这样赋值不行呢？分配到了局部存储器
请给出您的完整代码以便分析。
global void sgemmNN(const float* A,int lda,const float* B,int ldb,float* C,int ldc,int k)
{
A+=blockIdx.x64+threadIdx.x+threadIdx.y16;
B+=threadIdx.x+__mul24((__mul24(blockIdx.y,16)+threadIdx.y),ldb);
C+=blockIdx.x64+threadIdx.x+__mul24((threadIdx.y+__mul24(blockIdx.y,ldc)),16);
int num=0;
int id=blockIdx.y64+threadIdx.y*16+threadIdx.x;
__shared__ float bs[16][17];
float c[16]={0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0};
const float* Blast=B+k;
do
{
//num++;
#pragma unroll
for(int i=0;i<16;i+=4)
bs[threadIdx.x][threadIdx.y+i]=B[i*ldb];
B+=16;
__syncthreads();
#pragma unroll
for(int i=0;i<16;i++,A+=lda)
{
c[0]+=A[0]*bs[i][0];c[1]+=A[0]*bs[i][1];
c[2]+=A[0]*bs[i][2];c[3]+=A[0]*bs[i][3];
c[4]+=A[0]*bs[i][4];c[5]+=A[0]*bs[i][5];
c[6]+=A[0]*bs[i][6];c[7]+=A[0]*bs[i][7];
c[8]+=A[0]*bs[i][8];c[9]+=A[0]*bs[i][9];
c[10]+=A[0]*bs[i][10];c[11]+=A[0]*bs[i][11];
c[12]+=A[0]*bs[i][12];c[13]+=A[0]*bs[i][13];
c[14]+=A[0]*bs[i][14];c[15]+=A[0]*bs[i][15];
}
__syncthreads();
}while(B<Blast);
//pass[id]=num;
#pragma unroll
for(int i=0;i<16;i++)
{
(C+ildc)=c[i];
}
}
这种赋值我通过visual profiler 测出local memory有吞吐量
您可以参考profiler手册中的Matrics Reference章节的解释。
[attach]3448[/attach]
经过我试验，数组都是分配到局部存储器（for循环也测出局部存储器吞吐量），当分别申请16个变量时，并以此赋值，局部存储器吞吐量就为0，我的问题是，如果数组很大，我申请很多变量就相当不方便，怎么解决这个问题呢，版主大人？
我来说两句，首先说，
楼主的“数组都将分配在local memory中"的说法是错误的。前文ICE版主也多次说明了。
这个很容易给出反例的，例如这种：
global void wch(int *p)
{
int sum[4] = {0, 0, 0, 0}; //int a = 0, b = 0, c = 0, d= 0;
for (int i = 0; i < 10; i++)
{
sum[0] += p[i * 4 + 0]; //a += …
sum[1] += p[i * 4 + 1]; //b += …
sum[2] += p[i * 4 + 2]; //c += …
sum[3] += p[i * 4 + 3]; //d += …
}
p[0] = sum[0] + sum[1] + sum[2] + sum[3]; //p[0] = a + b + c + d;
}
这个例子在-O2下无local traffic
其次，如果你分配100个float的数组，那是必然放不下的。这个请绝望（除非你在用3.5的卡，还是有可能的）。
最后，无法针对你的特定例子给出评价（因为此例子的代码一些下标被论坛吃掉了，请您重新看下您的发帖是否有问题），但如果在给定了足够的编译器限制的情况下，外加-O2编译，和你单独的16个变量无区别的。请重新尝试。
感谢来访。
我发的上文有个文字错误，
最后一处的“编译器”应为“寄存器”。
特此更正。
版主您好：如您上面提到的—O2编译是什么意思，我是直接进行编译的，怎么进行—O2编译