CUDA_Samples.tex

\chapter{CUDA SDK Sample Codes}
\label{chap:cuda-sdk-sample}


\section{Starting CUDA programming}
\label{sec:start-cuda-progr}

\begin{verbatim}
1. Copy the template project

        cd <SDK_INSTALL_PATH>/projects
        cp -r template <myproject>

2. Edit the filenames of the project to suit your needs

        mv template.cu myproject.cu
        mv template_kernel.cu myproject_kernel.cu
  mv template_gold.cpp myproject_gold.cpp

3. Edit the Makefile and source files.  Just search and 
replace all occurrences of "template" with "myproject".

4. Build the project

        make

   You can build a debug version with "make dbg=1", 
        an emulation version with "make emu=1", and 
        a debug emulation with "make dbg=1 emu=1".
[Emulation is no longer supported in CUDA 3.1]

5. Run the program

        ../../bin/linux32/release/myproject

   (It should print "Test PASSED")

6. Now modify the code to perform the computation you require.  
   See the CUDA Programming Guide for details of 
   programming in CUDA.
\end{verbatim}

The library \verb!libcutil! is created to help you performing some
tasks. It is provided as is by the SDK examples (libcutil is simply
for convenience -- it is not a part of CUDA and is not required for
your own CUDA programs). You can copy it to your project to use it for
free.

\subsection{Compile a single sample}
\label{sec:comp-single-sample}

If you download only a single sample, then to compile it, e.g.
{\it reduction} application, you need to use the local Makefile
\begin{verbatim}
cd NVIDIA GPU Computing SDK/C/src/reduction
make Makefile
\end{verbatim}

If we use PGI C++ compiler {\bf pgcpp}
\begin{verbatim}
pgcpp simpleCUBLAS.c -I/opt/cuda/include -D__GNUC__ -L/opt/cuda/lib64 -lcublas
setenv LD_LIBRARY_PATH /opt/cuda/lib64
a.out
\end{verbatim}
or
\begin{verbatim}
pgcpp simpleCUBLAS.c -I/opt/cuda/include 
  -I/opt/cuda-3.0-beta/sdk/C/common/inc  -D__GNUC__ 
  obj/x86_64/release/MonteCarlo*.o 
  -L/opt/cuda/lib64 -lcudart
  -L/opt/cuda-3.0-beta/sdk/C/common/lib/linux 
  -L/opt/cuda-2.3/lib64/
  -L../../lib 
  -L../../common/lib/linux
  -L../../../shared/lib -lcudart
  -L/opt/cuda-2.3/lib64
  -L../../lib
  -L../../common/lib/linux
  -L../../shared/lib -lcudart -lcutil_x86_64 -shrutil_x86_64
setenv LD_LIBRARY_PATH /opt/cuda/lib64
a.out

\end{verbatim}

NOTE: We need to use the flag \verb!-D__GNUC__! since CUDA header
files don't recognize pgcpp. 


\textcolor{red}{There are many case that you need to include the file}
\verb!cutil_inline.h!
\textcolor{red}{locate at the ``common'' folder}.

\subsection{libshrutil\_x86\_64}
\label{sec:libshrutil_x86_64}

Since CUDA you have an additional shared utility, beside libcutil,
called \verb!libshrutil_x86_64! (for 64-bit) or \verb!libshrutil_x86!
(for 32-bit). 

If you compile a single project in CUDA SDK, you need to add
\begin{verbatim}
NVIDIA GPU Computing SDK/shared/lib
\end{verbatim}
to the link path
\begin{verbatim}
nvcc -I../../common/lib -L../../../shared/lib
\end{verbatim}
or the \verb!LD_LIBRARY_PATH! environment variable. 


\section{Reduction}
\label{sec:reduction}

\begin{itemize}
\item CUDPP support reductions for INT, UNIT, FLOAT, CHAR, UCHAR.

\item Thrust support reductions for all types, including DOUBLE

\item CUBLAS has: \verb!cublasIsamax()!, yet slower than Mark Harris'
  s reduction algorithm (CUDA SDK)
\begin{verbatim}
int cublasIsmax(int n, const float* x, int incx)
int cublasIdmax(int n, const double* x, int incx)
\end{verbatim}
  return the smallest index of the maximum magnitude element of vector
  $x$, with the pattern $x[1+i*incx]$, $i=0..n-1$. The result is
  1-based indexing for compatibility with Fortran.
\begin{verbatim}
int cublasIsmax(int n, const float* x, int incx)
int cublasIsmax(int n, const double* x, int incx)
\end{verbatim}
  returns the sum of absolute values of the elements in the vector
  $x$.

NOTE:
\textcolor{red}{CUBLAS 2.0 support both single-precision and
  double-precision.}

\item CUDA sample: \verb!reduction_kernel.cpp! by Mark Harris. Except
  reduce0, there are 6 versions which can be chosen via the
  \verb!whichKernel!  parameters (default is 6). NOTE: The code
  requires the array size $n$ to be {\it power of 2}, so you may need
  a code to set up the right array size first.
  \begin{enumerate}
  \item [reduce0] slowest, use modulo to choose the active thread
  \item [reduce1] use contiguous threads, yet has many shared
    memory bank conflicts
  \item [reduce2] use sequential addressing - no divergence or bank
    conflicts
  \item [reduce3] use $n/2$ threads (with $n$ is number of elements) -
    perform 1st level of reduction
  \item [reduce4] unroll the last warp - need a minimum of
    64*sizeof(T) bytes of shared memory (T is the type of data
    elements).
  \item [reduce5] completely unroll 
  \item [reduce6] add multiple elements per thread sequentially -
    previous versions process a single pair of elements per thread
  \end{enumerate}

  NOTE: The size of thread blocks must be the power of 2: 1, 2, 4, 8,
  16, 32, 64, 128, 256, 512. 
\end{itemize}


\subsection{Min/Max}
\label{sec:minmax}


\subsection{Sum}
\label{sec:sum}

\section{A sample header file}
\label{sec:header_file_gpumcml}


\url{http://gpumcml.googlecode.com/svn-history/r197/trunk/gpumcml/simple/gpumcml_kernel.h}
\begin{lstlisting}
/*****************************************************************************
 *
 *   Header file for GPU-related data structures and kernel configurations
 *
 ****************************************************************************/
/*   
 *   This file is part of GPUMCML.
 * 
 *   GPUMCML is free software: you can redistribute it and/or modify
 *   it under the terms of the GNU General Public License as published by
 *   the Free Software Foundation, either version 3 of the License, or
 *   (at your option) any later version.
 *
 *   GPUMCML is distributed in the hope that it will be useful,
 *   but WITHOUT ANY WARRANTY; without even the implied warranty of
 *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
 *   GNU General Public License for more details.
 *
 *   You should have received a copy of the GNU General Public License
 *   along with GPUMCML.  If not, see <http://www.gnu.org/licenses/>.
 */

#ifndef _GPUMCML_KERNEL_H_
#define _GPUMCML_KERNEL_H_

#include "gpumcml.h"

//////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////

/*  Number of simulation steps performed by each thread in one kernel call
*/
#define NUM_STEPS 50000  //Use 5000 for faster response time

//////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////

// Make sure __CUDA_ARCH__ is always defined by the user.
#ifdef _WIN32 
  #define __CUDA_ARCH__ 120
#endif

#ifndef __CUDA_ARCH__
#error "__CUDA_ARCH__ undefined!"
#endif

/**
 * Although this simple version of GPUMCML is not intended for high
 * performance, there are still a few parameters to be configured for
 * different GPUs.
 *
 * - NUM_BLOCKS:
 *      number of thread blocks in the grid to be launched
 *
 * - NUM_THREADS_PER_BLOCK:
 *      number of threads per thread block
 *
 * - EMULATED_ATOMIC:
 *      Enable this option for GPUs with Compute Capability 1.1,
 *      which do not support 64-bit atomicAdd to the global memory.
 *      In this case, we use two 32-bit atomicAdd's to emulate the
 *      64-bit version.
 *
 * - USE_TRUE_CACHE:
 *      Enable this option for GPUs with Compute Capability 2.0 (Fermi),
 *      which have a 64KB configurable L1 cache in each SM.
 *      If enabled, the L1 cache is configured to have 48KB of true cache
 *      and 16KB of shared memory, as opposed to 16KB of true cache and
 *      48KB of shared memory. Since the shared memory is not utilized
 *      in this simple version, you are encouraged to enable this option
 *      to cache more accesses to the absorption array in the global memory.
 */

/////////////////////////////////////////////
// Compute Capability 2.0
/////////////////////////////////////////////
#if __CUDA_ARCH__ == 200

#define NUM_BLOCKS 30
#define NUM_THREADS_PER_BLOCK 512
// #define EMULATED_ATOMIC
#define USE_TRUE_CACHE

/////////////////////////////////////////////
// Compute Capability 1.2 or 1.3
/////////////////////////////////////////////
#elif (__CUDA_ARCH__ == 120) || (__CUDA_ARCH__ == 130)

#define NUM_BLOCKS 30
#define NUM_THREADS_PER_BLOCK 256
#define EMULATED_ATOMIC

/////////////////////////////////////////////
// Compute Capability 1.1
/////////////////////////////////////////////
#elif (__CUDA_ARCH__ == 110)

#define NUM_BLOCKS 14       // should match the number of SMs on the GPUs
#define NUM_THREADS_PER_BLOCK 192
#define EMULATED_ATOMIC

/////////////////////////////////////////////
// Unsupported Compute Capability
/////////////////////////////////////////////
#else

#error "GPUMCML only supports compute capability 1.1 to 2.0!"

#endif

//////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////

/**
 * Derived macros
 */

#define NUM_THREADS (NUM_BLOCKS * NUM_THREADS_PER_BLOCK)

//////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////

typedef struct __align__(16)
{
  FLOAT init_photon_w;      // initial photon weight 

  FLOAT dz;                 // z grid separation.[cm] 
  FLOAT dr;                 // r grid separation.[cm] 

  UINT32 na;                // array range 0..na-1. 
  UINT32 nz;                // array range 0..nz-1. 
  UINT32 nr;                // array range 0..nr-1. 

  UINT32 num_layers;        // number of layers. 
} SimParamGPU;

typedef struct __align__(16)
{
  FLOAT z0, z1;             // z coordinates of a layer. [cm] 
  FLOAT n;                  // refractive index of a layer. 

  FLOAT muas;               // mua + mus 
  FLOAT rmuas;              // 1/(mua+mus) 
  FLOAT mua_muas;           // mua/(mua+mus)

  FLOAT g;                  // anisotropy.

  FLOAT cos_crit0, cos_crit1;
} LayerStructGPU;

// The max number of layers supported (MAX_LAYERS including 2 ambient layers)
#define MAX_LAYERS 100

__constant__ SimParamGPU d_simparam;
__constant__ LayerStructGPU d_layerspecs[MAX_LAYERS];

//////////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////////

// Thread-private states that live across batches of kernel invocations
// Each field is an array of length NUM_THREADS.
//
// We use a struct of arrays as opposed to an array of structs to enable
// global memory coalescing.
//
typedef struct
{
  // cartesian coordinates of the photon [cm]
  FLOAT *photon_x;
  FLOAT *photon_y;
  FLOAT *photon_z;

  // directional cosines of the photon
  FLOAT *photon_ux;
  FLOAT *photon_uy;
  FLOAT *photon_uz;

  FLOAT *photon_w;            // photon weight
  FLOAT *photon_sleft;        // leftover step size [cm]

  // index to layer where the photon resides
  UINT32 *photon_layer;

  UINT32 *is_active;          // is this thread active?
} GPUThreadStates;


typedef struct
{
  // cartesian coordinates of the photon [cm]
  FLOAT x;
  FLOAT y;
  FLOAT z;

  // directional cosines of the photon
  FLOAT ux;
  FLOAT uy;
  FLOAT uz;

  FLOAT w;            // photon weight

  FLOAT s;            // step size [cm]
  FLOAT sleft;        // leftover step size [cm]

  // index to layer where the photon resides
  UINT32 layer;

  // flag to indicate if photon hits a boundary
  UINT32 hit;
} PhotonStructGPU;

#endif // _GPUMCML_KERNEL_H_

\end{lstlisting}



%%% Local Variables: 
%%% mode: latex
%%% TeX-master: "gpucomputing"
%%% End: