mod 대신 bitwise shift를 사용하는 동안 CUDA의 중요한 시간 차이

Question

저는 CUDA 프로그래밍이 처음이었고 실행 시간의 차이를보기 위해 다른 간단한 커널을 테스트했습니다. 최적화 및 얻을 없애 내가 -g -G와 함께 컴파일

__global__ void mathKernel4(float *c) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    float ia, ib; 
    ia = ib = 0.0f; 

    int itid = tid >> 5; 

    if (itid & 0x01 == 0) 
    { 
     ia = 100.0f; 
    } 
    else 
    { 
     ib = 200.0f; 
    } 

    c[tid] = ia + ib; 
} 

:

__global__ void mathKernel2(float *c) 
{ 
int tid = blockIdx.x * blockDim.x + threadIdx.x; 
float ia, ib; 
ia = ib = 0.0f; 

    if ((tid/warpSize) % 2 == 0) 
    { 
     ia = 100.0f; 
    } 
    else 
    { 
     ib = 200.0f; 
    } 

    c[tid] = ia + ib; 
} 

다른 하나

은 나를 위해 동일한 작업을 수행 개의 커널을 가지고 :

mathKernel2 <<< 8192 32 >>> elapsed 0.000259 sec 
mathKernel4 <<< 8192 32 >>> elapsed 0.000103 sec 

왜 큰 차이가 있습니까?

Answer 1

이 라인 : false 및 분기의 하나에 컴파일시 평가

if (itid & (0x01 == 0)) 

제거된다

if (itid & 0x01 == 0) 

같이 읽는다.

비트 단위 이진 연산자는 C/C++ 및 CUDA에서 다소 우선 순위가 낮습니다.