キューで1トランザクションでスレッドごとに4文字を読み取る

Question

最近CUDAを学習しています。そして私はメモリ取引について質問します。 私が理解していることは、各トランザクションで、（同じブロック内の）32個の連続したスレッドが、連続した128バイト（32個の単精度ワード）のメモリに同時にアクセスできることです。これはワープと呼ばれます。 しかし、この例では、各スレッドは常に、変数全体として（4バイト）ワードにアクセスしています。だから私の質問は、グローバルメモリ内の配列がcharの型で定義されている場合、すべての32スレッドがこのメモリにアクセスし、同時に4つの連続するcharを読み取ることができるかどうかです。 、そして、

__global__ 
void kernel(char *d_mask) 
{ 
    extern __shared__ char s_tmp[]; 
    const unsigned int thId = threadIdx.x; 
    const unsigned int elementId = 4 * (threadIdx.x + blockDim.x * blockIdx.x); 

    s_tmp[thId_x] = d_mask[elementId]; 
    s_tmp[1 + thId_x] = d_mask[elementId + 1]; 
    s_tmp[2 + thId_x] = d_mask[elementId + 2]; 
    s_tmp[3 + thId_x] = d_mask[elementId + 3]; 
    __syncthreads(); 

    /* calculation */ 
} 

各スレッドが同時に4つのバイトを読み込みます：

だから、eaxmpleのために私はコードを書く、場合？もしそうでなければ、私はそれをどうすればできますか？私はmemcpyのようなAPIを使うべきですか？

Answer 1

効率的な読み込みを行うには、読み取るバイトを1つのトランザクションにまとめる必要があります。私たちは一般的に、いくつかのコード行を分割することでこれを行うことはできません。

複数の要素を1つのタイプに結合するベクトルタイプがあります。 pay attention to proper alignmentである限り、charまたはunsigned charの配列は、たとえば配列のように扱うことができます。 uchar4は、4つの文字を1つの（32ビット）タイプに結合するベクトルタイプです。あなたはcudaヘッダーファイルvector_types.hとvector_functions.hにもっとたくさんのおいしさを見つけることができます。

とにかく、私たちは「ベクトルロード」を活用するために、このようなあなたのサンプルを再記述することができます。

__global__ 
void kernel(char *d_mask) 
{ 
    extern __shared__ char s_tmp[]; 
    const unsigned int thId = threadIdx.x; 
    const unsigned int elementId = threadIdx.x + blockDim.x * blockIdx.x; 

    uchar4 *s_tmp_v = reinterpret_cast<uchar4 *>(s_tmp); 
    uchar4 *d_mask_v = reinterpret_cast<uchar4 *>(d_mask); 
    s_tmp_v[thId] = d_mask_v[elementId]; 
    __syncthreads(); 

    /* calculation */ 
}