CUDAでカスタムメモリー管理と推力を混ぜる

Question

私のプロジェクトでは、アプリケーションが "ウォームアップ"したら、cudaMallocへの不必要な呼び出しを避けるためにカスタムメモリーアロケーターを実装しました。さらに、配列の基本的な配列の埋め込みや算術演算などにカスタムカーネルを使用し、Thrustを使ってコードを単純化し、これらのカーネルを取り除きたいと考えています。デバイス上のすべての配列が作成され、未加工のポインタからアクセスされています（今のところ）device_vectorとThrustのメソッドをこれらのオブジェクトに使用したいのですが、自分自身が生ポインタとdevice_ptr<>をいつも変換していますコード。

私のむしろ漠然とした質問：カスタムメモリ管理の使用方法を整理するにはどうすればよいでしょうか？Thrustアレイメソッドとカスタムカーネルの呼び出しを最も読みやすい方法で呼び出す方法はありますか？

Answer 1

すべての標準的なC++コンテナと同様に、thrust::device_vectorは独自に"allocator"を指定してストレージを割り当てる方法をカスタマイズできます。デフォルトでは、thrust::device_vectorのアロケータはthrust::device_malloc_allocatorです。スラストのバックエンドシステムがCUDAの場合、cudaMalloc（cudaFree）でストレージを割り当て（割り当てを解除します）行います。

device_vectorは、OPの場合のように、プログラムの初期化時に実行される単一の大きな割り当て内でサブストレージを割り当てる方法など、メモリを割り当てる方法をカスタマイズすることが望ましいことがあります。これにより、基本的な割り当て方式（多くの場合、cudaMalloc）への個々の呼び出しによって発生するオーバーヘッドを回避できます。

device_vectorカスタムアロケータを提供する簡単な方法は、device_malloc_allocatorから継承することです。原則としてアロケータ全体を最初から作成することができますが、継承アプローチでは、allocateとdeallocateメンバ関数のみを提供する必要があります。カスタムアロケータが定義されると、それを第2のテンプレートパラメータとしてdevice_vectorに提供することができる。私たちはから聞くことに注意して、この例では

$ nvcc my_allocator_test.cu -arch=sm_20 -run 
my_allocator::allocate(): Hello, world! 
my_allocator::deallocate(): Hello, world! 

：

このコード例では、割り当てと解放時にメッセージを出力するカスタムアロケータを提供する方法を示しています。ここでは

#include <thrust/device_malloc_allocator.h> 
#include <thrust/device_vector.h> 
#include <iostream> 

template<typename T> 
    struct my_allocator : thrust::device_malloc_allocator<T> 
{ 
    // shorthand for the name of the base class 
    typedef thrust::device_malloc_allocator<T> super_t; 

    // get access to some of the base class's typedefs 

    // note that because we inherited from device_malloc_allocator, 
    // pointer is actually thrust::device_ptr<T> 
    typedef typename super_t::pointer pointer; 

    typedef typename super_t::size_type size_type; 

    // customize allocate 
    pointer allocate(size_type n) 
    { 
    std::cout << "my_allocator::allocate(): Hello, world!" << std::endl; 

    // defer to the base class to allocate storage for n elements of type T 
    // in practice, you'd do something more interesting here 
    return super_t::allocate(n); 
    } 

    // customize deallocate 
    void deallocate(pointer p, size_type n) 
    { 
    std::cout << "my_allocator::deallocate(): Hello, world!" << std::endl; 

    // defer to the base class to deallocate n elements of type T at address p 
    // in practice, you'd do something more interesting here 
    super_t::deallocate(p,n); 
    } 
}; 

int main() 
{ 
    // create a device_vector which uses my_allocator 
    thrust::device_vector<int, my_allocator<int> > vec; 

    // create 10 ints 
    vec.resize(10, 13); 

    return 0; 
} 

を出力ですmy_allocator::allocate()vec.resize(10,13)に一度。 my_allocator::deallocate()は、vecがその要素を破棄するときに範囲外になると一度呼び出されます。