__synthreads()同步出错

[zaviichen]

观察窗的变量vWatch在不同步的时候值是正确的，但是__syn以后的值完全变了，这是什么原因啊？

我看了论坛中好像也有人遇到相似的问题，好像没有解决办法啊。。。

代码如下，这里应该没有数组越界的问题。

__global__
void kernel_sim64_per_layer( uint   nNode,
                             ulong* vSim,
                             uint*  vTruth,
                             uint*  vNode,
                             uint*  vFaninOff,
                             uint*  vFanin,
                             ulong* vWatch )
{
    const unsigned int tid_g = blockDim.x * blockIdx.x + threadIdx.x;    
    
    if(tid_g >= nNode) return;
    
    uint nodeid  = vNode[tid_g];
    uint truth   = vTruth[nodeid];
    uint offset  = vFaninOff[tid_g];
    uint nFanins = vFaninOff[tid_g+1] - vFaninOff[tid_g];

    __syncthreads();

    // watch window for debug
    vWatch[tid_g*3 + 0] = nodeid;
    vWatch[tid_g*3 + 1] = tid_g;
    vWatch[tid_g*3 + 2] = offset;
}

[Cyrosly]

这一句： if(tid_g >= nNode) return; 如果一个block内的某些线程走这些路径，那么这些线程永远无法达到”同步点“，当然会死



另外，这里使用同步多次一举

如果非要使用,如下修改：
if(tid_g < nNode){

uint nodeid = vNode[tid_g];
uint truth = vTruth[nodeid];
uint offset = vFaninOff[tid_g];
uint nFanins = vFaninOff[tid_g+1] - vFaninOff[tid_g];
} __syncthreads();
即使这佯作毫无意义

[无心人_过过小日子]

[Quote=引用 1 楼 cyrosly 的回复:]
如果非要使用,如下修改：
  if(tid_g < nNode){
   
    uint nodeid  = vNode[tid_g];
    uint truth  = vTruth[nodeid];
    uint offset  = vFaninOff[tid_g];
    uint nFanins = vFaninOff[tid_g+1] - vFaninOff[tid_g];
  } __syncthreads();
即使这佯作毫无意义
[/Quote]

uint nodeid;
uint truth;
uint offset;
uint nFanins;

if(tid_g < nNode){

nodeid = vNode[tid_g];
truth = vTruth[nodeid];
offset = vFaninOff[tid_g];
nFanins = vFaninOff[tid_g+1] - vFaninOff[tid_g];
} __syncthreads();

呵呵.

[zaviichen]

谢谢你们，其实真正的程序不止这么多的，在前面的一堆赋值语句和__syn中还有一些要执行的代码。但是因为计算结果出错了，我把中间的代码都给注释了，最后发现问题是出在同步语句上的。
就以上这么简单的代码，不加同步语句的话vWatch中的值是正确的，加上同步以后值就混乱了。我的kernel没有停在等待同步的地方，它能够结束，就是结果不正确了。
我在CUDA的官方论坛上看到的是可以用return来提前结束无用的线程，这种情况下是不是不能用__syn()啊？
因为开启的线程数很有可能大于要计算的节点数目(nNode)，那这种情况该怎么处理好啊？如果把下面所有的程序代码都放到一个if的括号里的话，总觉得不太舒服。。。

[Cyrosly]

[Quote=引用 4 楼 l7331014 的回复:]
1)把nNode调整为32的倍数.
2)申请时适当扩大点计算用的空间(gmem等),让通过1)扩大后的线程如果计算的话也不越界.

这样你应该可以放心用return了.呵呵.

*:提示：__syn的有效范围是warp.
[/Quote]

1，2:这样做有些时候的确是可以的，但是LZ得程序中有用到动态索引(nodeid = vNode[tid_g])所以不敢确定是否索引也会在这个范围内。

3 同步的有效范围是CTA内的warp间，当然在warp内部使用同步也不会有问题，至于手册上所说的单独的warp不需要同步是需要仔细理解的，举个例子：

#if CTA_DIM>32
  #define __barrier { __syncthreads(); }
#else
  #define __barrier
#endif

inline __device__
void warp_reduce( TYPE* smem, uint lane )
{
    ...
}

inline __device__
void block_reduce( TYPE* smem )
{
    const uint lane=threadIdx.x&31u;
    const uint warpid=threadIdx.x>>5;

    warp_reduce( smem, lane ); __barrier

    if( lane==0 ){
        smem[ warpid ]=smem[ threadIdx.x ];
    } __barrier 
    //如果CTA_DIM>0,则必须同步

    但是对于上面的一段代码，我们可以通过如下修改去掉同步，当然会有很少的bank-conflicts:

    if( threadIdx.x<blockDim.x>>5 ){
        smem[ threadIdx.x ]=smem[ warpid<<5 ];
    }

    //最后的reduce-op (这里每个循环步都不需要同步以及上面两段代码是安全的是基于warp大小是32，而最大CTA尺寸是
    //512的事实，虽然事实上1024的block大小也是正确的（如果支持的话），但如果将来支持的最大CTA尺寸超过在各个大小
    //则还需要做些修改（当然，如果那时要用到那么大的尺寸的话，否则简单将CTA尺寸限制在合理的范围 就可以了
    for( uint n=blockDim.x>>6; n>0; n>>=1 ){
        if( threadIdx.x<n ){
            reduceOp( smem );
        }
    } 
    //如果后续的代码中CTA_DIM内的多个线程需要使用这个reduced结果则仍需要再次同步: __barrier
}
    

[Cyrosly]

另外LZ的那句话：“好像没有解决办法啊。。。 ”

你不贴完整代码怎么知道我们没办法解决

[无心人_过过小日子]

5楼的3是正确的.我一时搞错了.呵呵.

[zaviichen]

不好意思，我把相关部分的代码贴上来。我完全是个新手，写的很挫。。。希望大家帮帮忙啊！

我主要是利用CUDA来进行电路仿真，我在另一个帖中《kernel launch 请教》已经提过了。主要是想把一个电路分成多层，每层电路启动一次kernel进行计算。

这部分是kernel的代码：

__device__
void kernel_sim64_onenode( ulong& tmp_u64,
                           uint   truth,
                           uint   nFanins,
                           ulong* vSimFanin )
{
    tmp_u64 = 0;
    ulong minTerm_u64;

    for(int truthEntry = 0; truthEntry < (1<<nFanins); truthEntry++)
    {
        if(!TruthHasBit(truth, truthEntry)) continue;
        minTerm_u64 = 0xffffffffffffffff;
        for(int fanin = 0; fanin < nFanins; fanin++)
        {
            if(truthEntry & (1 << fanin))
                minTerm_u64 &= vSimFanin[fanin];            
            else
                minTerm_u64 &= ~vSimFanin[fanin];
        }
        tmp_u64 |= minTerm_u64;
    } 
}

__global__
void kernel_sim64_per_layer( uint   nNode,
                             ulong* vSim,
                             uint*  vTruth,
                             uint*  vNode,
                             uint*  vFaninOff,
                             uint*  vFanin,
                             ulong* vWatch )
{
    const unsigned int tid_g = blockDim.x * blockIdx.x + threadIdx.x;    
    
    if(tid_g >= nNode) return;

    uint nodeid  = vNode[tid_g];
    uint truth   = vTruth[nodeid];
    uint offset  = vFaninOff[tid_g];
    uint nFanins = vFaninOff[tid_g+1] - vFaninOff[tid_g];

    ulong tmp_u64 = 0;
    
    // assume fanin number <= 4
    __shared__ ulong vSimFanin[4];

    for(int fanin = 0; fanin < nFanins; fanin++) {
        vSimFanin[fanin] = vSim[FaninID(vFanin, offset, fanin)];
    }

    kernel_sim64_onenode(tmp_u64, truth, nFanins, vSimFanin);

    //__syncthreads();
    
    vSim[nodeid] = tmp_u64;

    // watch window for debug
    vWatch[tid_g*3 + 0] = nodeid;
    vWatch[tid_g*3 + 1] = vSimFanin[0];
    vWatch[tid_g*3 + 2] = vSimFanin[1];
}

这部分是调用的代码：

void cuda_simulate_once()
{
    int i, j, nSim;
 
    nBlock  = (nNode / 256) + 1;
    nThread = 256;

    dim3 dimGrid (nBlock,  1, 1);
    dim3 dimBlock(nThread, 1, 1);

    Abc_NtkForEachCi(pNtk, pNode, j)
    {
        h_vSim[pNode->Id] = h_vRandCi[j];
    }

#ifdef _CZH_DEBUG_  
    printf("\n ini per-layer h_vSim: \n");
    print_ulong_array(h_vSim, nObj);
#endif
        
    for(i = 0; i < (level + nSimIters - 1); i++)
    {
        // malloc the watch window
        sWatch = WATCH_NUM * nNode * sizeof(ulong);
        cutilSafeCall( cudaMallocHost((void**)&h_vWatch, sWatch) );
        cudaMalloc((void**)&d_vWatch, sWatch);

        cudaMemcpy(d_vSim, h_vSim, sSim, cudaMemcpyHostToDevice);
    
        kernel_sim64_per_layer <<< dimGrid, dimBlock >>>( nNode,
                                                          d_vSim,
                                                          d_vTruth,
                                                          d_vNode,
                                                          d_vFaninOff,
                                                          d_vFanin,
                                                          d_vWatch );
            
        CUDA_SAFE_CALL(cudaMemcpy(h_vSim,   d_vSim,   sSim,   cudaMemcpyDeviceToHost));
        CUDA_SAFE_CALL(cudaMemcpy(h_vWatch, d_vWatch, sWatch, cudaMemcpyDeviceToHost));

#ifdef _CZH_DEBUG_
        printf("\n Watch Window: \n");
        for(j = 0; j < nNode; j++){
            for(int k = 0; k < WATCH_NUM; k++){
                printf("w%d: %6ld\t", k, h_vWatch[j*WATCH_NUM + k]);
            }
            printf("\n");
        }
#endif

#ifdef _CZH_DEBUG_  
        printf("\n i=%d, nSim: %d, per-layer h_vSim: \n", i, nSim);
        print_ulong_array(h_vSim, nObj);
#endif

        if(i >= level-1) {
            nSim = i - (level-1);
            // record the simulation results for all COs in the golden case
            Abc_NtkForEachCo(pNtk, pNode, j)
            {
                h_vGoldenCo[nSim*nCo + j] = h_vSim[Abc_ObjFanin0(pNode)->Id];
            }
        }

        if(i+1 < nSimIters) {
            Abc_NtkForEachCi(pNtk, pNode, j)
            {
                h_vSim[pNode->Id] = h_vRandCi[(i+1) * nCi + j];
            }
#ifdef _CZH_DEBUG_  
            printf("\n i: %d,  h_vSim: \n", i);
            print_ulong_array(h_vSim, nObj);
#endif
        }
    }
}

其中电路每个节点的计算信息都是存放在global mem的vSim中，每一个线程对应电路的一个节点。

我的想法是有几层电路结构就开启几次kernel，后来要多次仿真，我就引入流水线的方法，即第一层电路的值不断改变。

初始时只有第一层电路有输入值，其他层的电路计算都是无效的。假设电路一共有n层，也就是在启动了n次kernel后会在最后一层节点处得到输出值。但是好像出乎我的意料，为什么只启动了一次kernel就把全部值给算出来了。。。但加上__syn以后，所有值都是初始值。

谢谢大家！

[Cyrosly]

一看就知道你这个程序（内核）本身就是错误的,这一句：

__shared__ ulong vSimFanin[4];

再看下面的代码，如果CTA内的所有线程都使用同一个vSimFainin[4]，则应该使用if

for(int fanin = 0; fanin < nFanins; fanin++) {
vSimFanin[fanin] = vSim[FaninID(vFanin, offset, fanin)];
}

if( threadIdx.x<nFanins ){
vSimFanin[fanin]=vSim[FaninID(vFanin, offset, fanin)];
}

否则，共享内存的大小应该为：
__shared__ vSimFainin[ 4*block_size ];

如果是第一种情况则只需在if之后同步一次，以后不需要再同步，对于第二种情况则根本就不需要同步，甚至共享内存的使用都是不必要的（如果寄存器压力适当的话）

[Cyrosly]

所以，事实上你这个程序的错误根本就和同步没有关系，是算法和对CUDA架构下实施理解上的错误

[zaviichen]

非常感谢！刚学习，对这些用法有些模棱两可。。。

但是我把那句改成了：

ulong vSimFanin[4];

这里是对于每个线程都是要开这4个空间的，用寄存器来处理。但是这样的话，加上__syn还是不能够正确计算，出来的结果都是初始值。

还有一个很莫名的情况，假设电路一共有7个节点，其中1,2号是电路输入，3号是电路输出，4-8号是电路的内部节点。其中4,5为第一层，6,7为第二层。则电路大概是这样的：

1 -- 4 -- 6
\
--3
/
2 -- 5 -- 7

我想启动2次kernel来分别仿真45和67节点，但是我传给kernel数据的时候是把整个电路传进去了，而且每次启动是4个线程（分别对应4-7节点），在第一次启动kernel的时候我觉得45的结果应该是正确的，而67则是读到的是45的初始值，结果应该不对。等第一次kernel结束后，45的值计算出来了，在第二次启动kernel才能得到67的正确值。但是实际运行结果是在第1次就把4-7的结果都算出来了，而且都是正确的。给我的感觉好像是CPU的顺序运行那样，我觉得GPU一起运行4个线程的话67的值在第1次应该是错的。

不好意思，刚学碰到一大堆的问题。。。

[天下第一好大人]

mark, Cyrosly学cuda多久了？

[无心人_过过小日子]

1)把nNode调整为32的倍数.
2)申请时适当扩大点计算用的空间(gmem等),让通过1)扩大后的线程如果计算的话也不越界.

这样你应该可以放心用return了.呵呵.

*:提示：__syn的有效范围是warp.

[无心人_过过小日子]

[Quote=引用 11 楼 cyrosly 的回复:]
所以，事实上你这个程序的错误根本就和同步没有关系，是算法和对CUDA架构下实施理解上的错误
[/Quote]

LZ的程序应该如下的构架较合理:
每个线程模拟一个器件.根据输入计算输出.(输入和输出存放在gmem中)
然后每一轮执行模拟一步脉冲,知道第一道脉冲通过最后的器件.

[diaolingdeyu]

关注中~
·