ギドラ用のNvidia Kepler（CUDAバイナリ、言語バージョンsm_30）の仕様を作成しています

Ghidraの仕様の多くは、通常のプロセッサ言語用に既に作成されていますが、グラフィカルな仕様はありません。シェーダーから継承された他のものを含め、パラメーターが渡される述語、定数など、固有の特性があるため、理解できます。さらに、コードを格納するために使用されるフォーマットは多くの場合独自仕様であり、自分で逆にする必要があります。



この記事では、2つの例の内容を理解します。



最初のプログラムは、最も単純なaxpy（GPGPUのHello Worldのアナログ）です。2番目は、GPUでの条件とジャンプの実装を理解するのに役立ちます。そこではすべてが異なります。



すべてのNvidia言語はリトルエンディアンエンコーディングを使用しているため、8進数の逆順で16進エディターから一部のノートブック（たとえば、メモ帳++）にバイトをすぐにコピーします（ここでは、命令の長さは一定です）。次に、プログラマー計算機（Microsoftの計算機が適しています）を使用して、それをバイナリコードに変換します。次に、一致を探し、命令マスクを作成してから、オペランドを作成します。マスクをデコードして検索するために、16進エディターとcuobjdump逆アセンブラーが使用されました。AMDGPUのように、アセンブラーが必要になる場合があります（逆アセンブラーはそこにはありませんが、これは別の記事のトピックです）。それはこのように機能します：電卓のすべての疑わしいビットを順番に反転させてから、バイトの新しい16進値を取得し、nvccまたはアセンブラを介してコンパイルされたバイナリに置き換えます（存在する場合）。次に、cuobjdumpを使用して確認します。



私はソースをフォーマット（主にCで、マシンのGPUコードとのより緊密な接続のためのOOPなしで）で展開し、一度にdisasm + bytesを無効にします。



それをaxpy.cuにコピーして、cmdを介してコンパイルします。nvcc axpy.cu --cubin --gpu-architecture sm_30

同じ場所にaxpy.cubinという名前の結果のELFファイルを分解します。cuobjdump axpy.cubin -sass



例1：

__global__ void axpy(float param_1, float* param_2, float* param_3) {
unsigned int uVar1 = threadIdx.x;
param_2[uVar1] = param_1 * param_3[uVar1];
}

例2：

__global__ void predicates(float* param_1, float* param_2) {
    unsigned int uVar1 = threadIdx.x + blockIdx.x * blockDim.x;
    if ((uVar1 > 5) & (uVar1 < 10)) param_1[uVar1] = uVar1;
    else param_2[uVar1] = uVar1;
}

コンパイラが最適化するものが何もないように、テストが元々マシンコードに合わせて調整されていたと推測するのは簡単です。それ以外の場合は、手動で最適化をキャンセルする必要があります。複雑な例では、これがまったくできない場合があるため、そのような場合は、逆コンパイラとフロントエンドに依存する必要があります。



一般に、ルールは次のとおりです。フロントエンドをテストするために、最初の適切な（エラーの再現）最初の（可能な最適化を最小限にした）単純な例を取り上げます。残りの部分については、逆コンパイルされたコードは既に最適化されています（または、リファクタリングによってそれをどうにか修正するだけです）。しかし今のところ、主要なタスクは、少なくともマシンコードと同じことを行う正しいコードです。これが「ソフトウェアモデリング」です。 「ソフトウェアモデリング」自体は、リファクタリング、CからC ++への変換、クラスの復元、さらにはテンプレートの識別などを意味するものではありません。



現在、ニーモニック、オペランド、修飾子のパターンを探しています。



これを行うには、疑わしい命令（または、そのように呼び出す方が便利な場合は文字列）間でビットを（バイナリ表現で）比較します。また、「バイナリ/ sass /マシンコードを理解するのに役立つ」など、stackoverflowに関する質問に他のユーザーが投稿したものを使用したり、チュートリアル（中国語を含む）やその他のリソースを使用したりすることもできます。したがって、メイン操作番号はビット58〜63に格納されますが、0〜2（無視する場合、3〜空の命令では4ビット、仕様では「UNK」とマークされています）。



レジスターと定数については、逆アセンブラーを試して、スポイラーの下に配置されているような、ダンプ出力に影響を与えるすべてのビットを見つけることができます。私が見つけたすべてのフィールドは、Githubの仕様、ファイルCUDA.slaspec、セクショントークンにあります。



次に、レジスタのアドレスを考え出す必要があります。これらもGithubにあります。これは、マイクロレベルでは、Sleighはレジスタを「register_space」タイプのスペース内のグローバル変数と見なしますが、それらのスペースは「推論可能」としてマークされていません（ほとんどの場合はそうではありません）。逆コンパイラでは、ローカル変数（ほとんどの場合「Var」インターフィックスが付いていますが、「ローカル」プレフィックスもある場合があります）、またはパラメータ（「 param_ "）。 SPは便利ではありませんでした。逆コンパイラが機能していることを確認するために、ほとんどが正式に必要です。エミュレートするには、PC（x86から​​のIPのようなもの）が必要です。



次に、フラグのような述語レジスタがありますが、オーバーフロー、ゼロでない（等しくない）など、予測された目的よりも「一般的な目的」があります。

次に、ISCADD .CCとIMAD.HIの一連の命令をモデル化するためのロックレジスタ。私の実装ではそれらの最初のものは、合計の一部を上位4バイトに転送することを避けるために、それ自体と2番目のものについてカウントを実行します。これは逆コンパイルを台無しにするでしょう。ただし、IMAD.HI操作が完了するまで、次のレジスタをロックする必要があります。類似したもの、つまり公式ドキュメントと逆コンパイラの予想される出力の不一致は、同じGhidraのSPUモジュールにすでにありました。



次に、これまでcpoolを介して実装された特別なレジスタがあります。将来的には、ある種の「推論可能な」スペースに対して、デフォルトで定義されている文字に置き換える予定です。これらは同じthreadIdx、blockIdxです。



次に、変数をdest、par0、par1、par2、resフィールドにバインドします。次に、サブテーブルがあり、その後に-それが何であるか-メイン（ルート）テーブルとメインの説明があります。



ここでは、「ニーモニックオペランド」の形式に厳密に従う必要がありますが、修飾子は除外されますが、修飾子はニーモニックまたはオペランドのあるセクションに添付する必要があります。他の形式は許可されていません。同じHexagon DSP asmでもこの構文に適合させる必要がありますが、それほど難しくはありません。



最終段階では、命令の実装をPcodeマイクロプログラミング言語で記述します。最初の例で注意したいのは、ISCADD .CCとIMAD.HI命令だけです。これらの最初の命令は、レジスターへのポインターを受け取り、それらを4ではなく8バイトへのポインターとして逆参照します。金額の一部の送金に関するNvidiaのドキュメントに記載されている内容にもかかわらず、その動作。



2番目の例では、「プレディケーションの簡略化」という表記の反対側の逆コンパイラ設定を確認することをお勧めします。重要なのは、述語は異なる命令に対してまったく同じ条件であり、本質的には、よく知られている「SIMD」またはその次の同等物に過ぎないということです。それら。述語ビットが設定されている場合、命令はさらに続けて実行されます。



また、デコンパイラー、エミュレーター、およびその他のアナライザーもあるので、プロトタイプ（ニーモニックオペランド）だけでなく、各命令の実装をすぐに作成する習慣をつける必要があります。

しかし、一般に、Pcodeで実装を作成することは、バイトデコーダーの文法を作成するよりもさらに簡単な作業です。非常に便利な中間言語、単一のミドルランド（オプティマイザー）、2つのバックエンド（ほとんどの場合C、代替としてJava / C＃、後者のように、つまり、 K. gotoは時々表示されますが、breakというラベルは付いていません。

次の記事では、DXBC、SPIR-Vなどのマネージ言語のフロントエンドもある可能性があり、Java / C＃バックエンドを使用します。しかし、これまでのところ計画にはマシンコードのみが含まれています。バイトコードには特別なアプローチが必要です。Ghidra Help



Project ：Pcode Sleigh