cuda - CUDA スレッドが同期していないように見える

Question

syncthreads を使用しているにもかかわらず、1 つのスレッドが残りのスレッドより遅れているように見えるという問題があります。次の抜粋は、大規模なプログラムから抜粋したものです。ここでは、できる限り多くを切り取っていますが、それでも問題が再現されます。私が見つけたのは、このコードを実行すると、test4 変数がすべてのスレッドに対して同じ値を返さないということです。私の理解では、TEST_FLAG 変数を使用すると、すべてのスレッドがif (TEST_FLAG == 2)条件に導かれ、配列 test4 のすべての要素が 43 の値を返す必要があります。すべてのスレッドが同じ同期スレッドに到達していないように見えます。多数のテストを実行したところ、次のようなコードをさらに削除することがわかりました。for (l=0; l<1; ++l)ループは問題を解決しますが、その理由がわかりません。私のスレッドがすべて同じ値を返さない理由についての助けをいただければ幸いです。

import numpy as np
import pycuda.driver as drv
import pycuda.compiler
import pycuda.autoinit
import pycuda.gpuarray as gpuarray
import pycuda.cumath as cumath
from pycuda.compiler import SourceModule

gpu_code=SourceModule("""
    __global__ void test_sync(double *test4, double *test5)
    {
        __shared__ double rad_loc[2], boundary[2], boundary_limb_edge[2];
        __shared__ int TEST_FLAG;
        int l;


        if (blockIdx.x != 0)
        {
            return;
        }

        if(threadIdx.x == 0)
        {
            TEST_FLAG = 2;
            boundary[0] = 1;
        }

        test4[threadIdx.x] = 0;
        test5[threadIdx.x] = 0;

        if (threadIdx.x == 0)
        {
            rad_loc[0] = 0.0;
        }
        __syncthreads();

        for (l=0; l<1; ++l)
        {
            __syncthreads();
            if (rad_loc[0] > 0.0)
            {
                test5[threadIdx.x] += 1;

                if ((int)boundary[0] == -1)
                {
                    __syncthreads();
                    continue;
                }
            }
            else
            {
                if (threadIdx.x == 0)
                {
                    boundary_limb_edge[0] = 0.0;
                }
            }
            __syncthreads();

            if (TEST_FLAG == 2)
            {
                test4[threadIdx.x] = 43;
                __syncthreads();

                TEST_FLAG = 99;
            }
            __syncthreads();

            return;
        }
        return;
    }

    """)


test_sync = gpu_code.get_function("test_sync")

DATA_ROWS=[100,100]

blockshape_data_mags    = (int(64),1, 1)
gridshape_data_mags     = (int(sum(DATA_ROWS)), 1)

test4 = np.zeros([1*blockshape_data_mags[0]], np.float64)
test5 = np.zeros([1*blockshape_data_mags[0]], np.float64)

test_sync(drv.InOut(test4), drv.InOut(test5), block=blockshape_data_mags, grid=gridshape_data_mags)

print test4
print test5