4

計算機能 1.3 GPU 上の CUDA カーネルでグローバル メモリのメモリ スループットを向上させるには、一般に 2 つの手法があります。メモリーは合体にアクセスし、少なくとも 4 バイトのワードにアクセスします。最初の手法では、同じハーフ ワープのスレッドによる同じメモリ セグメントへのアクセスは、より少ないトランザクションに結合されますが、少なくとも 4 バイトのワードにアクセスすると、このメモリ セグメントは 32 バイトから 128 バイトに効果的に増加します。

更新: talonmies answer に基づくソリューション。グローバル メモリに unsigned char が格納されている場合に 1 バイト ワードではなく 16 バイト ワードにアクセスするには、通常、メモリ配列を uint4 にキャストすることによって uint4 ベクトルが使用されます。uint4 ベクトルから値を取得するには、以下のように uchar4 に再キャストできます。

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>

__global__ void kernel ( unsigned char *d_text, unsigned char *d_out ) {

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ unsigned char s_array[];

    uint4 *uint4_text = reinterpret_cast<uint4 *>(d_text);
    uint4 uint4_var;

    //memory transaction
    uint4_var = uint4_text[0];

    //recast data to uchar4
    uchar4 c0 = *reinterpret_cast<uchar4 *>(&uint4_var.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&uint4_var.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&uint4_var.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&uint4_var.w);

    d_out[idx] = c0.y;
}

int main ( void ) {

    unsigned char *d_text, *d_out;

    unsigned char *h_out = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) );
    unsigned char *h_text = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) );

    int i;

    for ( i = 0; i < 16; i++ )
            h_text[i] = 65 + i;

    cudaMalloc ( ( void** ) &d_text, 16 * sizeof ( unsigned char ) );
    cudaMalloc ( ( void** ) &d_out, 16 * sizeof ( unsigned char ) );

    cudaMemcpy ( d_text, h_text, 16 * sizeof ( unsigned char ), cudaMemcpyHostToDevice );

    kernel<<<1,16>>>(d_text, d_out );

    cudaMemcpy ( h_out, d_out, 16 * sizeof ( unsigned char ), cudaMemcpyDeviceToHost );

    for ( i = 0; i < 16; i++ )
            printf("%c\n", h_out[i]);

    return 0;
}
4

2 に答える 2

3

char*へのキャストは問題なく機能します。試しましたか？もしそうなら、この質問を促したのは何が起こったのですか？

s_arrayあなたの例では、にキャストしてint*から1つのコピーを実行できるように見えます（ 16ではなく4をvar.x掛ける）。j

ワード内のバイトをより柔軟にシャッフルする必要がある場合は、__byte_perm()組み込み関数を使用できます。たとえば、整数のバイトの順序を逆にするには、次のようにしますx。__byte_perm(x, 0, 0x0123);

バイトを格納するためにベクトル型、または単一のintを使用しても、何も得られない可能性があります。Fermiでは、グローバルメモリトランザクションは128バイト幅です。したがって、ワープがグローバルメモリとの間でロード/ストアを実行する命令にヒットすると、GPUは32スレッドにサービスを提供するために必要な数の128バイトのトランザクションを実行します。パフォーマンスは、各スレッドがそのバイトをロードまたは格納する方法ではなく、必要な個別のトランザクションの数に大きく依存します。

于 2012-10-27T18:43:35.290 に答える
3

あなたがやろうとしていることを理解した場合、論理的なアプローチは、C++reinterpret_castメカニズムを使用してコンパイラに正しいベクトルロード命令を生成させ、CUDA組み込みのバイトサイズのベクトル型uchar4を使用して、4つのそれぞれの各バイト内の各バイトにアクセスすることですグローバル メモリからロードされた 32 ビット ワード。このアプローチを使用すると、各 32 ビット レジスタ内でバイト単位のアクセスを行う最適な方法を知っているコンパイラを本当に信頼できます。

完全に不自然な例は次のようになります。

#include <cstdio>
#include <cstdlib>

__global__
void kernel(unsigned int *in, unsigned char* out)
{
    int tid = threadIdx.x;

    uint4* p = reinterpret_cast<uint4*>(in);
    uint4  i4 = p[tid]; // vector load here

    uchar4 c0 = *reinterpret_cast<uchar4 *>(&i4.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&i4.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&i4.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&i4.w);

    out[tid*4+0] = c0.x;
    out[tid*4+1] = c4.y;
    out[tid*4+2] = c8.z;
    out[tid*4+3] = c12.w;
}

int main(void)
{
    unsigned int c[8] = { 
        2021161062, 2021158776, 2020964472, 1920497784, 
        2021161058, 2021161336, 2020898936, 1702393976 };

    unsigned int * _c;
    cudaMalloc((void **)&_c, sizeof(int)*size_t(8));
    cudaMemcpy(_c, c, sizeof(int)*size_t(8), cudaMemcpyHostToDevice);
    unsigned char * _m;
    cudaMalloc((void **)&_m, sizeof(unsigned char)*size_t(8));

    kernel<<<1,2>>>(_c, _m);

    unsigned char m[8];
    cudaMemcpy(m, _m, sizeof(unsigned char)*size_t(8), cudaMemcpyDeviceToHost);

    for(int i=0; i<8; i++)
        fprintf(stdout, "%d %c\n", i, m[i]);

    return 0;
}

これにより、カーネルに提供される符号なし整数の配列に埋め込まれた読み取り可能な文字列が生成されます。

1 つの注意点は、compute 1.x ターゲットに使用される open64 コンパイラは、ベクトル内のすべての単語が実際に使用されたわけではないことを検出できる場合、ベクトル ロードを生成しようとするこの戦略を無効にすることが多いということです。そのため、入力ベクトル型のすべての入力単語に触れて、コンパイラが適切に動作するようにしてください。

于 2012-10-28T12:27:06.780 に答える