공유 메모리가있는 CUDA에서 병렬 병합을 구현해도 성능이 향상되지 않습니다.

Question

CUDA에서 병렬 병합 알고리즘을 구현하려고합니다. 이 알고리즘은 하나의 스레드 블록에서 실행되도록 설계되었습니다. 기본 개념은 두 입력 시퀀스에서 각 요소의 전역 순위를 계산하는 것입니다. 이 두 입력 시퀀스가 ​​정렬되기 때문에 요소의 전역 순위는 원래 시퀀스의 인덱스와 이진 검색에 의해 계산 된 다른 시퀀스의 순위와 같습니다. 이러한 알고리즘을 구현하는 가장 좋은 전략은 두 개의 시퀀스를 공유 메모리에로드하여 전역 메모리 읽기를 줄이는 것입니다. 그러나 공유 메모리를 사용하지 않고 공유 메모리를 사용하는 구현 버전과 공유 메모리를 사용하지 않는 구현 버전을 비교했을 때 성능 향상을 볼 수 없었습니다. 내가 뭔가 잘못하고 있는지 궁금해.

하드웨어 : GeForce GTX 285, Linux x86_64. 두 구현에 대해 두 개의 1024 요소 시퀀스를 병합하는 시간은 약 0.068672ms입니다.

__global__ void localMerge(int * A, int numA,int * B,int numB,int * C){ 
extern __shared__ int temp[]; // shared memory for A and B; 
int tx=threadIdx.x; 
int size=blockDim.x; 
int *tempA=temp; 
int *tempB=temp+numA; 

int i,j,k,mid; 
    //read sequences into shared memory 
for(i=tx;i<numA;i+=size){ 
    tempA[i]=A[i]; 
} 
for(i=tx;i<numB;i+=size){ 
    tempB[i]=B[i]; 
} 
__syncthreads(); 
    //compute global rank for elements in sequence A 
for(i=tx;i<numA;i+=size){ 
    j=0; 
    k=numB-1; 
    if(tempA[i]<=tempB[0]){ 
     C[i]=tempA[i]; 
    } 
    else if(tempA[i]>tempB[numB-1]){ 
     C[i+numB]=tempA[i]; 
    } 
    else{ 
     while(j<k-1){ 
      mid=(j+k)/2; 
      if(tempB[mid]<tempA[i]){ 
       j=mid; 
      } 
      else{ 
       k=mid; 
      } 
     } 
     //printf("i=%d,j=%d,C=%d\n",i,j,tempA[i]); 
     C[i+j+1]=tempA[i]; 
    } 
} 
    //compute global rank for elements in sequence B 
for(i=tx;i<numB;i+=size){ 
    j=0; 
    k=numA-1; 
    if(tempB[i]<tempA[0]){ 
     C[i]=tempB[i]; 
    } 
    else if(tempB[i]>=tempA[numA-1]){ 
     C[i+numA]=tempB[i]; 
    } 
    else{ 
     while(j<k-1){ 
      mid=(j+k)/2; 
      if(tempA[mid]<=tempB[i]){ 
       j=mid; 
      } 
      else{ 
       k=mid; 
      } 
     } 
     //printf("i=%d,j=%d,C=%d\n",i,j,tempB[i]); 
     C[i+j+1]=tempB[i]; 
    } 
}  
} 

Answer 1

넌 __shared__ 메모리에 모두 입력리스트를 병렬 세밀한 이진 검색의 수집에 의존하기보다는 "merge path"알고리즘을 적용하여보다 운을 가질 수있다. 이 알고리즘에 재사용이 존재하기 때문에 캐시에서 꽤 잘 캡쳐 할 수 있기 때문에이 문제에 대해 __shared__ 메모리를 사용하는 것이 덜 중요합니다.

이 병합 알고리즘을 사용하면 CTA의 각 스레드가 병합 된 결과에서 k 출력을 생성하는 것이 좋습니다. 이것은 각 스레드의 작업이 대략 균일하고 좋은 결과를 가져 왔으며 관련 바이너리 검색은 상당히 거칠다.

스레드 i은 입력리스트 을 한 번에 번 검색하여 k*i 번째 출력 요소의 각 목록에서 위치를 찾습니다. 그런 다음 작업은 간단합니다. 각 스레드는 k 항목을 입력 목록에서 순차적으로 병합하고 출력에서 ​​위치 k*i에 복사합니다.

자세한 내용은 Thrust's implementation을 참조하십시오.