CUDA 추력 펑션 GMEM 액세스 : ctor 데이터 복사 대 ctor dev ptr arg

Question

추력 펑터가 벡터화 할 수없는 불균일하게 액세스되는 읽기 전용 상태에 액세스 할 수있게하는 두 가지 방법이 있습니다. 불행히도 커널 실행 시간에는 100 배의 차이가 있습니다. 왜 두 가지 전략에 어떤 차이가 있습니까?

더 일반적으로 : 이러한 종류의 전역에 액세스 할 수있는 추력 펑터를 제공하는 표준 방법이 있습니까?

내 첫 번째 방법은 글로벌 데이터 사본을 펑터에 넣는 것입니다.

// assuming barData is a float[] 
foo<N>(barData); 

내 두 번째 방법은 자신이 추력을 사용하여 장치에 업로드를 수행하는 것입니다

추력 ::의 for_each를 사용하여 호출

// functor containing a copy of array dependency 
template<size_t BARSIZE> 
struct foo1_func 
{ 
    __align__(16) float bar[BARSIZE]; 
    foo1_func(float _bar[BARSIZE]) { memcpy(bar,_bar,BARSIZE*sizeof(float)); } 
    __host__ __device__ operator()(float &t) { t = do_something(t, bar); } 
} 

... : 추력 기계 장치에 업로드 및 캐싱을 수행하기 위해 나타납니다 :: 업로드 한 데이터의 장치 메모리 포인터를 복사하여 내 펑터에 전달하십시오. 이 방법은 훨씬 느린 것으로 나타납니다 기꺼이 받아 정식 또는 고유 펑 패턴을 보여 소스에

추력 ::의 for_each를 사용하여 호출

// functor containing device pointers to array in GMEM 
struct foo2_func 
{ 
    float *bar; 
    foo2_func(float* _bar) { bar = bar; } 
    __host__ __device__ operator()(float &t) { t = do_something(t, bar); } 
} 

...

// assuming d_bar is a thrust::device_vector 
foo(thrust::raw_pointer_cast(d_bar.data())); 

링크.

Answer 1

첫 번째 방법은 커널 함수 매개 변수로 foo1_func 구조체를 전달하여 전체 배열 bar을 GPU 레지스터에 넣으려고하는 것입니다. bar의 크기가 거부자에 넣어 수있을 정도로 작은 경우

__global__ void kernel_generated_by_thrust(struct foo_func f, ...) { 
    float x = f.bar[3]; 
    ... 
} 

는 bar에 랜덤 액세스가 실제로 등록하는 랜덤 액세스이다.

하지만 두 번째 방법은 구조체를 통해 전역 메모리 포인터 만 전달했습니다. 따라서 bar에 대한 임의 액세스는 전역 메모리에 대한 임의 액세스입니다.

두 번째 방법이 훨씬 느립니다.

두 가지 방법 모두 사용 사례가 있습니다. 당신은 달성하고자하는 것에 따라 어느 하나를 선택할 수 있습니다. bar의 크기와 bar을 캐싱하는데 얼마나 많은 등록비를 쓸지 선택할 수 있습니다.