8

デバイスで何度も参照する必要がある float 配列があるため、格納するのに最適な場所は __ 定数 __ メモリ (この参照を使用) であると考えています。配列 (またはベクトル) は、初期化時に実行時に 1 回書き込む必要がありますが、複数の異なる関数によって何百万回も読み取られるため、関数呼び出しごとにカーネルに定数をコピーすることは悪い考えのように思えます。

const int n = 32;
__constant__ float dev_x[n]; //the array in question

struct struct_max : public thrust::unary_function<float,float> {
    float C;
    struct_max(float _C) : C(_C) {}
    __host__ __device__ float operator()(const float& x) const { return fmax(x,C);}
};
void foo(const thrust::host_vector<float> &, const float &);

int main() {
    thrust::host_vector<float> x(n);
    //magic happens populate x
    cudaMemcpyToSymbol(dev_x,x.data(),n*sizeof(float));

    foo(x,0.0);
    return(0);
}

void foo(const thrust::host_vector<float> &input_host_x, const float &x0) {
    thrust::device_vector<float> dev_sol(n);
    thrust::host_vector<float> host_sol(n);

    //this method works fine, but the memory transfer is unacceptable
    thrust::device_vector<float> input_dev_vec(n);
    input_dev_vec = input_host_x; //I want to avoid this
    thrust::transform(input_dev_vec.begin(),input_dev_vec.end(),dev_sol.begin(),struct_max(x0));
    host_sol = dev_sol; //this memory transfer for debugging

    //this method compiles fine, but crashes at runtime
    thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(dev_x);
    thrust::transform(dev_ptr,dev_ptr+n,dev_sol.begin(),struct_max(x0));
    host_sol = dev_sol; //this line crashes
}

グローバル推力::device_vector dev_x(n)を追加しようとしましたが、それも実行時にクラッシュし、__定数__メモリではなく__グローバル__メモリになります

スラスト ライブラリを破棄するだけで、これをすべて機能させることができますが、スラスト ライブラリをグローバルとデバイス定数メモリで使用する方法はありますか?

4

1 に答える 1

8

良い質問！__constant__配列を通常のデバイス ポインターのようにキャストすることはできません。

私はあなたの質問に答えます (下の行の後で)、しかし最初に: これは の不適切な使用であり__constant__、実際にはあなたが望むものではありません。CUDA の定数キャッシュは、ワープ内のスレッド間で均一にアクセスできるように最適化されています。これは、ワープ内のすべてのスレッドが同時に同じ場所にアクセスすることを意味します。ワープの各スレッドが異なる定数メモリ ロケーションにアクセスする場合、アクセスはシリアル化されます。したがって、連続するスレッドが連続するメモリ位置にアクセスするアクセス パターンは、均一なアクセスよりも 32 倍遅くなります。デバイスのメモリだけを使用する必要があります。データを 1 回書き込む必要があるが、何度も読み取る必要がある場合は、device_vector を使用します。一度初期化してから、何度も読み取ります。

あなたが要求したことを行うには、 athrust::counting_iteratorを入力として使用して、配列thrust::transformにインデックスの範囲を生成できます。__constant__次に、ファンクターは値オペランドではなくインデックスオペランドoperator()を取り、定数メモリへのルックアップを行います。intfloat

(これは、ファンクターが__device__コードのみであることを意味することに注意してください。移植性が必要な場合は、演算子を簡単にオーバーロードして float を取得し、ホスト データで別の方法で呼び出すことができます。)

例を変更してデータを初期化し、結果を出力して正しいことを確認しました。

#include <stdio.h>
#include <stdlib.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/iterator/counting_iterator.h>

const int n = 32;
__constant__ float dev_x[n]; //the array in question

struct struct_max : public thrust::unary_function<float,float> {
    float C;
    struct_max(float _C) : C(_C) {}

    // only works as a device function
    __device__ float operator()(const int& i) const { 
        // use index into constant array
        return fmax(dev_x[i],C); 
    }
};

void foo(const thrust::host_vector<float> &input_host_x, const float &x0) {
    thrust::device_vector<float> dev_sol(n);
    thrust::host_vector<float> host_sol(n);

    thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(dev_x);
    thrust::transform(thrust::make_counting_iterator(0),
                      thrust::make_counting_iterator(n),
                      dev_sol.begin(),
                      struct_max(x0));
    host_sol = dev_sol; //this line crashes

    for (int i = 0; i < n; i++)
        printf("%f\n", host_sol[i]);
}

int main() {
    thrust::host_vector<float> x(n);

    //magic happens populate x
    for (int i = 0; i < n; i++) x[i] = rand() / (float)RAND_MAX;

    cudaMemcpyToSymbol(dev_x,x.data(),n*sizeof(float));

    foo(x, 0.5);
    return(0);
}
于 2013-06-13T03:19:16.060 に答える