cudaThreadSynchronize & performance

Question

며칠 전 동일한 알고리즘의 매우 간단한 대체 및 스러스트 구현을 수행하는 광산 코드의 성능을 비교했습니다. 추력에 찬성하여 한 자리 수 (!)의 불일치를 발견 했으므로 디버거를 코드에 "서핑"하여 마법의 발생 위치를 발견하기 시작했습니다.

놀랍게도 필자는 모든 펑터를 제거하고 핵심 요소에 도달하면 매우 직접적인 구현이 실제로는 매우 유사하다는 것을 발견했습니다. 나는 Thrust가 블록 _size & grid_size (btw : 정확히 어떻게 작동합니까?)를 결정할 영리한 방법이 있다는 것을 알았습니다. 그래서 설정을 변경하고 코드를 다시 실행했습니다. 몇 마이크로 초가되었지만 거의 같은 상황이었습니다. 그런 다음, 결국, 왜 "단지"시도해보아야하는지 모르겠다. 나는 커널과 빙고 다음에 cudaThreadSynchronize()를 제거했다! 나는 그 격차 (그리고 더 나은)를 제로로 설정하고 실행 시간의 전체 등급을 얻었다. 배열의 값에 액세스하면 예상 한대로 정확하게 실행된다는 것을 알았습니다.

질문은 지금 다음과 같습니다. 언제 내가 cudaThreadSynchronize (et similia)를 제거 할 수 있습니까? 왜 그렇게 큰 오버 헤드가 발생합니까? 그 추력 그 자체가 (__THRUST_SYNCHRONOUS 매크로가 정의되어 있지 않고, 그렇지 않은 경우 NOP 인) synchronize_if_enabled (const char * message)가 동기화되지 않는 것을 볼 수 있습니다. 자세한 내용은 & 코드를 참조하십시오.

// my replace code 
template <typename T> 
__global__ void replaceSimple(T* dev, const int n, const T oldval, const T newval) 
{ 
    const int gridSize = blockDim.x * gridDim.x; 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 
    while(index < n) 
    { 
     if(dev[index] == oldval) 
      dev[index] = newval; 
     index += gridSize; 
    } 
} 

// replace invocation - not in main because of cpp - cu separation 
template <typename T> 
void callReplaceSimple(T* dev, const int n, const T oldval, const T newval) 
{  
    replaceSimple<<<30,768,0>>>(dev,n,oldval,newval); 
    cudaThreadSynchronize(); 
} 

// thrust replace invocation 
template <typename T> 
void callReplace(thrust::device_vector<T>& dev, const T oldval, const T newval) 
{ 
    thrust::replace(dev.begin(), dev.end(), oldval, newval); 
} 

파람 상세 : 배열 : N = 10,000,000 요소는 2로 설정 인 oldval = 2 된 newval = 3

• 스러스트 callReplace (추력) 실행 시간 : 0.057 밀리 실행할
• 시간 동기와 callReplaceSimple : 0.662 MS 동기화없이 callReplaceSimple을 실행하기
• 시간 : 0.011 밀리

나는 CUDA 5.0 재치를 사용 h 추력 포함, 내 카드는 GeForce GTX 570이며 2GB RAM의 쿼드 코어 Q9550 2.83GHz를 사용합니다.

Answer 1

커널 시작은 비동기입니다. cudaThreadSynchronize() 호출을 제거하면 커널이 완료 될 때까지의 시간이 아니라 커널 시작 시간 만 측정합니다.