2012-12-30 4 views
0

unary_op.operator에 임시 배열을 만들어야합니다.
나는 cudaMalloc 가야 할 길입니다.
하지만 성능이 효율적입니까 아니면 더 나은 디자인이 있습니까? 이 호스트 전용 기능 때문에추력 : transform_reduce : 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; 
}; 

답변

2

당신은 __device__ 기능에 cudaMalloc()를 컴파일 할 수 없습니다. 그러나 보통 malloc() 또는 new (컴퓨팅 기능> = 2.0 인 기기에서)을 사용할 수 있지만 기기에서 실행하는 경우에는 그다지 효율적이지 않습니다. 이것에는 두 가지 이유가 있습니다. 첫 번째는 동시에 실행중인 스레드가 메모리 할당 호출 중에 직렬화된다는 것입니다. 두 번째는 호출이 전역 메모리를 청크로 할당하고, 메모리로드 및 저장 명령어가 워프의 32 개 스레드에 의해 실행될 때 인접하지 않으므로 적절하게 병합 된 메모리를 얻지 못하도록 배열됩니다. 액세스.

__device__ 함수 (예 : int array[10];)에서 고정 크기 C 스타일 배열을 사용하여이 두 가지 문제를 해결할 수 있습니다. 작고 고정 된 크기의 배열은 매우 빠른 액세스를 위해 레지스터 파일에 저장되도록 컴파일러에서 최적화 할 수 있습니다. 컴파일러가이를 전역 메모리에 저장하면 로컬 메모리를 사용합니다. 로컬 메모리는 전역 메모리에 저장되지만, 워프의 32 개 스레드가로드 또는 저장 명령어를 실행할 때 각 스레드가 메모리의 인접한 위치에 액세스하여 트랜잭션을 완전하게 병합 할 수 있도록 인터리브됩니다.

런타임시 C 배열의 크기를 알지 못하는 경우 배열에 최대 크기를 할당하고 일부는 사용하지 않습니다.

고정 크기 배열에 사용되는 총 메모리 양은 커널이 실행 한 총 스레드 수가 아니라 GPU에서 동시에 처리되는 총 스레드 수에 따라 달라집니다. In this answer @mharris는 동시 스레드의 가능한 최대 수를 계산하는 방법을 보여줍니다. 이는 GTX580의 경우 24,576입니다. 따라서 고정 크기 배열이 16 32 비트 값인 경우 배열에서 사용하는 최대 메모리 양은 1536KiB가됩니다.

다양한 배열 크기가 필요한 경우 템플릿을 사용하여 다양한 크기의 커널을 컴파일 할 수 있습니다. 그런 다음 런타임에 필요한 크기를 수용 할 수있는 것을 선택하십시오. 그러나 필요한 경우 최대 값을 단순히 할당하면 메모리 사용량은 시작할 수있는 스레드 수를 제한하는 요소가되지 않습니다.