使用して、両方の動的に割り当てられたと静的に割り当てられた共有メモリ

Question

それぞれが次のローカル変数を持つ、私は2つの__device__ CUDA機能があるとします。

__shared__ int a[123]; 

と別の関数（__global__機能、すなわち、それは私のカーネルだと言います） 、with：

extern __shared__ int b[]; 

これはnVIDIAによって明示的に許可/禁止されていますか？ （programming guideセクションB.2.3の__shared__には表示されません）サイズはすべて共に共有メモリの制限に合わせて集計されますか、または最大で1回で使用されますか？それとも何か他のルール？

_{これは、this oneへのフォローアップの問題と考えることができます。}

Answer 1

共有メモリは、静的に割り当てられ、動的に割り当てられる2つの部分に分割されます。最初の部分は、コンパイル時に計算され、各宣言は、実際の配分である - コンパイル時の情報ptxas活性化ここでそれを示しています。ここ

ptxas info : Used 22 registers, 384 bytes smem, 48 bytes cmem[0] 

、我々は32 intの3配列である384バイトを、持っています。 （下のサンプルコードを参照）。

Keplerから共有メモリへのポインタを、別の関数に渡して、デバイスのサブ関数が別の共有メモリ宣言にアクセスできるようにすることができます。

次に、動的に割り当てられた共有メモリがあります。予約されたサイズはカーネルコール中に宣言されます。

ここでは、2つの機能のさまざまな用途の例を示します。各共有メモリ領域のポインタ値に注意してください。

__device__ void dev1() 
{ 
    __shared__ int a[32] ; 
    a[threadIdx.x] = threadIdx.x ; 

    if (threadIdx.x == 0) 
     printf ("dev1 : %x\n", a) ; 
} 

__device__ void dev2() 
{ 
    __shared__ int a[32] ; 
    a[threadIdx.x] = threadIdx.x * 5 ; 

    if (threadIdx.x == 0) 
     printf ("dev2 : %x\n", a) ; 
} 

__global__ void kernel(int* res, int* res2) 
{ 
    __shared__ int a[32] ; 
    extern __shared__ int b[]; 

    a[threadIdx.x] = 0 ; 
    b[threadIdx.x] = threadIdx.x * 3 ; 

    dev1(); 
    __syncthreads(); 
    dev2(); 
    __syncthreads(); 

    res[threadIdx.x] = a[threadIdx.x] ; 
    res2[threadIdx.x] = b[threadIdx.x] ; 

    if (threadIdx.x == 0) 
     printf ("global a : %x\n", a) ; 
    if (threadIdx.x == 0) 
     printf ("global b : %x\n", b) ; 
} 

int main() 
{ 
    int* dres ; 
    int* dres2 ; 

    cudaMalloc <> (&dres, 32*sizeof(int)) ; 
    cudaMalloc <> (&dres2, 32*sizeof(int)) ; 

    kernel<<<1,32,32*sizeof(float)>>> (dres, dres2); 

    int hres[32] ; 
    int hres2[32] ; 

    cudaMemcpy (hres, dres, 32 * sizeof(int), cudaMemcpyDeviceToHost) ; 
    cudaMemcpy (hres2, dres2, 32 * sizeof(int), cudaMemcpyDeviceToHost) ; 

    for (int k = 0 ; k < 32 ; ++k) 
    { 
     printf ("%d -- %d \n", hres[k], hres2[k]) ; 
    } 
    return 0 ; 
} 

このコードは384 bytes smemを用いptxas情報を出力し、それは、グローバルaアレイ、DEV1方法に対する第aアレイ、及びDEV2方法aアレイのための第三のための1つのアレイです。合計3*32*sizeof(float)=384 bytes。

ダイナミック共有メモリを使用してカーネルを実行する場合、32*sizeof(float)に等しい場合、これらの3つのアレイの直後にbへのポインタが開始されます。

EDIT： このコードによって生成されたPTXファイルは、それが方法

の本体内に定義されているエントリポイントを除い

.shared .align 4 .b8 _ZZ4dev1vE1a[128]; 
.shared .align 4 .b8 _ZZ4dev2vE1a[128]; 
.extern .shared .align 4 .b8 b[]; 

を静的に定義された共有メモリの宣言を保持します

：

// _ZZ6kernelPiS_E1a has been demoted 

メモリの共有スペースは、PTXのドキュメントhereに定義されています210

共有（.shared）状態空間は、CTA内のスレッドがデータを共有するためのCTA単位の領域です。共有メモリ内のアドレスは、CTA内の任意のスレッドによって読み書きできます。共有変数にアクセスするには、ld.sharedとst.sharedを使用します。

実行時には詳細はありませんが、プログラミングガイドhereには、2つのミキシングの詳細はありません。

PTXコンパイル時に、コンパイラは静的に割り当てられた共有メモリの量を知ることがあります。いくつかの補足的な魔法があるかもしれません。 SASSを見て、最初の命令は、（非常にstackallocの形のように見える）静的に割り当てられた共有メモリに異なる値を割り当てる逆順でSR_LMEMHIOFF

1    IADD32I R1, R1, -0x8; 
2    S2R R0, SR_LMEMHIOFF; 
3    ISETP.GE.U32.AND P0, PT, R1, R0, PT; 

および関数を呼び出すを使用します。

ptxasコンパイラは、すべてのメソッドを呼び出すことができる最悪の場合（メソッドの1つを使用せず、関数ポインタを使用し、bのアドレスが変更されず、割り当てられていない共有メモリ領域は決してアクセスされない）。

最後に、einpoklumはコメントで示唆しているように、これは実験的であり、ノルム/ API定義の一部ではありません。