Hi,
我安装了VS2012和CUDA-5.5,内带了很多工程,我的GPU 是Quadro 2000。我看里面Code Generation 里有compute_10,sm_10; compute_20,sm_20;compute_30,sm_30;compute_35,sm_35, 这是准备产生这么多种类型的PTX文件,但真心不知道执行时是选择哪个?
谢谢。
LZ您好:
您保留原样即可,无需手工修改,执行的时候会选择最合适的一个版本执行的。
如果您打算小小纠结一下,请参考如下的讨论帖:
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=7685&extra=page%3D1
祝您编码顺利~
赞斑竹了。
还有一个问题:Code Generation 里还是compute_10,sm_10; compute_20,sm_20; compute_30, sm_30; compute_35,sm_35时,编译出错: “uses too much shared data (0x401c bytes, 0x4000 max)”.
所以,我就去掉了“compute_10,sm_10;”, 因为我的gpu是Quadro 2000。编译是通过了,但执行时 cudaEventRecord(start,0); cudaEventRecord(stop,0); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsedTime, start,stop);出来的时间是负数。 假如加上“compute_10,sm_10;”, 执行时间就正常。请教一下斑竹。谢谢。
LZ您好:
编译的时候编译成哪些个版本的代码和您的实际硬件是无关的,以及多个版本之间是不影响的。
如果您用到了高计算能力版本的硬件特性,那么是无法编译为低版本的执行代码的,会报错。
Quadro 2000应该是SM2.1的fermi核心,照说您去掉“compute_10,sm_10”应该完全不影响的,尚不清楚是怎么回事,不过您可以试下“compute_20,sm_21”,看看。
未及之处,请其他网友/斑竹/总版主/NV原厂支持予以补充。
祝您好运~
:)谢谢回复。期待大家能发表些看法。哈哈。
关于cudaEventElapsedTime得到负数的问题,请您查看您之前的所有操作都没有发生错误。
如果都是成功执行,还得到负数,则此问题我无法解释。
哈,我过来想推荐一下前两天我提问的那个帖子的,结果版主已经链接过去了~:)
我定位出出错的地方了:cudaEventSynchronize(stop). 不知道为啥没有compute_10,sm_10;就不能正常执行?
楼主您好:
您这可能是kernel在非compute_10下编译会运行出错,因为
cudaEventSynchronize将返回之前存在的异步错误的,请您发下具体错误是什么?
只有有了具体错误,才能说为何不能正确运行。
(以及,建议发点代码)
感谢深夜来访。
把compute_10加上,cudaEventSynchronize等一整个测时间都对的。代码如下:
void runBandwidthTest(int argc, char **argv)
{
findCudaDevice(argc, (const char **)argv);
// Allocate host memory for the signal
complex *h_signal = (complex *)malloc(sizeof(complex) * SIGNAL_SIZE);
// Initalize the memory for the signal
for (unsigned int i = 0; i < SIGNAL_SIZE; ++i)
{
h_signal[i].r = (float) (2i);
h_signal[i].i = (float) (2i+1);
}
int new_size = SIGNAL_SIZE;
int mem_size = sizeof(complex) * new_size;
// Allocate device memory for signal
complex *d_signal;
checkCudaErrors(cudaMalloc((void **)&d_signal, mem_size));
// Copy host memory to device
checkCudaErrors(cudaMemcpy(d_signal, h_signal, mem_size, cudaMemcpyHostToDevice));
int blocks = SIGNAL_SIZE/1024;
cudaEvent_t start, stop;
float elapsedTime;
checkCudaErrors( cudaEventCreate(&start) );
checkCudaErrors( cudaEventCreate(&stop) );
printf("first glance\n");
//////////////////ddr_stride_16_bandwidth
checkCudaErrors( cudaEventRecord(start,0) );
system("pause");
ddr_stride_16_bandwidth <<<blocks, 256>>> (d_signal, new_size, 20);
checkCudaErrors( cudaEventRecord(stop,0) );
printf(“before synchronize pause \n”);
system(“pause”);
printf(“before synchronize \n”);
checkCudaErrors( cudaEventSynchronize(stop) );
printf(“after synchronize \n”);
system(“pause”);
checkCudaErrors( cudaEventElapsedTime(&elapsedTime, start,stop) );
printf(“ddr_stride_16_bandwidth Elapsed time : %f ms\n” ,elapsedTime);
system(“pause”);
}
Then, the elapsedTime 是负数。 同样的代码在linux下没有compute_10没有出错。
错误代码呢?
请配合点。
以及,请发下kernel的代码。
不好意思。kernel 代码如下:
global void ddr_stride_16_bandwidth( complex *A, int numElements, int log_n) //tid = 256 threads for a block.
{
int burst = 4;
int burst_mask = (2<<burst)-1;
int bid = blockIdx.x;
int tid = threadIdx.x;
int p_bid = bid << burst;
int t_3_0 = tid & burst_mask;
int t_7_4 = (tid >> burst);
int p_mem = p_bid + (t_7_4 <<(log_n-7)) + t_3_0; // consecutive
int inc_256 = 1 << (log_n-3);
float *p_A = (float *)A;
shared float s_2k_mem[4096];
s_2k_mem[tid ] = p_A[p_mem ];
s_2k_mem[tid+256] = p_A[p_mem+inc_256 ];
s_2k_mem[tid+512] = p_A[p_mem+inc_256*2];
s_2k_mem[tid+768] = p_A[p_mem+inc_256*3];
s_2k_mem[tid+1024] = p_A[p_mem+inc_256*4];
s_2k_mem[tid+1280] = p_A[p_mem+inc_256*5];
s_2k_mem[tid+1536] = p_A[p_mem+inc_256*6];
s_2k_mem[tid+1792] = p_A[p_mem+inc_256*7];
s_2k_mem[tid+2048] = p_A[p_mem+inc_256*8];
s_2k_mem[tid+2304] = p_A[p_mem+inc_256*9];
s_2k_mem[tid+2560] = p_A[p_mem+inc_256*10];
s_2k_mem[tid+2816] = p_A[p_mem+inc_256*11];
s_2k_mem[tid+3072] = p_A[p_mem+inc_256*12];
s_2k_mem[tid+3328] = p_A[p_mem+inc_256*13];
s_2k_mem[tid+3584] = p_A[p_mem+inc_256*14];
s_2k_mem[tid+3840] = p_A[p_mem+inc_256*15];
__syncthreads();
p_A[p_mem ] = s_2k_mem[tid ];
p_A[p_mem+inc_256 ] = s_2k_mem[tid+256];
p_A[p_mem+inc_256*2] = s_2k_mem[tid+512];
p_A[p_mem+inc_256*3] = s_2k_mem[tid+768];
p_A[p_mem+inc_256*4] = s_2k_mem[tid+1024];
p_A[p_mem+inc_256*5] = s_2k_mem[tid+1280];
p_A[p_mem+inc_256*6] = s_2k_mem[tid+1536];
p_A[p_mem+inc_256*7] = s_2k_mem[tid+1792];
p_A[p_mem+inc_256*8] = s_2k_mem[tid+2048];
p_A[p_mem+inc_256*9] = s_2k_mem[tid+2304];
p_A[p_mem+inc_256*10] = s_2k_mem[tid+2560];
p_A[p_mem+inc_256*11] = s_2k_mem[tid+2816];
p_A[p_mem+inc_256*12] = s_2k_mem[tid+3072];
p_A[p_mem+inc_256*13] = s_2k_mem[tid+3328];
p_A[p_mem+inc_256*14] = s_2k_mem[tid+3584];
p_A[p_mem+inc_256*15] = s_2k_mem[tid+3840];
__syncthreads();
}
我找到错误了,是我的kernel出错了,只要调用我的kernel, 测时间就不准了。但还不知道具体错误,谢谢大家的提醒。