NCHC00210201F.ppt

* This should be emphasized! Maybe another slide on “This is the first thing” * * * * Try to update tool slides with the tutorial? * * * * * What are the options that the students can use? * ? David Kirk/NVIDIA and Wen-mei W. Hwu Taiwan, June 30-July 2, 2008 * Compiling a CUDA Program NVCC C/C++ CUDA Application PTX to Target Compiler G80 … GPU Target code PTX Code Virtual Physical CPU Code Parallel Thread eXecution (PTX) Virtual Machine and ISA Programming model Execution resources and state float4 me = gx[gtid]; me.x += me.y * me.z; ld.global.v4.f32 {$f1,$f3,$f5,$f7}, [$r9+0]; mad.f32 $f1, $f5, $f3, $f1; ? David Kirk/NVIDIA and Wen-mei W. Hwu Taiwan, June 30-July 2, 2008 * Compilation Any source file containing CUDA language extensions must be compiled with NVCC NVCC is a compiler driver Works by invoking all the necessary tools and compilers like cudacc, g++, cl, ... NVCC outputs: C code (host CPU Code) Must then be compiled with the rest of the application using another tool PTX Object code directly Or, PTX source, interpreted at runtime ? David Kirk/NVIDIA and Wen-mei W. Hwu Taiwan, June 30-July 2, 2008 Linking Any executable with CUDA code requires two dynamic libraries: The CUDA runtime library (cudart) The CUDA core library (cuda) ? David Kirk/NVIDIA and Wen-mei W. Hwu Taiwan, June 30-July 2, 2008 Debugging Using theDevice Emulation Mode An executable compiled in device emulation mode (nvcc -deviceemu) runs completely on the host using the CUDA runtime No need of any device and CUDA driver Each device thread is emulated with a host thread Running in device emulation mode, one can: Use host native debug support (breakpoints, inspection, etc.) Access any device-specific data from host code and vice-versa Call any host function from device code (e.g. printf) and vice-versa Detect deadlock situations caused by improper usage of __syncthreads ? David Kirk/NVIDIA and Wen-mei W. Hwu Taiwan, June 30-July 2, 2008 Device Emulation Mode Pitfalls