AMDデバイス上の物理メモリ：ローカル対プライベート

Question

私はOpenLCLでアルゴリズムを書いていますが、かなりの量のデータを記憶する必要があります。long[70]とlong[200]かそれ以上のカーネルがあります。

最近のAMDデバイスには32 KiB __localメモリがあります。これは、20〜58個の作業単位の情報を格納するのに十分な量です（カーネルあたりの所与のデータ量）。しかし、私がアーキテクチャから理解しているもの（特にthis drawing）から、各シェーダコアには専用メモリが専用に用意されています。私はしかし、そのサイズを見つけることに失敗します。

各カーネルが持つプライベートメモリの量を知る方法を教えてもらえますか？

私は特にこれらのいくつかをすぐに購入する予定であるので、私はHD7970について特に興味があります。

編集：問題解決、答えは付録のhereあるD.

Answer 1

コメントはユーザーの嫌悪感によってコメントに書かれていますので、私は質問を閉じるためにここで新しい答えに書きます。

これらの値は、AMD APP OpenCLプログラミングガイドhttp://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdfの付録Dに記載されています（nVidiaにも同様のドキュメントがあります）。明らかに、レジスタはAMDデバイス用に128ビット（4x32）であり、すべての最新のハイエンドデバイス用に16384個のレジスタがあるため、計算単位あたり256KBという顕著なものです。

Answer 2

私はあなたが__localメモリーを探していると思います。これが、32KBのローカルデータストレージが参照しているものです。私はあなたが私的なメモリ量を得るためにデバイスをポーリングできるとは思わない。

NULLの長い* cl_mem参照を渡して、メモリを割り当てることができます。私は、WIごとに静的な量のメモリを使用することが最善であると思います。各作業項目に長い[20​​0]が必要であると仮定すると、以下のコードを使用します。 LDSメモリを最大限に活用するために、同じ（または同様の）メモリ要件を持つグループに作業を分けることも良い考えです。

void __kernel(__local long* localMem, const int localMemPerItem 
     //more args... 
     ) 
{ 
    //host has 'passed' localMemPerItem*get_local_size() long values in as locamMem 
    //this work item has access to all of it, but can choose to restrict 
    //itself to only the portion it needs. 
    //work group size will be limited to CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem) 
    int startIndex=localMemPerItem*get_local_id(0); 
    //use localMem[startIndex+ ... ] 
} 

Answer 3

そのGCNアーキテクチャに基づいて、それはアナンドテックでイメージどおり64キロバイトであることから、の79xxシリーズのカードにファイルを登録する方法を大に答えるために：http://www.anandtech.com/print/5261

方法を見つけるためにどのようにあなたの質問に答えるためにあなたがカーネルでAMD APPプロファイラを実行するのを見ることができます。それは、カーネルがどのくらいのスペースを利用しているかをカーネル占有セクションで教えてくれます。