有关程序测试的问题

版主你好,我在使用NVIDIA Visual Profiler 对写完的CUDA程序做测试的时候 他的"console" 选项卡里显示一下信息

========Warning: Application returned non-zero code -1073741819
======== Warning: No CUDA application was profiled, exiting

我不太明白是什么原因,麻烦版主和大家帮忙分析。:slight_smile:

楼主您好,

您可能使用了类似CUDA_CHECKERROR之类的宏,导致您的代码在cuda runtime自我初始化之前就退出了。
(因为您的该宏的用法导致您提前的退出的)。而因为无runtime建立的context, 此时profiler就会报告这是非cuda程序的。

建议的解决方案:
因为贵宏将错误代码转换为了返回值,而-1073741819这个是众所周知的臭名昭著的0xc0000005, 即windows的"Access Violation"(即Linux的Segment fault), 请仔细检查代码,看看是哪行的问题。

您可以逐行跟踪代码,看看哪里会挂。

在您修正了代码无BUG后,请再尝试进行profiling.

代码如下

#include <cuda_runtime.h>
#include <stdio.h>
#include <assert.h>

#define THREADBLOCK_SIZE 1024

#define w 32

#define d_n  8 

__global__ void scan(float *c, float *a,float *e, int n)  //n是一个线程块也就是一个block处理的元素个数,是2倍的THREADBLOCK_SIZE
{  
   __shared__ float b[2*THREADBLOCK_SIZE+THREADBLOCK_SIZE/32];   
   int x = threadIdx.x;
   int y = blockIdx.x; 
   int m = 1;  
   b[2*x+x/w] = a[n*y+2*x];     // w是bank数,由显卡型号决定。可以宏定义。
   b[2*x+1+x/w] = a[n*y+2*x+1];
   for (int d = n>>1; d > 0; d >>= 2)  //要循环展开 ,考虑展开的系数,如果N是2的某次方可以是2,依次考虑,当然不整除也行
  {
   __syncthreads();
   if (x < d) 
   {  int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   b[j]+=b[k];
   }
   m*=2;
   d >>= 1;
   __syncthreads();
   
   if (x < d) 
   {  int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;   //33行
   b[j]+=b[k];
   }
   m*=2;
  }
   if (x == 0) { e[y] = b[n - 1];          ; 
   b[n - 1] = 0;            
   }  
   
   for (int d = 1; d < n; d *= 4)
  {
   __syncthreads();
   m >>= 1; 
   if (x < d) 
   {
   int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   float t = b[k]; 
   b[k]=b[j];
   b[j]+=t;
   }
   d *= 2;
   __syncthreads();
   m >>= 1; 
   if (x < d) 
   {
   int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   float t = b[k]; 
   b[k]=b[j];
   b[j]+=t;
   }
   }
   __syncthreads();
   c[n*y+2*x]= b[2*x+x/w];
   c[n*y+2*x+1]= b[2*x+1+x/w];
}

__global__ void scan1(float *c, float *a, int n)  
{  
   __shared__ float b[2*THREADBLOCK_SIZE+THREADBLOCK_SIZE/32];   
   int x = threadIdx.x;
   
   int m = 1;  
   b[2*x+x/w] = a[2*x];     // w是bank数,由显卡型号决定。可以宏定义。
   b[2*x+1+x/w] = a[2*x+1];
   for (int d = n>>1; d > 0; d >>= 2)  //要循环展开 ,考虑展开的系数,如果N是2的某次方可以是2,依次考虑,当然不整除也行
  {
   __syncthreads();
   if (x < d) 
   {  int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   b[j]+=b[k];
   }
   m*=2;
   d >>= 1;
   __syncthreads();
   if (x < d) 
   {  int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   b[j]+=b[k];
   }
   m*=2;
   }
   if (x == 0) {           
   b[n - 1] = 0;            
   }     
   for (int d = 1; d < n; d *= 4)
  {
   __syncthreads();
   m >>= 1; 
   if (x < d) 
   {
   int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   float t = b[k]; 
   b[k]=b[j];
   b[j]+=t;
   }
   d *= 2;
   __syncthreads();
   m >>= 1; 
   if (x < d) 
   {
   int k=m*(2*x+1)-1+x/w;
   int j=m*(2*x+2)-1+x/w;
   float t = b[k]; 
   b[k]=b[j];
   b[j]+=t;
   }
   }
   __syncthreads();
   c[2*x]= b[2*x+x/w];
   c[2*x+1]= b[2*x+1+x/w];
}


__global__ void add (float *c, float *ee, int n) 
{
 
 int x = threadIdx.x;
 int y = blockIdx.x;

 __syncthreads();


 __syncthreads();
 
 c[n*y+2*x] += ee[y];
 c[n*y+2*x+1] += ee[y];

}   

void constantInit(float *data, int size1, float val)
{
   for (int i = 0; i < size1; i++)
   {
   data[i] = val;
   }
}

 
 int mscan(int m_w)

{
   
   unsigned int size = sizeof(float) * m_w;  //

   
   
   float *h_A = (float *)malloc(size);     //CPU里分配数组A的空间首地址和具体空间

   
   constantInit(h_A, size, 1.0f);      //数组A初始化
   
   
   float *h_C = (float *)malloc(size);     //CPU里分配数组C的空间首地址和具体空间

   float *d_A, *d_C,*d_E,*d_EE;

   
   cudaMalloc((void **) &d_A, size);  //设备(GPU)里分配数组A的空间首地址和具体空间

   cudaMalloc((void **) &d_C, size);   // 设备(GPU)里分配数组C的空间首地址和具体空间

   
   unsigned int size1 = sizeof(float) * d_n;
   
   cudaMalloc((void **) &d_E, size1);   // 设备(GPU)里分配数组E的空间首地址和具体空间
   
   cudaMalloc((void **) &d_EE, size1);  // 设备(GPU)里分配数组EE的空间首地址和具体空间

   
   cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);  //拷贝数据去gpu

   
   scan <<< d_n, THREADBLOCK_SIZE >>>(d_C, d_A,d_E, 2*THREADBLOCK_SIZE);

   scan1 <<< 1, d_n/2 >>>(d_EE, d_E,d_n);

   add <<< d_n, THREADBLOCK_SIZE >>>(d_C, d_EE, 2*THREADBLOCK_SIZE);

   cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);   //拷贝数据回cpu
   
   free(h_A);
   free(h_C);
   cudaFree(d_A);
   cudaFree(d_C);
   cudaFree(d_E);
   cudaFree(d_EE);
   return 0;
   
}

int main()

{ 

   
   int m_w=d_n*THREADBLOCK_SIZE;

   int me=mscan(m_w);
   exit(me);
   
 }  

是不是203行返回值的问题?我看到大多数例子中都是对计算结果核查之后得出正确或是错误 作为函数的返回值,交回给主函数main () 我因为不熟悉所以没有核查结果 ,直接返回了一个0,感觉好像不太正确,请教版主指点,好久不用C语言了

楼主您好,您的代码直接看无异常返回,但是执行的结果却是0xc0000005(access violation), 请仔细step您的代码,一步步的调用,看哪一步出现的异常。

access violation / segmentation fault如果没有被捕获,可能程序将随时返回这个并结束运行。

所以,代码在您手中,您只需要简单的step运行一遍,直接就知道哪里的问题了。而不是让版主们人肉编译和人肉调试。

以及,根据人肉调试结果,您的这里将导致access violation:
函数constantInit()中的行150:
data[ i ]= val;
因为您的循环范围越界了。亲。

根据给出的建议自己调试是您的本分,请考虑这点。这次依然为您指出。请下次注意。

太感谢了 :slight_smile:

按着版主说的调试,没有理解为什么会越界:'(

明白了,汗,好久不写C程序好多不习惯,麻烦版主了多谢多谢:)

您客气了。服务您是我们的荣幸。

感谢来访,节日快乐!