在请问下哈，我编程示意图是这样 for（时间循环）{ 核函数1（源区）； 核函数2（散射区）}。函数2需要一个函数1的数据；
1函数测试过，完全正确。但是当1,2函数同时运行时，不仅整个循环结果是错误的，而且第一个函数1的输出也跟着变错了。当把第二个函数屏蔽掉时，只运行第一个函数1。这时1函数的结果就正确了。按道理说，第二个核函数的运行是不会影响第一个核函数1的。而且，两个函数的唯一联系就是1函数仅仅传输一个数据给2函数，这样也没理由函数2会导致函数1错误运行啊？
按照您的叙述，这是不应该发生的。
但如果kernel2和kernel1都会改写同一个数组，比如ex，那么kernel2的错误会传递给kernel1。
同时，还有一种比较少见的情况，就是如果kernel2如果写越界了，恰好改掉了kernel1的数据，（注意此时未必会报错），也可能连带造成kernel1也错误。
上述情况无论哪种，都表明kernel2一定是有问题的，请努力检查之。
祝您早日修复BUG~
非常感谢ice版主！
额，别着急感谢，先去看程序吧，还指不定问题在哪里呢~
就是啊，寻找中……
cudaMalloc( (void **) &pEin, mem_size1 );
cudaMalloc( (void ** ) &pHin, mem_size1 );
cudaMalloc( (void ** ) &pOb, mem_ob );
cudaMalloc( ( void** ) &pEx, mem_size1 );
cudaMalloc( ( void** ) &pHy, mem_size1 );
//一下四个是分配二维数组空间
cudaMalloc( ( void ** ) &pFE1, mem_size2 );
cudaMalloc( ( void ** ) &pFE2, mem_size2 );
cudaMalloc( ( void ** ) &pFH1, mem_size2 );
cudaMalloc( ( void ** ) &pFH2, mem_size2 );
错误人在延续，请ice版主帮忙看看，这样的空间划分，不同核函数里的数据会相互影响么？谢谢！
您好，如上的显存申请是没问题的，至于kernel里面是否会互相影响，要看kernel写的如何，我猜测，您需要实现的逻辑是不会影响的。
但是您的kernel有没有问题，包括有没有越界，这些就无从判断了，请您还是淡定地继续检查程序吧，目前别无他法。
祝您调试愉快~
[
for( TimeStep = 0; TimeStep <= TimeStop; TimeStep++ )
{
FDTD_SOURCE<<< 1,1000 >>> ( pEin, pHin, TimeStep, WaveLength, WL, Isource );
CUT_CHECK_ERROR( "Kernel failed" );
cudaThreadSynchronize();
cudaMemcpy( Ein, pEin, mem_size1, cudaMemcpyDeviceToHost );
cudaMemcpy( Hin, pHin, mem_size1, cudaMemcpyDeviceToHost );
FDTD_SCATTER<<< 1, 1000 >>>( 11 );
CUT_CHECK_ERROR( "Kernel failed" );
cudaThreadSynchronize();
}
第二个核函数内容如下：
__global__ void FDTD_SCATTER( int c )
{
int tid = threadIdx.x;
__shared__ float ex[1000];
ex[tid] = 11.0;
}
现在问题是这样，循环中的第一个核函数在下面两种情况下可以正确运行。1：将第二个核函数ex[1000]的__shared__去掉；第2种情况：保留shared，但把ex[tid]的赋值去掉；这两种条件，满足任意一种，第一个核函数都可以正确运行。当然，如果直接把第二个核函数给注释掉，循环中的第一个核函数FDTD_SOURCE<<< 1,1000 >>> 也也能正确运行。
我感觉问题真的很蹊跷。block之间不是相互隔离的么！而且，以上我的两个核函数，除了是在同一个循环里工作，其它根本没有一点联系了啊。为什么第二个核函数的存在，会影响第一个核函数的运行啊？而且还会存在以上两种情况？ 求解惑啊！
请恕我直言，您的理解，就目前而言还是混乱的。
1：您在近期的一直说的是block之间的隔离问题，而您的kernel只启动了一个block，所以您的说法和做法是矛盾的。其实我很早之前就建议过，用kernel分隔电场更新，磁场更新和源更新。每个kernel里面开多个block是没有关系的。
2：您在第二个kernel里面仅仅完成了对shared memory的赋值。应该说这么做是没有什么意义的。您的__shared__ ex只在第二个kernel里面用，并且考虑到shared memory是按block可见的。那么，其他kernel是无法访问这个数组的，以及您的第二个kernel如果开两个block，那么这两个block也是互相不能访问对方的shared memory的。同时，shared memory太小了，一个SM上满打满算也就几十KB的容量，根本不够用的，如果真是这个规模的计算，CPU瞬间就秒杀了。第三，您在给shared memory赋值的时候，没有加__syncthreads()。
3：在您的循环中，不用每步都把值拷回host端，直接在device端搞定即可。况且您每次似乎都回拷到了host端的同样的位置，并且也没有打印输出啥的，似乎是无用功。
4：您上述代码的第12行怎么冒出来一个TimeStep和半个括号？
5：您的第二个kernel的参数 int c，根本就没有被用到。当然，这个也不算错误。
总之，如果说有什么建议的话，建议您重新学习/复习一下CUDA的基本内容，以及FDTD的基本内容。可以先写个模块清晰，结果正确的CPU代码，再考虑并行化。
祝您编码愉快~
ice版主，我就是按您的建议做的啊。循环里的第一个核函数是源更新，第二个核函数是
场更新。现在我是把第二个核函数给简化了，目的是为了发现第二个核函数到底是哪里影响了
第一个核函数的运行。第二个核函数影响就像上面说的那样。它为什么会影响第一个的运行
呢？请您明示啊。这个程序，我cpu已经验证过了，现在就是把它运行到gpu上了。
谢谢您不厌其烦地解答啊，ice版主！
1：您在28#中的说法依然是block的隔离问题，当然这个可能是笔误。
2：您在kernel中用shared memory来存储场值，个人认为是瞎用。具体分析见上面，不再赘述。
3：如果您的电场更新和磁场更新都在第二个kernel里面，建议拆分为两个kernel。如果您坚决地只使用一个block的话，请加上__syncthreads()，以保证更新逻辑的正确性。
您给出的信息并不足以找出您的具体问题所在，我不是神，无法猜到。
以及，您判断的“第二个kernel影响了第一个kernel”，该判断也可能不是实情而只是表象。
所以，您就是再追问我1000次，我也无法回答的。
唯一的建议依然是，请学习/复习CUDA的基本内容，FDTD的基本内容，理清逻辑，仔细调试您的程序。
大致如上，祝您调试顺利~
“您在kernel中用shared memory来存储场值，个人认为是瞎用。”您觉得应该用什么来存储
场值呢？真心求教啊。
显然是global memory的。
我感觉，global memory的速度没有shared memory快，所以就用后者了。
ice版主，如您所说，我更换了存储器之后，第二个核函数就不影响第一个核函数的运行了。
但为什么用第二个核函数使用shared memory 就会影响第一个核函数的运行，依然困惑。再次感
谢您的帮助啊，ice版主！
刚才出去听报告去了，才回来。
shared memory根本就不是干这个事情的。
理由如下：
1：A kernel里面定义的shared memory，B kernel是无法访问到的。
2：同一个kernel里面，a block 的shared memory，b block也是无法访问的。
3：kernel此次运行结束的时候，shared memory里面的东西都没有了。（实际上在一个kernel有多个block的时候，如果当前block运行结束，那么shared memory里面的东西也就清掉了。）
4：shared memory太小，一个SM最多几十KB，FDTD算法稍微上一点规模，也远远比这个大了，所以即使前面各条限制都不存在，shared memory的规模也使得完全不具备实用性。
那么，shared memory是神马呢？
shared memory是一个block内部可以共享的，可以手工操作的，高速cache。
祝您编码愉快~
稍微补充一下，如果在kernel里面直接 int ex[100];这样定义一个数组，那么这定义的实际上是一个线程私有的数组，每个线程都会维护一个私有的这样的数组。
一般会放到local memory里面，local memory也是在显存上的，但是和global memory不同，每个线程只能访问自己的那份。
而存放场值的话，一般都是用global memory的，一般在host端cudaMalloc一下，然后把指针作为参数压给kernel函数使用。
大致如上，祝您调试愉快~
ice版主，谢谢您的详细解答！看君一席话，获益匪浅啊！
提前祝您元旦快乐，新年好运！
客气客气，也祝您新年快乐！常来论坛转转~
[
__global__ void FDTD_SCATTER_Hy( float *Ex, float *Hy, float (*pFE1)[2], float (*pFE2)[2], float (*pFH1)[2], float (*pFH2)[2] ,int *Ob)
{
int tid = threadIdx.x;
int Imin = 230, Imax = 769, Itmin = 350, Itmax = 649;
float FH = 0.5;
if ( (Itmin <= tid) && (tid <= ( Itmax - 1 )) )
{
hy[tid] = pFH1[Ob[tid]][Ob[tid+1]]*hy[tid] - pFH2[Ob[tid]][Ob[tid+1]]*(ex[tid+1] - ex[tid]);
}
if ( tid == ( Itmin - 1 ) )
{
hy[tid] = hy[tid] - FH*( ex[tid+1] - ex[tid] - ein[Itmin]);
}
if ( tid == Itmax )
{
hy[tid] = hy[tid] - FH*( ex[tid+1] - ex[tid] );
}
if ( (Imin <= tid) && (tid <= ( Itmin -2 )) )
{
hy[tid] = hy[tid] - FH*( ex[tid+1] - ex[tid] );
}
if ( (( Itmax + 1 ) <= tid) && (tid <= ( Imax - 1 )) )
{
hy[tid] = hy[tid] - FH*( ex[tid+1] - ex[tid] );
}
__syncthreads();
Ex[tid] = ex[tid];
}
调用函数为： FDTD_SCATTER_Hy<<< 1, 1000 >>>( pEx, pHy, pFE1, pFE2, pFH1, pFH2, pOb );
请问版主，我这段程序如何改进能才能提高它的运行效率啊？还有，我写的这段程序有哪些缺点，可以如何改进啊？请指教。 非常感谢啊！