调用一次内核函数执行的时间限制

编了一个程序验证cuda的加速效果,在num_threads=32时,执行结果和cpu的差不多,但是时间是cpu的二倍,如果改为64,花3000ms,而且在执行下面红色部分会报错,按照以前问版主的同样问题的解释,说是内核崩溃,但是如果是把num_threads改为512,直接就说执行时间过长了,我尝试把内核的四重循环拿一重出来,但是结果还是一样

#include <stdlib.h>
#include <stdio.h>
#include <string.h>

#include
#define H 5
#define Rr 5
#define R 10
#define Pi 3.1415

// includes CUDA
#include <cuda_runtime.h>

// includes, project
#include <helper_cuda.h>
#include <helper_functions.h> // helper functions for SDK examples
using std::cout;
using std::endl;
using std::cin;

////////////////////////////////////////////////////////////////////////////////
// declaration, forward
void runTest(int argc, char **argv);

extern “C”
void computeGold(float *reference, float *idata, const unsigned int len);

////////////////////////////////////////////////////////////////////////////////
//! Simple test kernel for device functionality
//! @param g_idata input data in global memory
//! @param g_odata output data in global memory
////////////////////////////////////////////////////////////////////////////////
device float Find1(int i,int j,int k ,int n)
{
float num=blockDim.x;
float x1=R/numicos(2Pi/numj);
float y1=R/numisin(2Pi/numj);
float x2=Rr/numkcos(2Pi/numn);
float y2=Rr/numksin(2Pi/numn);
float S=(x1-x2)(x1-x2)+(y1-y2)(y1-y2)+HH;
return(H
H/(PiSS));
}

global void
testKernel(int i,float * save)
{

// shared memory
// the size is determined by the host application
extern shared float sdata1;
float s1=0,s2=0;

// access thread id
const int tid = threadIdx.x;
if(i==1&&tid==0)
*save=0;
// access number of threads in this block
const float num_threads = blockDim.x;

for(int j=1;j<=num_threads;j++)//R圆弧度对应的数 j
{
for(int k=1;k<=num_threads;k++)//r圆半径对应的数 k
{
sdata1[tid]=Find1(i,j,k,tid+1)(2Pi/num_threads);
__syncthreads();
for(int n=num_threads/2;n>=1;n/=2)
{
if(tid<n)
{

sdata1[tid]+=sdata1[tid+n];
}
__syncthreads();
}

if(tid==0)
{
s1+=sdata1[0](Rrk/num_threads)*(Rr/num_threads);
}
}

if(tid==0)
{
s2+=s1*(2*Pi/num_threads);
s1=0;
}
}

if(tid==0)
{
save+=s2(R/num_threadsi)(R/num_threads);
s2=0;
}

// write data to global memory

}

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main(int argc, char **argv)
{
runTest(argc, argv);
}

////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for CUDA
////////////////////////////////////////////////////////////////////////////////
void
runTest(int argc, char **argv)
{
bool bTestResult = true;

printf(“%s Starting…\n\n”, argv[0]);

// use command-line specified CUDA device, otherwise use device with highest Gflops/s
int devID = findCudaDevice(argc, (const char **)argv);

unsigned int num_threads ;
cout<<“请输入线程数”<<endl;
cin>>num_threads;
unsigned int mem_size = sizeof(float) * num_threads;

// allocate host memory
//float *h_idata = (float *) malloc(mem_size);

// initalize the memory
/* for (unsigned int i = 0; i < num_threads; ++i)
{
h_idata[i] = (float) i;
}*/

// allocate device memory
//float *d_idata;
//checkCudaErrors(cudaMalloc((void **) &d_idata,sizeof( float)));
// copy host memory to device
// checkCudaErrors(cudaMemcpy(d_idata, h_idata, mem_size,
// cudaMemcpyHostToDevice));

// allocate device memory for result

float * save;
checkCudaErrors(cudaMalloc((void **) &save, sizeof( float)));
StopWatchInterface *timer = 0;
sdkCreateTimer(&timer);
sdkStartTimer(&timer);

// setup execution parameters
dim3 grid(1, 1, 1);
dim3 threads(num_threads, 1, 1);

// execute the kernel
for(int i=1;i<=num_threads;i++)
{

testKernel<<< grid, threads, mem_size >>>( i,save);
}

// check if kernel execution generated and error
getLastCudaError(“Kernel execution failed”);

// allocate mem for the result on host side
float *h_odata = (float *) malloc(sizeof( float));
// copy result from device to host
checkCudaErrors(cudaMemcpy(h_odata, save, sizeof(float) ,cudaMemcpyDeviceToHost));

sdkStopTimer(&timer);
printf(“Processing time: %f (ms)\n”, sdkGetTimerValue(&timer));
sdkDeleteTimer(&timer);
float *h_out = (float *) malloc(sizeof( float));
h_out=h_odata/(PiRR);
std::cout<<*h_out<<endl;

// compute reference solution
// float *reference = (float *) malloc(mem_size);
// computeGold(reference, h_idata, num_threads);

// check result
/if (checkCmdLineFlag(argc, (const char **) argv, “regression”))
{
// write file for regression test
sdkWriteFile(“./data/regression.dat”, h_odata, num_threads, 0.0f, false);
}
else
{
// custom output handling when no regression test running
// in this case check if the result is equivalent to the expected soluion
bTestResult = compareData(reference, h_odata, num_threads, 0.0f, 0.0f);
}
/

// cleanup memory
free(h_odata);
free(h_out);
checkCudaErrors(cudaFree(save));

cudaDeviceReset();
exit(bTestResult ? EXIT_SUCCESS : EXIT_FAILURE);
}

因为num_threads=32时,结果正确,所以应该不会有什么逻辑错误,请问对一次内核函数的执行,有没有时间限制,要怎么修改时间限制?
[/i]

楼主您好,首先一般情况下(WDDM驱动)是会有运行时间限制的。

您可以如下修改时间限制:
(1)安装nsight visual edition,
(2)从开始菜单找到nsight monitor, 并运行, 并等待3秒钟。
(3)在右下角的任务栏通知区域找到nsight monitor图标. 右键点击,选择options
(4)在弹出的窗口中窗口左侧选择general, 右侧选择TDR Delay选项. 改成60, 然后点击确定。
(5)重新启动计算机。
(6)重新运行.

注意!请一定要重新启动计算机!谢谢配合!

LZ您好,横扫斑竹已经向您详细解释了如何修改TDR。

此外,如果您使用的是telsa卡,那么切换为TCC驱动模式就不再有kernel运行时间的限制了,也无需nsight修改。如果您用的不是telsa卡或者虽然使用telsa卡但是不便修改驱动模式,那么请无视此建议。


此外,
1:您顶楼中说的“红色部分”文中没有体现,可能是编辑错误。
2:斑竹对您所有的回答,请以原帖,原情况,原上下文,原回答为准,请勿让斑竹无辜躺枪。
3:一般来说,将原有串行算法改为GPU实现,需要重新改写为合适GPU的算法,铺上大量线程计算。如果只是将串行算法稍加改动直接扔进kernel里面,上少量线程的话,可能效率难以提升。

祝您编码愉快~

嗯,谢谢

我拿了两重循环出来,可以运行,Num_threads=128了

LZ您好,上面建议您修改TDR,您是否修改,效果如何?

以及,针对GPU特性重新修改算法实现是推荐的做法,但是如果您只有128个线程的话,一般来说,效率还是不行的。

祝您编码顺利~

正在和导师汇报,希望能导师能买telsa卡

LZ,只需修改TDR,即可立即解决您的kernel运行超时问题,您不想先试试么?

祝您编码愉快!

学习CUDA,主要还是想做比较复杂的科学计算,用消费市场的geforce卡,满足不了需求,顺便问一下,如果购买Tesla的卡,是需要安装在服务器上,还是装在一般的台式机就行了?

学习CUDA,主要还是想做比较复杂的科学计算,用消费市场的台式机,满足不了需求,顺便答一下,如果购买Tesla的卡,是需要安装在服务器的,我不能装在一般的台式机就行了。

版主幽默啊,在下什么都不懂,惭愧啊

对您建议过,让您修改TDR delay以便在WDDM下也能运行大kernel, 但是被您拒绝。您非得认为必须上tesla才可以。所以我心里很伤,但是不能责备你,就按照您的逻辑给予了评价。您看是不是重新考虑下?

如果修改了TDR delay,会烧显卡吗?

我建议您还是不要上本论坛寻求建议为好。

询问了,却又反复质疑"geforce不行", 以及"即使行,修改了也能烧毁"。那您还是直接考虑别的吧。

LZ您好,修改TDR不会烧显卡,不会使显卡生锈,不会使显卡断裂,更不会引起机器爆炸等不良反应。即使您关掉TDR,也不过是kernel死循环不退出的时候需要重启机器而已。

geforce卡可以插在台式机上,geforce卡也可以插在工作站上;
telsa卡可以插在台式机上,telsa卡也可以插在工作站上。

他们都能进行通用数值计算。
如果您需要telsa的特性,请选择购买。
如果您要求特别高或者有现成的工作站,可以把telsa插在工作站上使用。

明白了,谢谢版主

我在科学计算的时候也遇到了这样的问题。
需要计算时间很长。
想请问下用nsight把TDR Delay改的很长会不会有问题。
单显卡在长时间计算时候,桌面显示会不会有问题?

您好,请参考15#的说法。修改TDR不会对显卡本身造成任何伤害。

另外,如果您的显卡既要计算又要显示,那么在计算的时候,可能产生卡顿等现象,这是正常的。