\documentclass{article}
\usepackage{amsmath}
\usepackage[mathletters]{ucs}
\usepackage[utf8x]{inputenc}
\setlength{\parindent}{0pt}
\setlength{\parskip}{6pt plus 2pt minus 1pt}
\usepackage{graphics}
\usepackage[ps,dvips,xdvi,arrow,frame,matrix,graph,curve,cmtip]{xy}
\usepackage[pdftex]{graphicx}

\usepackage[breaklinks=true]{hyperref}

\setcounter{secnumdepth}{0}
\title{Week 4 Progress Report}
\author{Abdulmajed Dakkak}
\date{July 4th, 2008}
\begin{document}
\maketitle

\section{Saturday/Sunday}

Looked over the Schaber method, and figured out how to Schaberize
simpleGL. Wrote a rule 90 1D cellular automata in openGL to
experiment with C's boolean operators. The 1D cellular automata was
then ported to CUDA, although the graphics are off, since I do not
know anything about \emph{vertex buffer objects}. Finally, I
implemented the FHP model.

Read a little bit about the FCHC model.

\subsection{Code}

\begin{itemize}
\item 
  \href{code/fhp.c}{FHP}
\item 
  \href{code/rule90.c}{Open GL Rule 90}
\item 
  \href{code/CA.tar.gz}{rule90\_cuda.tar.gz}
  (\href{code/CA/ca.cu}{Cuda's Rule 90 CA},
  \href{code/CA/ca_kernel.cu}{Kernel})
\end{itemize}
\section{Monday}

Schaberized the \verb!simpleGL! example included in the CUDA
projects, and worked with Chase and Will to migrate fluidsGL from
Unix Sockets to Winsocks.

\subsection{Code}

\begin{itemize}
\item 
  \href{code/dynamic_simpleGL.tar.gz}{Schaberized SimpleGL}(
  \href{code/dynamic_simpleGL/dyn_simpleGL_client.c}{client},
  \href{code/dynamic_simpleGL/dyn_simpleGL_server.cu}{server},
  \href{code/dynamic_simpleGL/dyn_simpleGL_kernel.cu}{kernel})
\end{itemize}
\section{Tuesday}

Decided that I do not care much about Windows, and started my
effort to more the \verb!fhp! code from openGL to CUDA. Converted
my representation of the lattice from a 2D object to a 1D flat
array. It turns out that this might not have been a good choice.

\subsection{Code}

\begin{itemize}
\item 
  \href{code/fhp-cuda.tar.gz}{FHP CUDA} (
  \href{code/fhp-cuda/fhp.cu}{fhp.cu},
  \href{code/fhp-cuda/fhp_kernel.cu}{fhp\_kernel.cu},
  \href{code/fhp-cuda/types.h}{types.h})[\textbf{Does not work}]
\end{itemize}
\section{Wednesday}

Joined Will and Chase, again, in researching for ways of making the
Schaber method work in Windows. This meant going through the
Windows documentation, forums, and some Winsock books trying to
find why Winsock does not implement \verb!MSG_WAITALL!.

Formed a pact to stay all night (if we have to) to get fluidGL into
Sysgz. In the end we had to and stayed. Around 9AM the next morning
the next night we got something that should work in the CAVE.

\section{Thursday}

\subsection{CUDA/FHP}

Made the \verb!fhp! program work on the CUDA chip. The problem was
with the following piece of code

\begin{verbatim}
dim3 block(32, 32, 1);
dim3 grid(LATTICE_WIDTH / block.x, LATTICE_HEIGHT / block.y, 1);
collision<<<grid, block>>>(gpu_c1, gpu_c2, gpu_c3, gpu_c4, gpu_c5, gpu_c6,
                           gpu_k1, gpu_k2, gpu_k3, gpu_k4, gpu_k5, gpu_k6,
                           gpu_boundary);   
propagate<<<grid, block>>>(gpu_c1, gpu_c2, gpu_c3, gpu_c4, gpu_c5, gpu_c6,
                           gpu_k1, gpu_k2, gpu_k3, gpu_k4, gpu_k5, gpu_k6);
\end{verbatim}
From my understanding of CUDA, a few grids
(\verb!LATTICE_WIDTH /  block.x! to be precises) containing
\verb!LATTICE_HEIGHT  / block.y! blocks which in turn contain
\verb!block.x * block.y! or \verb!32*32! threads would be allocated
on the GPU. If that fails, then CUDA will figure everything out and
everyone would be happy. This is not the case, however, and the
results were not seen in the form of an error in this function.
Rather, when I tried to copy the results after computation.using
\verb!cudaMemcpy!

\begin{verbatim}
cudaMemcpy(c1, gpu_c1, ARRAY_LENGTH, cudaMemcpyDeviceToHost);
cudaMemcpy(c2, gpu_c2, ARRAY_LENGTH, cudaMemcpyDeviceToHost);
cudaMemcpy(c3, gpu_c3, ARRAY_LENGTH, cudaMemcpyDeviceToHost);
cudaMemcpy(c4, gpu_c4, ARRAY_LENGTH, cudaMemcpyDeviceToHost);
cudaMemcpy(c5, gpu_c5, ARRAY_LENGTH, cudaMemcpyDeviceToHost);
cudaMemcpy(c6, gpu_c6, ARRAY_LENGTH, cudaMemcpyDeviceToHost);
\end{verbatim}
I was getting the same array as I had before; i.e.
\verb!gpu_c[1-6] = c[1-6]!. This meant that I only saw the initial
positions of particles, and nothing changed. I knew, however, that
something must have changed, because, at the very least, the
positions should have propagated.

After a few hours (\verb!4! to be more precise) I figured out that
the problem can be fixed by setting the block size to \verb!16*16!
rather than \verb!32*32!.

\begin{verbatim}
dim3 block(32, 32, 1);
\end{verbatim}
This was found on the forums, among many other places. This is the
first time I also saw the immaturity of CUDA. Take, for example,
the following piece of code:

\begin{verbatim}
int x = powf(2.0, 3.0);
\end{verbatim}
In regular mathematics we are taught that $2^3 = 8$, but in CUDA's
opinion the answer is $7$. This phenomenon continues for other odd
powers of 2.

\subsection{Schaberization}

I also stuck around trough the night with Chase and Will helping
them in making the Phleet and CUDA talk. Will and Chase did the
majority of the work, and I only helped a bit with debugging.
Incidentally, the problem they were encountering, namely with some
dimensions working while others not, might be related to the
problem with odd powers of 2 found in CUDA. Their programs fail,
have interesting behavior, or are slow if you use odd powers of two
(although we are limited by what we could try).

\subsection{Code}

\begin{itemize}
\item 
  \href{code/fhp-cuda2.tar.gz}{FHP CUDA} (
  \href{code/fhp-cuda2/fhp.cu}{fhp2.cu},
  \href{code/fhp-cuda2/fhp_kernel.cu}{fhp\_kernel2.cu},
  \href{code/fhp-cuda2/types.h}{types2.h})
\end{itemize}
\section{Friday}

After the long night (and subsequents sleep), I noticed that the
toroidal boundary is not working in \verb!fhp-cuda!. Thus one of my
aims today is to fix that issue. Possibly the most important thing,
however, is to optimize the code. This means that I have to learn
more about CUDA's memory architecture and use Vertex Buffer Objects
(or \verb!vbo!); something I dreaded from the beginning.

I also noticed that CUDA's profiler which I downloaded and tried
two weeks ago (ans is invaluable for the optimization) now gives a
segfault. I will talk to Johnathan on Monday to see if this is
related to the update done last week.

\end{document}
