CUDA 워프 투표 기능으로 인해 코드가 느려 집니까?

Question

2D의 점 집합에서 볼록 엔벌 로프를 계산하는 CUDA 함수를 작성했습니다. 하지만 은 매우 CPU 코드보다 느립니다!

나는 워프 투표 기능과 __syncronisation()을 사용하고 있습니다. 상당히 많은 횟수. 그러면 코드가 느려 집니까?

덕분에 코드를 추가

:

__global__ void find_edges_on_device(TYPE * h_x, TYPE * h_y, int *h_edges){ 

int tidX = threadIdx.x; 
int tidY = threadIdx.y; 
int tid = tidY*blockSizeX + tidX; 
int i = threadIdx.x+blockIdx.x*blockDim.x; 
int j = threadIdx.y+blockIdx.y*blockDim.y; 

int hxi = h_x[i]; 
int hxj = h_x[j]; 
int hyi = h_y[i]; 
int hyj = h_y[j]; 

long scalarProduct = 0; 
TYPE nx; 
TYPE ny; 

bool isValid = true; 

__shared__ int shared_X[blockSizeX*blockSizeY]; 
__shared__ int shared_Y[blockSizeX*blockSizeY]; 
__shared__ bool iswarpvalid[32]; 
__shared__ bool isBlockValid; 

if (tid==0) 
{ 
    isBlockValid=true; 
} 
if (tid<(blockSizeX*blockSizeY-1)/32+1) 
{ 
    iswarpvalid[tid]=true; 
} 
else if (tid<32) 
{ 
    iswarpvalid[tid]=false; 
} 

//all the others points should be on the same side of the edge i,j 
//normal to the edge (unnormalized) 
nx = - (hyj- hyi); 
ny = hxj- hxi; 
int k=0; 
while ((k==i)||(k==j)) 
{ 
    k++; 
} //k will be 0,1,or 2, but different from i and j to avoid 
scalarProduct=nx* (h_x[k]-hxi)+ny* (h_y[k]-hyi); 
if (scalarProduct<0) 
{ 
    nx*=-1; 
    ny*=-1; 
} 

for(int count = 0; count < ((NPOINTS/blockSizeX*blockSizeY) + 1); count++){ 

    int globalIndex = tidY*blockSizeX + tidX + count*blockSizeX*blockSizeY; 

    if (NPOINTS <= globalIndex){ 
     shared_X[tidY*blockSizeX + tidX] = -1; 
     shared_Y[tidY*blockSizeX + tidX] = -1; 
    } 
    else { 
     shared_X[tidY*blockSizeX + tidX]= h_x[globalIndex]; 
     shared_Y[tidY*blockSizeX + tidX]= h_y[globalIndex]; 
    } 
    __syncthreads(); 

    //we have now at least one point with scalarProduct>0 
    //all the other points should comply with the same condition for 
    //the edge to be valid 
    //loop on all the points 

    if(i < j){ 
     for (int k=0; k < blockSizeX*blockSizeY; k++) 
     { 
      if((count * blockSizeX*blockSizeY + k < NPOINTS)&&(isValid)) { 
       scalarProduct=nx* (shared_X[k]-hxi)+ny* (shared_Y[k]-hyi); 
       if(__all(scalarProduct) < 0){ 
        iswarpvalid[(tidY*blockSizeX + tidX)/32] = false; 
        break; 
       } 
       else if(0 > (scalarProduct)){ 
        isValid = false; 
        break; 
       } 
      } 
     } 
    } 

    __syncthreads(); 
    if (tid<32) 
    { 
     isBlockValid=__any(iswarpvalid[tid]); 
    } 
    __syncthreads(); 
    if(!isBlockValid) break; 
} 

if ((i<j) && (true == isValid)){ 
      int tmp_i = i; 
      int tmp_j = j; 

      if(-1 != atomicCAS(&h_edges[2*i], -1, tmp_j)) 
       h_edges[2*i+1]=j; 

      if(-1 != atomicCAS(&h_edges[2*j], -1, tmp_i)) 
       h_edges[2*j+1]=i; 

} 
} 

Answer 1

당신은 NVIDIA CUDA C 프로그래밍 가이드에서 찾을 수 있습니다 찾고있는 답변.

섹션 5.4.3 진술한다 :

처리량 __syncthreads()를 계산 능력 1.x에서의 장치 및 계산 능력 (2)의 장치에 대한 클록 사이클 당 16 작업 클럭 사이클 당 8 작업이기 .엑스.

워프 투표 기능은 PTX ISA 매뉴얼의 섹션 B.12 및 표 109에서 다루어집니다. 후자는 워프 투표를 수행하기 위해 두 가지 지침이 필요하다는 것을 나타냅니다. 그러나 참조 문서에서 워프 투표 기능에 대한 클록 사이클 수치를 파악할 수 없었습니다.