7

さて、私はかなりデリケートな質問があります:)

私が持っているものから始めましょう：

  1. Data、データの大規模な配列、GPU にコピー
  2. Program、CPU (ホスト) によって生成され、その配列内のすべてのデータについて評価する必要があります
  3. プログラムは非常に頻繁に変更され、CUDA 文字列、PTX 文字列、またはその他 (?) として生成される可能性があり、変更のたびに再評価する必要があります。

私が欲しいもの：基本的には、これをできるだけ効果的（高速）にしたいだけです。CUDA から PTX へのコンパイルを回避します。ソリューションは完全にデバイス固有であってもかまいません。ここでは大きな互換性は必要ありません:)

私が知っていること:ファイルに保存された PTX コードからカーネルをロードおよび作成できる関数cuLoadModuleを既に知っています。しかし、最初にファイルに保存せずに、カーネルを直接作成する方法が他にあるに違いないと思います。あるいは、バイトコードとして保存することは可能でしょうか?

私の質問: どうやってそれをしますか? 例を投稿するか、同様のトピックの Web サイトへのリンクを投稿していただけますか? タイ

編集: OK、PTX カーネルはPTX文字列 (文字配列) から直接実行できます。とにかく、私はまだ疑問に思っていますが、これに対するより良い/より速い解決策はありますか? 文字列から一部の PTX バイトコードへの変換がまだありますが、これはおそらく回避する必要があります。また、PTX からデバイス固有の Cuda バイナリを作成する巧妙な方法が存在するのではないかと考えています。これにより、JIT コンパイラのラグが解消されます (小さいですが、実行するカーネルが膨大な数に達すると、合計される可能性があります) :)

4

1 に答える 1

6

彼のコメントで、Roger Dahl は次の投稿をリンクしています。

PTX プログラムを CUDA ドライバーに直接渡す

cuModuleLoadここでは、との 2 つの関数の使用cuModuleLoadDataExが取り上げられています。前者は、ファイルから PTX コードをロードし、それをnvccコンパイラ ドライバに渡すために使用されます。後者は I/O を回避し、PTX コードを C 文字列としてドライバーに渡すことができます。いずれの場合も、CUDA カーネルのコンパイルの結果として (C 文字列にロードまたはコピーして貼り付ける)、または手書きのソースとして、PTX コードを自由に使用できるようにしておく必要があります。

しかし、CUDA カーネルから開始してオンザフライで PTX コードを作成する必要がある場合はどうなるでしょうか? CUDA Expression templatesのアプローチに従って、CUDA カーネルを含む文字列を次のように定義できます。

ss << "extern \"C\" __global__ void kernel( ";
ss << def_line.str() << ", unsigned int vector_size, unsigned int number_of_used_threads ) { \n";
ss << "\tint idx = blockDim.x * blockIdx.x + threadIdx.x; \n";
ss << "\tfor(unsigned int i = 0; i < ";
ss << "(vector_size + number_of_used_threads - 1) / number_of_used_threads; ++i) {\n";
ss << "\t\tif(idx < vector_size) { \n";
ss << "\t\t\t" << eval_line.str() << "\n";
ss << "\t\t\tidx += number_of_used_threads;\n";
ss << "\t\t}\n";
ss << "\t}\n";
ss << "}\n\n\n\n";

次に、システムコールを使用してコンパイルします

int nvcc_exit_status = system(
         (std::string(NVCC) + " -ptx " + NVCC_FLAGS + " " + kernel_filename 
              + " -o " + kernel_comp_filename).c_str()
    );

    if (nvcc_exit_status) {
            std::cerr << "ERROR: nvcc exits with status code: " << nvcc_exit_status << std::endl;
            exit(1);
    }

最後に と を使用cuModuleLoadしcuModuleGetFunctionてファイルから PTX コードをロードし、それを次のようにコンパイラ ドライバに渡します。

    result = cuModuleLoad(&cuModule, kernel_comp_filename.c_str());
    assert(result == CUDA_SUCCESS);
    result =  cuModuleGetFunction(&cuFunction, cuModule, "kernel");
    assert(result == CUDA_SUCCESS);

もちろん、式テンプレートはこの問題とは何の関係もありません。この回答で報告しているアイデアのソースを引用しているだけです。

于 2013-11-07T22:08:41.403 に答える