爲什麼整數除法和模數不NVRTC

優化掉了我編了內核NVRTC：爲什麼整數除法和模數不NVRTC

__global__ void kernel_A(/* args */) { 
    unsigned short idx = threadIdx.x; 
    unsigned char warp_id = idx/32; 
    unsigned char lane_id = idx % 32; 
    /* ... */ 
}

我知道整數除法和模數都在CUDA的GPU非常昂貴。不過，我認爲這種劃分按冪-2應該被優化成位運算，直到我發現它是不是：

__global__ void kernel_B(/* args */) { 
    unsigned short idx = threadIdx.x; 
    unsigned char warp_id = idx >> 5; 
    unsigned char lane_id = idx & 31; 
    /* ... */ 
}

似乎kernel_B只是運行速度更快。當省略內核的所有其他代碼，大小爲1024的1024塊發射，nvprof顯示kernel_A在平均15.2us運行，而kernel_B運行7.4us平均。我推測NVRTC沒有優化整數除法和模數。

結果是在GeForce 750 Ti，CUDA 8.0平均100次調用中獲得的。給nvrtcCompileProgram()的編譯器選項是-arch compute_50。

這是預期嗎？

來源

2017-06-01 Kh40tiK

你不需要推測。通過'cuobjdump -sass'運行可執行文件以查找。 – tera

@tera我正在通過NVRTC進行JIT編譯，所以沒有可執行文件。任何組裝方式傾銷？ – Kh40tiK

NVRTC發出PTX並將其傳遞給驅動程序進行JIT編譯。你可以用'nvrtcGetPTX'提取PTX。那麼你不需要猜測 – talonmies

在代碼庫中做了徹底的bugsweep。原來我的應用程序內置於DEBUG模式。這會導致額外的標誌-G和-lineinfo傳遞給nvrtcCompileProgram()

從nvcc手冊頁：

--device-debug(-G)

生成設備代碼調試信息。關閉所有優化。請勿用於分析;改用-lineinfo。

來源

2017-06-10 13:37:20 Kh40tiK

爲什麼整數除法和模數不NVRTC

回答

相關問題