我在Visual Profiler下运行完一个程序，然后选择显卡GeForce GTX 670,在Settings的Metrics/Events里勾选任何参数，在Run中选Collect Metrics/Events命令跑
最后都会告诉我：
Unable to collect metric and values
The order of kernel execution does not match the timeline. To associate events and metrics with the correct kernel, the application must behave identically on each run. Discarding all collected events and metrics .
请问下大侠们，这是咋回事啊，我的程序现在不快，本想通过这个信息看看我程序哪里有毛病呢。。。。现在啥也看不到，呜呜~~~~(>_<)~~~~
試著在主程式最後面加上cudaThreadExit();看看
加了cudaThreadExit();还是不行啊，跑完还是啥都收集不到呢:'(
在程序退出之前加个cudaDeviceReset()试试
在程序的尾巴上加上cudaDeviceReset()，结果也还是啥也搜集不到，最后的提升也一样：Unable to collect metric and values…
LZ您好，
您的问题很可能是因为多次启动之间，没有维护同样的流、kernel发布的线程grid形状和kernel顺序造成的。因此，NVVP无法通过多次运行去积累相关的数据。
因此，请LZ提供一下您代码的情况，以便具体分析。
祝您好运~
我没有使用多流呢，这个kernel确实是循环调用的，我的kernel的并行size是提前定好的
dim3 grid(GRID_X,(chmax+1),1);
dim3 block(BLOCK_X,1,1);
在调用这个kernel之前我还掉过cufft的函数，不过之前不调用cufft的时候，visual profiler也一样啥参数都统计不出来呢
我截图我的程序哈
global void corrKernel_BBS( char* d_localCaTable ,VIRTUAL_CHAR * d_productTable , char* d_pData ,
GNSS_FPGA_CORRELATOR_STRUCT* d_correlator)
{
GNSS_FPGA_CORRELATOR_STRUCT* d_correlatorT = d_correlator + blockIdx.y;//blockIdx.y是通道号
INDEX* d_Index_SharonT = (INDEX*)&(d_correlatorT->Index);
if(d_correlatorT->state!=CHN_ON)//此通道不可用
return;
if(d_Index_SharonT->IsFinish )//此通道数据处理完了
return;
long* d_pChar;
__int64 temp_d_tau ;
long temp_d_phaseLocal;
__int64 TEMP_d_phaseLocal;//载波相位一次加N个点的增量超出32位的表示范围，用一个64位临时变量缓存处理
int manThreadIndex;
int lastThreadNum ;
int i; int j; int k;
int correspond_index;
long d_sEI_temp=0; long d_sER_temp=0; long d_sPI_temp=0; long d_sPR_temp=0 ;
int d_localCaTable_index = (d_correlatorT->sv-1)*(sizeof(char)*CODE_LENGTH*CAExtTableResolution);//10230*32
__shared__ float tempS_sER[BLOCK_X]; //block内规约，BLOCK_X = blockDim.x，即1024
__shared__ float tempS_sEI[BLOCK_X];
__shared__ float tempS_sPR[BLOCK_X];
__shared__ float tempS_sPI[BLOCK_X];
manThreadIndex = (d_Index_SharonT->num_threads)/(d_Index_SharonT->samplePerThread);
lastThreadNum = d_Index_SharonT->num_threads - manThreadIndex * (d_Index_SharonT->samplePerThread);
if( blockIdx.x*blockDim.x + threadIdx.x < manThreadIndex )//一个通道内的线程Id
{
for(int mm=0;mm<(d_Index_SharonT->samplePerThread) ;mm++)
{
temp_d_tau = d_correlatorT->tau + (d_correlatorT->fCode)*((blockIdx.x*blockDim.x + threadIdx.x)*(d_Index_SharonT->samplePerThread)+mm +1);
TEMP_d_phaseLocal = d_correlatorT->phaseLocal+ (__int64)(d_correlatorT->fCarr)*((blockIdx.x*blockDim.x + threadIdx.x)*(d_Index_SharonT->samplePerThread)+mm +1);
temp_d_phaseLocal = TEMP_d_phaseLocal&0x3FFFFFFF;
// pChar = productTable[data[counter]][phaseLocal>>26][localCaTable[*pTau>>3]];
i = d_pData[d_Index_SharonT->indexStart + ((blockIdx.x*blockDim.x + threadIdx.x))*(d_Index_SharonT->samplePerThread)+mm];
j= temp_d_phaseLocal >>26;
k=d_localCaTable[d_localCaTable_index + (*((unsigned int *)(&temp_d_tau) +1) >>3)];
correspond_index = i*(CARRIER_TABLE_LENGTH*CAExtOutput*ADNumber) + j*(CAExtOutput*ADNumber)+k*ADNumber;//correspond_index = i*(16*8*4) + j*(8*4)+k*4;//
d_sER_temp += tex1Dfetch(texRef1D,correspond_index);
d_sEI_temp += tex1Dfetch(texRef1D,correspond_index+1);
d_sPR_temp += tex1Dfetch(texRef1D,correspond_index+2);
d_sPI_temp += tex1Dfetch(texRef1D,correspond_index+3);
}
}
else if( blockIdx.x*blockDim.x + threadIdx.x == manThreadIndex )
{
for(int mm=0;mm<lastThreadNum ; mm++)
{
temp_d_tau = d_correlatorT->tau + (d_correlatorT->fCode)*((blockIdx.x*blockDim.x + threadIdx.x)*(d_Index_SharonT->samplePerThread)+mm +1);
TEMP_d_phaseLocal = d_correlatorT->phaseLocal+ (__int64)(d_correlatorT->fCarr)*((blockIdx.x*blockDim.x + threadIdx.x)*(d_Index_SharonT->samplePerThread)+mm +1);
temp_d_phaseLocal = TEMP_d_phaseLocal&0x3FFFFFFF;
// pChar = productTable[data[counter]][phaseLocal>>26][localCaTable[*pTau>>3]];
i = d_pData[d_Index_SharonT->indexStart + ((blockIdx.x*blockDim.x + threadIdx.x))*(d_Index_SharonT->samplePerThread)+mm];
j= temp_d_phaseLocal >>26;
k=d_localCaTable[d_localCaTable_index + (*((unsigned int *)(&temp_d_tau) +1) >>3)];
correspond_index = i*(CARRIER_TABLE_LENGTH*CAExtOutput*ADNumber) + j*(CAExtOutput*ADNumber)+k*ADNumber;//correspond_index = i*(16*8*4) + j*(8*4)+k*4;
d_sER_temp += tex1Dfetch(texRef1D,correspond_index);
d_sEI_temp += tex1Dfetch(texRef1D,correspond_index+1);
d_sPR_temp += tex1Dfetch(texRef1D,correspond_index+2);
d_sPI_temp += tex1Dfetch(texRef1D,correspond_index+3);
}
}
tempS_sER[threadIdx.x] = d_sER_temp;//一个线程处理的采样数据点的累加值
tempS_sEI[threadIdx.x] = d_sEI_temp;
tempS_sPR[threadIdx.x] = d_sPR_temp;
tempS_sPI[threadIdx.x] = d_sPI_temp;
__syncthreads();
for(unsigned int s=blockDim.x/2; s>0 ;s>>=1)//一个block内并行规约
{
if(threadIdx.x <s)
{
tempS_sER[threadIdx.x] += tempS_sER[threadIdx.x +s];
tempS_sEI[threadIdx.x] += tempS_sEI[threadIdx.x +s];
tempS_sPR[threadIdx.x] += tempS_sPR[threadIdx.x +s];
tempS_sPI[threadIdx.x] += tempS_sPI[threadIdx.x +s];
}
__syncthreads();
}
if(threadIdx.x==0)
{
d_correlatorT->i_eNlBlock[blockIdx.x] = tempS_sER[threadIdx.x];//blockIdx.y是通道号
d_correlatorT->q_eNlBlock[blockIdx.x] = tempS_sEI[threadIdx.x] ;
d_correlatorT->i_promptBlock[blockIdx.x] = tempS_sPR[threadIdx.x] ;
d_correlatorT->q_promptBlock[blockIdx.x] = tempS_sPI[threadIdx.x];
}
__syncthreads();
}
dim3 grid(8,121);
dim3 block(1024,1,1);
我尝试在调用这个kernel一次之后程序就exit（0）；但是也统计不出信息呢
无论是ICE还是profiler, 都指出了你每次启动的kernel顺序和/或形状不同，导致无法组合数据。
profiler:
Unable to collect metric and values
The order of kernel execution does not match the timeline.
这个其实是常见的。以及此提示足够明显了。
但是你反复强调每次都是同样的次序同样的kernel启动，显然要么机器在撒谎，要么你在撒谎。
你的第二个问题，请单独测试一个kernel的时候, <<<>>>完毕后使用同步+cudaDeviceReset()一次。谢谢合作。
以及，请楼主给出自己每次都是用同样的配置和同样的顺序启动kernel的证据（例如您启动kernel的代码），而不是反复用嘴说，“我是一样的”。谢谢。
版主大人啊，你觉得我会撒谎吗，在闲也不这样折腾自己啊，为这个问题很头疼，渴望能解决，我是CPU是单线程循环调的，我测kernel执行一次的时候是这样测的:
…cpu balabala…
cudaMemcpyAsync(d_correlator, pCorr,sizeof(GNSS_FPGA_CORRELATOR_STRUCT)(chmax+1),cudaMemcpyHostToDevice);
corrKernel_BBS<<<grid,block>>>(d_localCaTable ,d_productTable , d_pData ,d_correlator);
cudaMemcpyAsync(pCorr,d_correlator,sizeof(GNSS_FPGA_CORRELATOR_STRUCT)(chmax+1),cudaMemcpyDeviceToHost);
cudaThreadSynchronize();
cudaDeviceReset();
exit(0);
然后还是同样的提示,搜集不到信息,难道是我在这个kernel前面调用了cufft的原因吗？谢谢版主大人！！
你没有给出你的原始代码，依然只是说“每次循环”如此。
为了能尽量为您服务，建议您给出没有修改过的原始代码。
以及，问几个问题：
（1）请问每次循环您都使用cudaDeviceReset()么？
（2）请问每次循环您的grid和block这2个变量都是什么？
我还是建议您直接给出原始代码，这样您就不用费心回答这些问题了。请考虑。
以及，也不排除是cufft的问题。因为无法知道每次它(cufft)是否以同样的形状和次序启动它的kernel们。
如果这个假设成立，我建议楼主可以考虑单独抽出此kernel, 然后单独做profiling. 单独抽出，然后测试，也是常见做法。
非常感谢版主大人的帮助啊，经您提示，我把cufft撤掉，然后我自己写的kernel就可以搜集到信息了，确实是cufft用的grid和block siaze跟我的不一样导致的。。。。。。。。。。那个我之前没用cufft也测不出是因为我还有其他的size的kernel。。。。
虽然我的问题解决了，但是版主热心的问了我还是回答哈，
（1）不是每次循环都使用cudaDeviceReset()
（2）每次循环的grid是（8，12，1）block是（1024，1，1），不变的
捣乱的是cufft，找到原因O(∩_∩)O~谢谢各版主热心帮助:loveliness:
嗯嗯。这个不是cufft和你的形状不同，而是它未能在多次启动间使用同样的形状（以及kernel调用顺序）。考虑到这个是个闭源的库，无法知道内部实现，也就这样了吧。
感谢您的来访。