1

このカーネル関数で共有メモリを利用しようとしていますが、期待したほどパフォーマンスが良くありません。この関数は、私のアプリケーションでは何度も (約 1000 回以上) 呼び出されるため、共有メモリを活用してメモリ レイテンシを回避することを考えていました。しかし、共有メモリを使用しているため、アプリケーションが非常に遅くなったため、明らかに何かが間違っています。
これはカーネルです：

__global__ void AndBitwiseOperation(int* _memory_device, int b1_size, int* b1_memory, int* b2_memory){
int j = 0;

// index GPU - Transaction-wise
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int tid = threadIdx.x;

// shared variable
extern __shared__ int shared_memory_data[];
extern __shared__ int shared_b1_data[];
extern __shared__ int shared_b2_data[];

// copy from global memory into shared memory and sync threads
shared_b1_data[tid] = b1_memory[tid];
shared_b2_data[tid] = b2_memory[tid];
__syncthreads();

// AND each int bitwise
for(j = 0; j < b1_size; j++)
    shared_memory_data[tid] = (shared_b1_data[tid] & shared_b2_data[tid]);

// write result for this block to global memory
_memory_device[i] = shared_memory_data[i];
}

b1 と b2 のサイズは、実行時にしか知ることができない顧客の数に依存するため、共有変数はexternとして宣言されます (ただし、どちらも常に同じサイズです)。
これは私がカーネルを呼び出す方法です：

void Bitmap::And(const Bitmap &b1, const Bitmap &b2)
{

int* _memory_device;
int* b1_memory;
int* b2_memory;

int b1_size = b1.getIntSize();

// allocate memory on GPU
(cudaMalloc((void **)&b1_memory,  _memSizeInt * SIZE_UINT));
(cudaMalloc((void **)&b2_memory,  _memSizeInt * SIZE_UINT));
(cudaMalloc((void **)&_memory_device,  _memSizeInt * SIZE_UINT));

// copy values on GPU
(cudaMemcpy(b1_memory, b1._memory, _memSizeInt * SIZE_UINT, cudaMemcpyHostToDevice ));
(cudaMemcpy(b2_memory, b2._memory, _memSizeInt * SIZE_UINT, cudaMemcpyHostToDevice ));
(cudaMemcpy(_memory_device, _memory, _memSizeInt * SIZE_UINT, cudaMemcpyHostToDevice ));

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

AndBitwiseOperation<<<dimGrid, dimBlock>>>(_memory_device, b1_size, b1_memory, b2_memory);

// return values
(cudaMemcpy(_memory, _memory_device, _memSizeInt * SIZE_UINT, cudaMemcpyDeviceToHost ));

// Free Memory
(cudaFree(b1_memory));
(cudaFree(b2_memory));
(cudaFree(_memory_device));
}

b1 と b2 は、各要素が 4 ビットのビットマップです。要素の数は、顧客の数によって異なります。また、カーネルのパラメーターに問題があります。ブロックまたはスレッドを追加すると、AndBitwiseOperation() が正しい結果を返さないためです。ブロックごとに 1 つのブロックと 1 つのスレッドだけを使用すると、結果は正しくなりますが、カーネルは並列ではありません。
あらゆるアドバイスを歓迎します:)
ありがとう

4

2 に答える 2

4

カーネルが何をしたいのかよくわかりませんでした。

CUDA と GPU プログラミングの詳細を読む必要があります。

間違いを指摘してみました。

  1. 共有メモリ (sm) は、グローバル メモリの読み取りを減らす必要があります。グローバル メモリ (gm) の読み取り操作と書き込み操作をスレッドごとに分析します。

    a. a.グローバル メモリを 2 回読み取り、sm を 2 回書き込み
    ます。(無意味なループは無視され、インデックスは使用されません) sn を 2 回読み取り、sm を 1 回書き込みます
    c. sm を 1 回読み、gm を 1 回書きます。

    したがって、合計で何も得られません。グローバルメモリを直接使用できます。

  2. すべてのスレッドを使用して、ブロック インデックス "i" に 1 つの値を書き出します。このデータを書き出すには、1 つのスレッドのみを使用する必要があります。
    シリアル化される複数のスレッドで同じデータを出力しても意味がありません。

  3. ループを使用し、ループ カウンターをまったく使用しません。

  4. 「tid」で書き込み、「i」でランダムに読み取ります。

  5. この割り当てはオーバーヘッドです。

    unsigned int tid = threadIdx.x;

  6. 1 つのブロック tid = i!
    すべての間違ったインデックス作成により、複数のブロックを使用して間違った計算が行われる

  7. 「i」の共有メモリは書き込まれませんでした!

    _memory_device[i] = shared_memory_data[i];

カーネルが何をすべきかという私の仮定

/*
 * Call kernel with x-block usage and up to 3D Grid
 */
__global__ void bitwiseAnd(int* outData_g, 
    const long long int inSize_s, 
    const int* inData1_g, 
    const int* inData2_g)
{
    //get unique block index
    const unsigned long long int blockId = blockIdx.x //1D
        + blockIdx.y * gridDim.x //2D
        + gridDim.x * gridDim.y * blockIdx.z; //3D

    //get unique thread index
    const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x; 

    //check global unique thread range
    if(threadId >= inSize_s)
        return;

    //output bitwise and
    outData_g[thread] = inData1_g[thread] & inData2_g[thread];
}
于 2012-04-22T02:52:20.330 に答える
4

配列を宣言するときextern __shared__は、カーネル呼び出しでそのサイズも指定する必要があります。

カーネル構成は次のとおりです。

<<< Dg , Db , Ns , S >>>

Nsはextern __shared__配列のサイズで、デフォルトは 0 です。

extern __shared__カーネルで複数の配列を定義できるとは思いません。プログラミング ガイドの例では、単一の配列を定義し、そのextern __shared__中にオフセットを含む配列を手動で設定します。

extern __shared__ float array[]; 
__device__ void func()      // __device__ or __global__ function 
{ 
    short* array0 = (short*)array;  
    float* array1 = (float*)&array0[128]; 
    int*   array2 =   (int*)&array1[64]; 
}
于 2012-04-21T23:32:07.057 に答える