## NVIDIA nvcc compiler 

The NVIDIA nvcc  compiler driver is responsible for generating the final executable file, which is a combination of host (CPU) and device (GPU) codes. 

IBM XL and GCC compilers can be used to generate the host code and their flags. As described in 7.1, “Compiler options” on page 162, this process produces an optimized code to run in the POWER9 processor. By default, nvcc  uses the GCC, unless the -ccbin  flag is passed to set another back-end compiler. 

Note:  The nvcc  compiler uses the GCC C++ compiler by default to generate the host code. If wanted, the IBM XL C++ compiler can instead use the -ccbin xlC  option. 

The NVIDIA Tesla V100 GPUs are built on Tesla GV100 architecture, which provides CUDA compute capability V7.0. Use the -gencode  compiler option to set the virtual architecture and binary compatibility. For example, -gencode arch=compute_70,code=sm_70  generates code that is compatible with the Tesla GV100 architecture (virtual architecture V7.0). 

Example 7-13 shows a CUDA program that prints to standard output the index values for 2D thread blocks. The executable file is built with nvcc , which uses xlC as the host compiler (by using the -ccbin xlC  option). Parameters to xlC are passed by the -Xcompiler option. 

$ cat -n hello.cu #include<cuda_runtime.h> #include<stdio.h> __global__ void helloKernel() { int tidx = threadIdx.x; int tidy = threadIdx.y; int blockidx = blockIdx.x; int blockidy = blockIdx.y; printf("I am CUDA thread (%d, %d) in block (%d, %d)
", tidx, tidy, blockidx, blockidy); }; int main(int argc, char* argv[]) { dim3 block(16,16); dim3 grid(2,2); helloKernel<<<grid, block>>>(); cudaDeviceSynchronize(); return 0; } $ nvcc -ccbin xlC -Xcompiler="-qarch=pwr9 -qtune=pwr9 -qhot -O3" -gencode arch=compute_70,code=sm_70 hello.cu -o helloCUDA $ ./helloCUDA I am CUDA thread (0, 10) in block (1, 0) I am CUDA thread (1, 10) in block (1, 0) I am CUDA thread (2, 10) in block (1, 0) I am CUDA thread (3, 10) in block (1, 0) I am CUDA thread (4, 10) in block (1, 0) I am CUDA thread (5, 10) in block (1, 0) <... Output omitted ...> 

