cuda - CUDA 共有メモリに関する奇妙な動作

Question

使用可能な共有メモリを 16384 バイトと報告する GTX9800 を使用しています

次のカーネル コードを指定して、T = int (4 バイト) で実行します。

template <typename T>
__global__ void foo(unsigned n, T *x)
{
    unsigned idx    = blockIdx.x * blockDim.x + threadIdx.x;
    __shared__ T sx[4096];
    x[idx] = 0;
}

期待される結果が得られます。つまり、最初はゼロではない配列 x がゼロで埋められます。

ただし、何もしないコード行を追加します。

template <typename T>
__global__ void foo(unsigned n, T *x)
{
    unsigned idx    = blockIdx.x * blockDim.x + threadIdx.x;
    __shared__ T sx[4096];
    sx[0] = 0;
    x[idx] = 0;
}

カーネルを呼び出した後、x にはゼロがまったく含まれなくなりました。

ただし、sx のサイズを <= 4088 に変更すると、期待どおりの結果が得られます。

どうしたの？私はかなり混乱しています。

編集：

タイプミスの修正: 16384 "KB" を "bytes" に変更