스레드

Question

이 질문은 몇 주 전에 내 옆에 게시 기존 질문과 관련이 사이의 최소 찾기 : TERCOM algorithm - Changing from single thread to multiple threads in CUDA

간단히 커널의 스레드 각각 MAD 값을 계산하고 내가 알고 싶습니다 설명을 최소치와 그 위치.

나는이

__global__ void kernel (int m, int n, int h, int N, int *f, float heading, float *measurements, int *global_min) 
{ 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    int idy = blockIdx.y * blockDim.y + threadIdx.y; 

    float MAD=0; 
    float pos[2]; 
    float theta=heading*(PI/180); 
    float fval = 0; 

    // Calculate how much to move in x and y direction 
    float offset_x = h*cos(theta); 
    float offset_y = -h*sin(theta); 

    //Calculate Mean Absolute Difference 
    if(idx < n && idy < m) 
    { 
     for(float g=0; g<N; g++) 
     { 
      float fval = tex2D (tex, idx+(g-2)*offset_x+0.5f, idy+(g-2)*offset_y+0.5f); 
      MAD += abs(measurements[(int)g]-fval); 
     } 
    } 
    cuPrintf("%.2f \n",MAD); 

    atomicMin(global_min, MAD); 
    pos[0]=idx; 
    pos[1]=idy; 

    f[0]=*global_min; 
    f[1]=pos[0]; 
    f[2]=pos[1]; 
} 

같은 atomicMin을 사용하려고했습니다 그리고 그것은 올바른 결과를 생성하지만, atomicMin 최소의 위치를 ​​찾을 수 없습니다.

는 또한 추력 라이브러리

__global__ void kernel (int m, int n, int h, int N, int *f, float heading, float *measurements, int *global_min, float *dev_MAD) 
{ 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    int idy = blockIdx.y * blockDim.y + threadIdx.y; 

    float theta=heading*(PI/180); 
    float fval = 0; 

    // Calculate how much to move in x and y direction 
    float offset_x = h*cos(theta); 
    float offset_y = -h*sin(theta); 

    //Calculate Mean Absolute Difference 
    if(idx < n && idy < m) 
    { 
     for(float g=0; g<N; g++) 
     { 
      float fval = tex2D (tex, idx+(g-2)*offset_x+0.5f, idy+(g-2)*offset_y+0.5f); 
      *dev_MAD += abs(measurements[(int)g]-fval); 
     } 
    } 
    cuPrintf("%.2f \n",MAD); 
} 

을 사용하려고 그리고이

kernel <<< dimGrid,dimBlock >>> (m, n, h, N, dev_results, heading, dev_measurements, global_min, dev_MAD); 

thrust::device_ptr<float> dev_ptr(dev_MAD); 
thrust::device_ptr<float> min_pos = thrust::min_element(dev_ptr, dev_ptr + n*m); 
int abs_pos = min_pos - dev_ptr; 
float min_val=min_pos[0]; 

cudaMemcpy(&min_val, dev_MAD+abs_pos, sizeof(float), cudaMemcpyDeviceToHost); 

// Print out the result 
printf("Min=%.2f pos=%d\n",min_val,abs_pos); 

같은 커널하지만이 프로그램을 프린트 아웃 호출 : 최소 = -207521258711807190000000000000000000000.00 POS = 0

I 많은 축소 예제를 살펴 보았지만 모든 사람들이 배열에 저장된 값을 갖고 있고 각 스레드가 아닌 것처럼 보입니다.

질문에 대한 그래서 :

1. 는 atomicMin 기능의 위치를 ​​반환 할 수 있습니까?
2. 추력 라이브러리의 문제를 해결하는 방법에 대한 힌트를 누군가 줄 수 있습니까?

Answer 1

스러스트 코드의 경우 dev_MAD [0]에 쓰고 있지만 아직 전체 배열에 작성한 것처럼 계산됩니다.

IIUC, 최소값과 해당 위치를 찾으려고하면 각 스레드에서 변수로 값을 가지지 만 메모리에 저장되지는 ​​않습니다.

내가 생각할 수있는 몇 가지 쉬운 방법이 있지만 둘 다 값을 메모리에 저장하고 두 번째 패스에서 최소값/위치를 계산하는 것이 포함됩니다.

우선, 이미 시도한 것처럼 Thrust의 min_element을 사용할 수 있지만 커널의 device_vector에 값을 저장 한 다음 thrust :: min_element를 독립적으로 호출합니다.

둘째, 먼저 스레드 블록 내의 최소/위치를 계산하여 메모리 공간과 대역폭을 절약 할 수 있습니다 (나중에 thrust :: min_element를 사용하십시오). 이를 위해 사용자 정의 reduce 연산자로 CUB의 감소를 사용할 수 있습니다 (값과 비교, 데이텀은 {value, index} 임).