0

私のunary_op.operatorでは、一時配列を作成する必要があります。
私cudaMallocは行く方法だと思います。
しかし、それは効率的なパフォーマンスですか、それともより良い設計がありますか?

struct my_unary_op
{
    __host__ __device__ int operator()(const int& index) const
    {
        int* array;
        cudaMalloc((void**)&array, 10*sizeof(int));

        for(int i = 0; i < 10; i++)
            array[i] = index;

        int sum=0;
        for(int i=0; i < 10 ; i++)
            sum += array[i];

        return sum;
    };

};
int main()
{
    thrust::counting_iterator<int> first(0);
    thrust::counting_iterator<int> last = first+100;

    my_unary_op unary_op = my_unary_op();

    thrust::plus<int> binary_op;

    int init = 0;
    int sum = thrust::transform_reduce(first, last, unary_op, init, binary_op);

    return 0;
};
4

1 に答える 1

2

ホストのみの関数であるため、関数でcudaMalloc()コンパイルすることはできません。__device__ただし、通常のmalloc()or new(コンピューティング機能 >= 2.0 のデバイス上) を使用できますが、デバイス上で実行する場合、これらはあまり効率的ではありません。これには 2 つの理由があります。1 つ目は、同時に実行されているスレッドが、メモリ割り当ての呼び出し中にシリアル化されることです。2 つ目は、メモリのロード命令とストア命令がワープ内の 32 のスレッドによって実行されるときにそれらが隣接しないように配置されるチャンクで呼び出しがグローバル メモリを割り当てるため、適切に結合されたメモリを取得できないことです。アクセスします。

__device__関数で固定サイズの C スタイルの配列 (つまり)を使用することで、これらの問題の両方に対処できますint array[10];。小さい固定サイズの配列は、非常に高速なアクセスのためにレジスターファイルに格納されるように、コンパイラーによって最適化される場合があります。コンパイラがそれらをグローバル メモリに格納する場合、ローカル メモリが使用されます。ローカル メモリはグローバル メモリに格納されますが、ワープ内の 32 のスレッドがロードまたはストア命令を実行するときに、各スレッドがメモリ内の隣接する場所にアクセスし、トランザクションを完全に結合できるようにインターリーブされます。

実行時に C 配列のサイズがわからない場合は、配列に最大サイズを割り当て、その一部を未使用のままにします。

固定サイズの配列で使用されるメモリの総量は、カーネルによって起動されるスレッドの総数ではなく、GPU で同時に処理されるスレッドの総数に依存すると思います。この回答で @mharris は、可能な最大同時スレッド数 (GTX580 の場合は 24,576) を計算する方法を示しています。したがって、固定サイズの配列が 16 個の 32 ビット値である場合、配列が使用できるメモリの最大量は 1536KiB になります。

幅広い配列サイズが必要な場合は、テンプレートを使用して、さまざまなサイズのカーネルをコンパイルできます。次に、実行時に、必要なサイズに対応できるものを選択します。ただし、必要な最大数を単純に割り当てると、メモリ使用量が起動できるスレッド数の制限要因にならない可能性があります。

于 2012-12-30T22:10:32.173 に答える