2

次々に実行される2つのcudaカーネルがあります：

__global__
void calculate_histo(const float* const d_logLuminance,
        unsigned int* d_histogram,
        float min_logLum,
        float lumRange,
        int numBins,
        int num_elements){
    extern __shared__ float sdata[];
    int tid = threadIdx.x;
    int bid = blockIdx.x;
    int gid = tid * blockDim.x + bid;

    // load input into __shared__ memory
    if(gid < num_elements)
    {
        sdata[tid] = d_logLuminance[gid];
        __syncthreads();

        //compute bin value of input
        int bin = static_cast <int> (floor((d_logLuminance[gid]-min_logLum)/ lumRange * numBins));
        //increment histogram at bin value
        atomicAdd(&(d_histogram[bin]), 1);
    }
}

__global__
void blelloch_scan(unsigned int* const d_cdf, unsigned int* d_histogram, int numBins) {
    extern __shared__ unsigned int sdata[];// allocated on invocation
    int thid = threadIdx.x;
    //printf("%i \n", thid);
    //printf("%i \n", d_histogram[thid]);

    int offset = 1;


    sdata[2*thid] = d_histogram[2*thid]; // load input into shared memory
    sdata[2*thid+1] = d_histogram[2*thid+1];

    // build sum in place up the tree
    for (int d = numBins>>1; d > 0; d >>= 1) {
        __syncthreads();
        if (thid < d) {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            sdata[bi] += sdata[ai];
        }
        offset *= 2;
    }
    if (thid == 0) { sdata[numBins - 1] = 0; } // clear the last element
    // traverse down tree & build scan
    for (int d = 1; d < numBins; d *= 2) {
        offset >>= 1;
        __syncthreads();
        if (thid < d) {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            float t = sdata[ai];
            sdata[ai] = sdata[bi];
            sdata[bi] += t;
        }
        __syncthreads();
        d_cdf[2*thid] = sdata[2*thid]; // write results to device memory
        d_cdf[2*thid+1] = sdata[2*thid+1];
    }

}

どちらも共有メモリを使用します。2 つ目は、共有メモリとして unsigned int 配列を持っています。1 つ目は float 配列です。カーネルが起動するたびに共有メモリがクリアされるため、両方の配列に同じ変数名 sdata を再利用できるはずだと思っていましたが、次のエラーが表示されます。

declaration is incompatible with previous 'sdata'

カーネルごとに異なる変数名を使用すると、問題が解決するようです。同じ変数名を再利用できない理由を知っている人はいますか?

4

1 に答える 1

1

CUDA は、標準 C 言語のルールに従っているだけです。Kernighan と Ritchie の「The C Programming Language」の本を引用します。

外部変数は、関数の外で一度だけ定義する必要があります。これにより、ストレージが確保されます。変数は、それにアクセスする各関数でも宣言する必要があります。これは、変数の型を示します。[...]定義とは、変数が作成される場所またはストレージが割り当てられる場所を指します。宣言は、変数の性質が述べられているが、ストレージが割り当てられていない場所を指します。

プログラムのどこかに次のようなものがあるはずです

extern __shared__ unsigned int sdata[];

その場所で、 へのという名前のポインタを作成しています。関数内での型を宣言しているため、関数はそれを認識することができます。の中にsdataunsigned int__global__sdata__global__

kernel<<<blocks,threads,numbytes_for_shared>>>(...);

が指す配列のスペースを割り当てていますsdata。

于 2013-10-12T20:42:49.350 に答える