Main Content

coder.gpu.constantMemory

変数を GPU の定数メモリにマッピングするプラグマ

ページ内をすべて折りたたむ

構文

coder.gpu.constantMemory(v)

説明

例

coder.gpu.constantMemory(v) は、変数 v を GPU デバイスの定数メモリ領域にマッピングします。このプラグマは並列化可能なループ内に配置します。GPU Coder™ でそのループのカーネルが生成される場合、v はデバイスの定数メモリ変数に読み込まれます。カーネル内のこの変数へのアクセスと、定数メモリ変数へのアクセスが置き換えられます。カーネル内では、変数 v は読み取り専用でなければなりません。そうでない場合、GPU Coder はこのプラグマを無視します。すべてのスレッドでパラメーター配列または行列のすべての要素にアクセスする場合、このプラグマを使用します。

この関数はコード生成関数です。MATLAB® では効果がありません。

例

すべて折りたたむ

この例では、coder.gpu.constantMemory プラグマを使用して入力を GPU の定数メモリ領域にマッピングする方法を示します。

サイズ 256x256 の a とサイズ 1x3 の定数 k という 2 つの入力を受け入れるエントリポイント関数 myFun を記述します。この関数には、定数を a の各要素に加算する入れ子にされた for ループがあります。カーネルを作成するには、coder.gpu.kernel() プラグマを入れ子にされた for ループの外に配置します。coder.gpu.constantMemory(k) は、読み取り専用入力 k を GPU の定数メモリに配置します。

function b = myFun(a,k)
  b = coder.nullcopy(zeros(size(a)));
  coder.gpu.kernel();
    for j = 1:256
      for i = 1:256
        coder.gpu.constantMemory(k);  
        b(i,j) = a(i,j) + k(1) + k(2) + k(3);
      end
    end
end

MEX コード生成用の構成オブジェクトを作成します。

cfg = coder.gpuConfig('mex');

関数 myFun の入力 a,k のサイズとデータ型を宣言する cell 配列 input を定義します。

input = {ones(256),ones(1,3)}

構成の指定、入力引数の指定、およびコード生成レポートの生成を行うために -config、-args、および -report オプションを使用して、MEX 関数 myFun_mex を生成します。

codegen -config cfg -args input -report myFun

レポートの [C コード] タブで、myFun.cu をクリックします。

次のコード スニペットに示すように __constant__ 修飾子を使用することで、読み取り専用変数 k が const_k として宣言されます。 

/* Variable Definitions */
__constant__ real_T const_k[3];

cudaMemcpyToSymbol 呼び出しは、k の値をホストからデバイスの定数メモリ const_k にコピーします。

  cudaMemcpyToSymbol(const_k, k, 24U, 0U, cudaMemcpyHostToDevice);
  cudaMemcpy(gpu_a, a, 524288U, cudaMemcpyHostToDevice);
  myFun_kernel1<<<dim3(128U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_a, gpu_b);
  cudaMemcpy(b, gpu_b, 524288U, cudaMemcpyDeviceToHost);

カーネル本体から定数 const_k にアクセスし、それを a の各要素に加算します。

static __global__ __launch_bounds__(512, 1) void myFun_kernel1(const real_T *a,
 real_T *b)
{
  int32_T i;
  int32_T j;
  int32_T threadIdX;
  threadIdX = (int32_T)(blockDim.x * blockIdx.x + threadIdx.x);
  i = threadIdX / 256;
  j = threadIdX - i * 256;
  if ((!(j >= 256)) && (!(i >= 256))) {
    b[i + (j << 8)] = ((a[i + (j << 8)] + const_k[0]) + const_k[1]) + const_k[2];
  }
}

入力引数

すべて折りたたむ

GPU デバイスの定数メモリ領域にマッピングされなければならない変数の名前。

バージョン履歴

R2017b で導入

参考

アプリ

関数

オブジェクト

トピック