CUDA如何选择kernel二进制文件？

system · 2013 年11 月 12 日 09:05

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文件，但真心不知道执行时是选择哪个？
谢谢。

system · 2013 年11 月 12 日 09:08

LZ您好：

您保留原样即可，无需手工修改，执行的时候会选择最合适的一个版本执行的。

如果您打算小小纠结一下，请参考如下的讨论帖：

http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=7685&extra=page%3D1

祝您编码顺利~

system · 2013 年11 月 12 日 09:25

赞斑竹了。
还有一个问题：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;”，执行时间就正常。请教一下斑竹。谢谢。

system · 2013 年11 月 12 日 09:37

LZ您好：

编译的时候编译成哪些个版本的代码和您的实际硬件是无关的，以及多个版本之间是不影响的。

如果您用到了高计算能力版本的硬件特性，那么是无法编译为低版本的执行代码的，会报错。

Quadro 2000应该是SM2.1的fermi核心，照说您去掉“compute_10,sm_10”应该完全不影响的，尚不清楚是怎么回事，不过您可以试下“compute_20,sm_21”，看看。

未及之处，请其他网友/斑竹/总版主/NV原厂支持予以补充。

祝您好运~

system · 2013 年11 月 12 日 10:38

:)谢谢回复。期待大家能发表些看法。哈哈。

system · 2013 年11 月 12 日 12:05

关于cudaEventElapsedTime得到负数的问题，请您查看您之前的所有操作都没有发生错误。

如果都是成功执行，还得到负数，则此问题我无法解释。

system · 2013 年11 月 12 日 12:16

哈，我过来想推荐一下前两天我提问的那个帖子的，结果版主已经链接过去了～:)

system · 2013 年11 月 12 日 13:50

我定位出出错的地方了：cudaEventSynchronize(stop). 不知道为啥没有compute_10,sm_10;就不能正常执行？

system · 2013 年11 月 12 日 13:55

楼主您好：

您这可能是kernel在非compute_10下编译会运行出错，因为
cudaEventSynchronize将返回之前存在的异步错误的，请您发下具体错误是什么？

只有有了具体错误，才能说为何不能正确运行。
（以及，建议发点代码）

感谢深夜来访。

system · 2013 年11 月 12 日 14:04

把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没有出错。

system · 2013 年11 月 12 日 14:08

错误代码呢？

请配合点。
以及，请发下kernel的代码。

system · 2013 年11 月 13 日 01:26

不好意思。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();

}

system · 2013 年11 月 13 日 06:28

我找到错误了，是我的kernel出错了，只要调用我的kernel，测时间就不准了。但还不知道具体错误，谢谢大家的提醒。