cuda - 共有メモリを使用せずに行列要素の動径平均を効率的に実行する方法

Question

CUDA で行列要素の放射平均を実装しようとしています。ここでは、すべての行列要素のすべての隣接要素 (それ自体を含む) の平均を見つけて出力する必要があります。以下は私が最終的に得たものです（半径= 1の場合）：

__global__ void matrix_avg(float *ad,float *bd)
{
    float t1, t2, t3, t4, t5, t6, t7, t8, t9, avg;
    int i = blockIdx.y * blockDim.y + threadIdx.y;
    int j = blockIdx.x * blockDim.x + threadIdx.x;
    int k = (i*N)+(j);

    if(i==0)
        {
            if(j==0)
            {
                bd[k]=(ad[k+1]+ad[k+N]+ad[k+N+1]+ad[k])/4;
            }
            else if(j==N-1)
            {
                bd[k]=(ad[k-1]+ad[k+N]+ad[k+N-1]+ad[k])/4;
            }
            else
            {
                bd[k]=(ad[k-1]+ad[k+1]+ad[k+N-1]+ad[k+N]+ad[k+N+1]+ad[k])/6;
            }
        }
        else if(i==N-1)
        {
            if(j==0)
            {
                bd[k]=(ad[k+1]+ad[k-N]+ad[k-N+1]+ad[k])/4;
            }
            else if(j==N-1)
            {
                bd[k]=(ad[k-1]+ad[k-N]+ad[k-N-1]+ad[k])/4;
            }
            else
            {
                bd[k]=(ad[k-1]+ad[k+1]+ad[k-N-1]+ad[k-N]+ad[k-N+1]+ad[k])/6;
            }
        }
        else if(j==0)
        {
            bd[k]=(ad[k-N]+ad[k-N+1]+ad[k+1]+ad[k+N]+ad[k+N+1]+ad[k])/6;
        }
        else if(j==N-1)
        {
            bd[k]=(ad[k-N-1]+ad[k-N]+ad[k-1]+ad[k+N-1]+ad[k+N]+ad[k])/6;
        }
        else
        {
            t1=ad[k-N-1];
            t2=ad[k-N];
            t3=ad[k-N+1];
            t4=ad[k-1];
            t5=ad[k+1];
            t6=ad[k+N-1];
            t7=ad[k+N];
            t8=ad[k+N+1];
            t9=ad[k];
            avg=(t1+t2+t3+t4+t5+t6+t7+t8+t9)/9;
            bd[k]=avg;
        }
}

上記のコードは、一番上の行、一番下の行、一番右と一番左の列要素の条件をチェックし、6 要素の平均を計算する必要があります。4つの要素の平均を計算する必要がある4つのコーナー要素についても答えます。残りの内部要素については、9 つ​​の要素の平均を計算する必要があります。上記のコードは、単純に C を CUDA プログラムに変換したものです。共有メモリを使用せずにプログラムを作成する最も効率的な方法を探しています。任意の半径について。任意のアルゴリズム、疑似コード、または提案で十分です。前もって感謝します。