2

Nvidiaのスラストライブラリについて知りました。試してみるだけで、一連のベクトルを正規化することになっている小さな例を書きました。

#include <cstdio>

#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

struct normalize_functor: public thrust::unary_function<double4, double4>
{
    __device__ __host__ double4 operator()(double4 v)
    {
        double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
        v.x /= len;
        v.y /= len;
        v.z /= len;
        printf("%f %f %f\n", v.x, v.y, v.z);
    }
};

int main()
{
    thrust::host_vector<double4> v(2);
    v[0].x = 1; v[0].y = 2; v[0].z = 3;
    v[1].x = 4; v[1].y = 5; v[1].z = 6;

    thrust::device_vector<double4> v_d = v; 
    thrust::for_each(v_d.begin(), v_d.end(), normalize_functor());

    // This doesn't seem to copy back
    v = v_d;

    // Neither this does..
    thrust::host_vector<double4> result = v_d;

    for(int i=0; i<v.size(); i++)
        printf("[ %f %f %f ]\n", result[i].x, result[i].y, result[i].z);

    return 0;
}

上記の例は機能しているようですが、データをコピーして戻すことができません。単純な割り当てでcudaMemcpyが呼び出されると思いました。ホストからデバイスにデータをコピーするように機能しますが、元に戻すことはできません???

第二に、これを正しい方法で行うかどうかはわかりません。for_eachのドキュメントには次のように書かれています。

for_eachは、関数オブジェクトfを[first、last）;の範囲の各要素に適用します。fの戻り値がある場合は、無視されます。

ただし、unary_function構造体テンプレートは2つのテンプレート引数（1つは戻り値用）を想定しており、operator（）も値を返すように強制します。これにより、コンパイル時に警告が発生します。戻り値のない単項ファンクターをどのように書くべきかわかりません。

次はデータの配置です。double4を選択したのは、2つのフェッチ命令ld.v2.f64とld.f64IIRCが生成されるためです。ただし、thrustが内部でデータをフェッチする方法（およびcudaスレッド/ブロックの数）が作成されるのではないかと思います。代わりに4つのベクトルの構造体を選択すると、合体した方法でデータをフェッチできるようになります。

最後に、推力はタプルを提供します。タプルの配列はどうですか？この場合、データはどのように配置されますか。

例を調べましたが、一連のベクトルに対してどのデータ構造を選択するかを説明する例は見つかりませんでした（dot_products_with_zip.cuの例では、「構造の配列」ではなく「配列の構造」について説明していますが、この例では構造は使用されていません。

アップデート

上記のコードを修正し、より大きな例を実行しようとしました。今回は10kのベクトルを正規化します。

#include <cstdio>

#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

struct normalize_functor
{
    __device__ __host__ void operator()(double4& v)
    {
        double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
        v.x /= len;
        v.y /= len;
        v.z /= len;
    }
};

int main()
{
    int n = 10000;
    thrust::host_vector<double4> v(n);
    for(int i=0; i<n; i++) {
        v[i].x = rand();
        v[i].y = rand();
        v[i].z = rand();
    }

    thrust::device_vector<double4> v_d = v;

    thrust::for_each(v_d.begin(), v_d.end(), normalize_functor());

    v = v_d;

    return 0;
}

Computeprofを使用したプロファイリングでは、占有率が低く、非合体のメモリアクセスが示されます。

Kernel Occupancy Analysis

Kernel details : Grid size: 23 x 1 x 1, Block size: 448 x 1 x 1
Register Ratio      = 0.984375  ( 32256 / 32768 ) [24 registers per thread] 
Shared Memory Ratio     = 0 ( 0 / 49152 ) [0 bytes per Block] 
Active Blocks per SM        = 3 / 8
Active threads per SM       = 1344 / 1536
Potential Occupancy     = 0.875  ( 42 / 48 )
Max achieved occupancy  = 0.583333  (on 9 SMs)
Min achieved occupancy  = 0.291667  (on 5 SMs)
Occupancy limiting factor   = Block-Size

Memory Throughput Analysis for kernel launch_closure_by_value on device GeForce GTX 470

Kernel requested global memory read throughput(GB/s): 29.21
Kernel requested global memory write throughput(GB/s): 17.52
Kernel requested global memory throughput(GB/s): 46.73
L1 cache read throughput(GB/s): 100.40
L1 cache global hit ratio (%): 48.15
Texture cache memory throughput(GB/s): 0.00
Texture cache hit rate(%): 0.00
L2 cache texture memory read throughput(GB/s): 0.00
L2 cache global memory read throughput(GB/s): 42.44
L2 cache global memory write throughput(GB/s): 46.73
L2 cache global memory throughput(GB/s): 89.17
L2 cache read hit ratio(%): 88.86
L2 cache write hit ratio(%): 3.09
Local memory bus traffic(%): 0.00
Global memory excess load(%): 31.18
Global memory excess store(%): 62.50
Achieved global memory read throughput(GB/s): 4.73
Achieved global memory write throughput(GB/s): 45.29
Achieved global memory throughput(GB/s): 50.01
Peak global memory throughput(GB/s): 133.92

どうすればこれを最適化できるのだろうか？

4

2 に答える 2

4

シーケンスをインプレースで変更する場合for_eachは、ファンクターで参照して引数を取る必要があります。

struct normalize_functor
{
    __device__ __host__ void operator()(double4& ref)
    {
        double v = ref;
        double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
        v.x /= len;
        v.y /= len;
        v.z /= len;
        printf("%f %f %f\n", v.x, v.y, v.z);
        ref = v;
    }
};

normalize_functorまたは、transformアルゴリズムでの定義を使用してv_d、送信元と宛先の両方の範囲として指定することもできます。

thrust::transform(v_d.begin(), v_d.end(), v_d.begin(), normalize_functor());

私の個人的な好みはtransformこの状況で使用することですが、パフォーマンスはどちらの場合も同じである必要があります。

于 2011-04-21T02:59:20.277 に答える
1

最適化の問題に関しては、Thrustで実行できることはあまりありません。これは、実際にはライブラリの意図ではありません。Thrustの作者の1人であり、すでにこのスレッドに投稿しているNathan Bellについて話すことを望まずに、GPUのさまざまなデータ並列アルゴリズムを、多くのことを書かなくても、シンプルで直感的な方法で利用できるようにすることを目的としています。 、もしあれば、CUDAコード。そして、私の意見では、それは見事に成功しています。多くのスラストカーネルのカーネルパフォーマンスは最先端に近いものですが、一般的なテンプレートコードでは簡単に実行できない特定の場合に実行できる最適化が常にあります。これは、Thrustが提供する使いやすさと柔軟性のために支払う価格の一部です。

そうは言っても、演算子関数を試してみると、改善される可能性のある調整がいくつかあると思います。私は通常、次のようなものを書きます。

struct normalize_functor
{
    __device__ __host__ void operator()(double4& v)
    {
        double4 nv = v;
        double len = sqrt(nv.x*nv.x + nv.y*nv.y + nv.z*nv.z);
        nv.x /= len;
        nv.y /= len;
        nv.z /= len;
        (void)nv.h;
        v = nv;
    };
};

これで、元のファイルほどきれいではありませんが、コンパイラーがベクトル化されたロードおよびストア命令を発行するようにする必要があります。過去に、コンパイラーがベクトル型の未使用メンバーのロードとストアを最適化し、その結果、PTXジェネレーターがスカラーのロードとストアを発行し、結果として合体を中断するケースを見てきました。明確なfloat4のロードとストアを行い、構造の各要素が使用されていることを確認することで、少なくとも2.xおよび3.xnvccリリースに存在していたこの不要な「最適化」を回避できます。これが4.0コンパイラでも当てはまるかどうかはわかりません。

于 2011-04-21T09:57:35.367 に答える