text:
\section{Keywords:} Large Hadron Collider, High Performance Computing, GPU, CUDA, Kokkos, Performance Portability, Particle Physics}
\end{abstract}

\section{Introduction}
The ATLAS detector at the Large Hadron Collider (LHC) relies heavily on massive samples of simulated proton-proton collision events for detector modeling, upgrade studies and to deliver high-quality physics results.
However, standard simulations — using Geant4~\cite{Agostinelli:2002hh} for accurate and detailed modeling of physics processes — are extremely CPU-intensive, and its use for all Monte Carlo-based studies is impractical due to the formidable size of the LHC dataset.
In particular, simulating electromagnetic and hadronic shower developments in the ATLAS calorimeters is extremely CPU-intensive, comprising more than 90\% of the overall simulation time \cite{ATLAS_Sim_Infrastructure}.

To reduce the significant processing time required to simulate the calorimeters,  ATLAS developed FastCaloSim~\cite{ATL-SOFT-PUB-2018-002}\cite{ATLAS_PUB_REPO}, a primarily C++ package that uses a simplified detector geometry and parameterizations of shower development initiated by particles traversing the calorimeter volumes.
The parameterizations are derived using single particles — electrons, photons and charged pions — with initial energy between 1~GeV and 4~TeV,  uniformly distributed in polar angle.
During a simulation, the optimal parameterization is selected based on the particle type, energy deposit and location in the detector in order to best model the corresponding electromagnetic or hadronic shower.
The parameterizations are then used to create calorimeter \textit{hits}, which contain information about the location, and amount of energy deposited in the detector.
Thus, the more initial and secondary particles produced in a shower, the larger the number of hits.
Following the simulation, hits are then processed and used for event reconstruction.
Therefore, it is the number of hits, not necessarily the number of initial particles, that determines the length of time required for a particular simulation.

A parameterized simulation can reduce the CPU time by a factor of 10 to 25 — depending on the process being modeled — relative to that of standard Geant4-based simulations. However,
it should be noted that, with respect to the full detector simulation, FastCaloSim is not the speed-limiting component.
Nevertheless, there are a number of advantages of porting FastCaloSim to accelerators, in particular, GPUs.
First, an intelligent port of the serial CPU code to GPU-friendly code permits parallelization over events (event-level parallelization) and also over particles within a single event (intra-event, or particle-/hit-level parallelization).
The ability to parallelize in these ways becomes exceedingly important when simulating not only a single proton-proton collision (the ``hard-scatter'' event) but when additional collisions within the same bunch-crossing are also simulated (``pile-up'' events).
For example, at the LHC there are on average 30\cite{ATLAS_RUN_CONDITIONS} pile-up events for every one hard-scatter;
in the conditions of the High-Luminosity LHC (HL-LHC)\cite{HLLHC_RUN_CONDITIONS}, there are expected to be more than 200 pile-up events for every hard-scatter. ATLAS foresees\cite{ATLAS-TDR} utilizing fast simulation for 60-90\% of its simulation workload for various HL-LHC scenarios, making highly-parallelizable codes  of this nature extremely desirable.
Lastly, an initial port of a set of codes such as those comprising FastCaloSim, that makes significant use of GPUs and demonstrates their computational ability, can be invaluable to the high-energy physics community in general.
It can serve as a proof of concept and motivation for other groups and collaborations to design and implement their own codes to execute on accelerator hardware, and prompt investigations into alternative portable solutions to exploit the wide range of modern architectures that are becoming increasingly pervasive.

\section{Test Data}\label{sec:test-data}
The inputs to FastCaloSim are tuples containing information about the incoming particle: the particle type and associated four-vector,
which consists of the particle's energy and momentum.
Additional details -- e.g. calorimeter layer, cell and amount of energy deposited -- are also included to determine the correct parameterization as the particle traverses the detector.

For the studies presented here, several different inputs were used; single-particle samples -- electrons, photons and charged pions -- with various initial energies between 16~GeV and 4~TeV and polar angle within the acceptance of the calorimeter; and a sample consisting of $t\bar{t}$ -- top-anti-top quark -- production at a center-of-mass proton-proton collision energy of 13~TeV.
This is a configuration similar to that as was used in Reference \cite{ATL-SOFT-PUB-2018-002}.
Single-particle inputs contain only one primary initial particle that results in a relatively small number, $O(1)$,
of secondaries being produced per event as the primary particle interacts with the calorimeter material.
In contrast, the $t\bar{t}$ simulation results in more realistic complex final states, with $O(1000)$ secondaries produced per event.
Such large per-event particle multiplicities could potentially utilize more GPU compute -- provided they have sufficient energy to traverse, and not stop in, the calorimeter -- and can therefore represent a good benchmark for algorithm performance.

\section{Performance Profiling of the Unaltered FastCaloSim}
In the ATLAS production runs, FastCaloSim is part of the Athena software framework~\cite{ATLAS_PUB_REPO}. In order to make it easy to experiment with the GPU porting, we used the standalone version of FastCaloSim to reduce the number of dependencies that are typically needed in the Athena framework. The version of the standalone FastCaloSim package we used consists of roughly 40,000 lines of C/C++ code.

In order to determine where the computational hotspots are located in FastCaloSim, we ran performance profiling using the test data described in Section~\ref{sec:test-data} for electrons with an energy of 65.536~GeV and $\eta_{\rm min}=0.2$, where $\eta$ is the pseudorapidity, describing the angle between a particle and the beam axis. The test was run using a single core on Intel Xeon CPU E5-2695 v4 at 2.10 GHz. We generated call tree information using \texttt{valgrind --tool=callgrind}, and visualized the results with  \texttt{KCachegrind}. We found that more than 40\% of time was spent reading and unpacking the input ROOT data files, while the actual simulations consumed about 57\% of the total execution time. The focus of this work and subsequent discussions is on accelerating the simulation part, which is amenable to parallelization and  is contained in about O(1000) lines of code. The I/O overhead will limit the overall speedup we can achieve on the GPUs, which will also be discussed.

\section{CUDA Implementation}
For porting to the GPU we focused on the function which, other than data unpacking, consumed the majority of the execution time: \texttt{TFCSLateralShapeParametrizationHitChain::Simulate\_hit()}  (see Listing \ref{hitsiml}). This function performs the hit simulation  in a calorimeter layer for a passing particle. In this function there is a loop over the number of hits which can range from 1 to 15000  depending on particle type, energy, angle and calorimeter layer. For the input sample file of  electrons with  $E=65.536$ GeV and $\eta_{\rm min}=0.2$, the number of hits typically range from 4000 to 6500.   Each hit  loop entry will decide into which calorimeter cell unit this simulated hit will deposit energy. The calculation is based on particle position, energy, angle and a random parametrization.  The hit simulation loop is only offloaded to the GPU when the number of hits is over 500 for a particular layer. Each GPU thread will simulate one hit.  If the number of hits is less than 500, or it is in a  forward detector layer (layer 21-24) \cite{ATLAS_Detector}, the calculation is left the on host. This is to ensure that the GPU can offer computational speedup compared to the CPU, which can only be achieved if there are a large enough number of hits (GPU threads) to keep the GPU busy.

\begin{lstlisting}[language=C++, frame=single, caption=Hit simulation function, basicstyle=\small, label=hitsiml]
TFCSLateralShapeParametrizationHitChain::simulate(
                                TFCSSimulationState& simulstate
                                const TFCSTruthState* truth,
                                const TFCSExtrapolationState* extrapol ) {
    int nhit = get_number_of_hits( simulstate, truth, extrapol );
    float Ehit = simulstate.E( calosample() ) / nhit;

    for ( int i = 0; i < nhit; ++i ) {  //  hit simulation loop
      TFCSLateralShapeParametrizationHitBase::Hit hit;
      hit.E() = Ehit;
      for ( TFCSLateralShapeParametrizationHitBase* hitsim : m_chain ) {
         hitsim->simulate_hit( hit, simulstate, truth, extrapol );
      }
    }
}

\end{lstlisting}

The hit simulation calculation requires detector geometry data and  parametrization tables. To offload  the calculation, this data first needs to be loaded onto the GPU.

 For the  data used, the detector geometry included  187,652  calorimeter elements and 24 layers or samples, located in  78 regions.
 However, on the host side, the detector geometry data structures make heavy use of the C++ Standard Template Library (STL)  and pointers.  For example, the host region data structures include  \texttt{std::vector<std::vector<CaloDetDescrElements *>>}.  These kinds  of data structures are not supported by CUDA - in general, STL data structures will automatically re-allocate memory as their size changes, a feature not supported on GPUs due to their very limited memory resources. Furthermore, copying many small  chunks  of  non-contiguous  data generated with \texttt{new()} from CPU to GPU is  very inefficient.   We thus flattened the data, using arrays instead of pointers, and moved scattered  data together  before copying them to the GPU.
 After being copied to the GPU, all calorimeter detector element data as well as all region data were kept together as array of \texttt{structs}. The total size of the geometry data was about 20MB.

 Similarly, the parametrization tables, which are held in an \texttt{std::vector}, were transformed into a simple array before being copied to the GPU. Since the complete parametrization data is very large, only the needed tables were loaded onto the GPU during simulation. Furthermore, the same table is only loaded once, so if subsequent events require the same tables, they will not be loaded again. The total size of tables loaded to GPU depends on the type of simulations. Single particle simulations have a small footprint, while the data size for multi-particle simulations can be big. In our benchmark study for electrons with $E=65.5$~GeV, a simulation for 10,000 events only needs 1.2MB of table data loaded to the GPU,  while in the $t\bar{t}$ simulation with $E=13$~TeV, 500 multi-particle events need to have 1.1GB of tables loaded.

 Hit simulation requires a large amount of random numbers. Each hit needs 3 random numbers.  The \texttt{cuRand}  library function was used to generate random numbers on the GPU.  However this will only produce statistically equivalent and not binary identical result when compared to CPU only calculations.  For code and physics validation, random numbers were also generated on the host, and copied to the GPU for calculations.  This technique was only used for validation, and not for performance benchmarks, as it is much slower.

 Other than the random number generator, there are 3  CUDA kernels in the code. The first, \texttt{clean}, resets the workspace. The number of threads is the same as the total number of detector cells.  The second or main kernel, \texttt{sim\_a}, does the hit simulation,  which calculates which cell gets the hit, and deposits energy into it. Each thread simulates one hit.  Atomic operations were chosen to accumulate the energy deposits, as the number of hits is relatively small. A block reduction method  was also tried, but did not demonstrate any performance gain, and introduced more kernels.  The last kernel,   \texttt{sim\_ct},   counts the hit cells and gathers the result for more efficient data transfer back to the host. The total number of threads is again the same as the total number of detector cells.
 The kernels are launched each time the code enters  the hit simulation loop,  i.e. for each calorimeter sample layer a particle is passing through. After the third kernel executes, the results, an  array of   \texttt{struct \{long long  cell\_id;  float  E \} } are copied back to host.  For all 3 kernels, we used block size 256, i.e. 256 threads per block, in our benchmark.

To avoid many small kernel launches that would be required if the random numbers were generated on demand, GPU memory was pre-allocated, and a large number of random numbers were generated before hand, so that they would not need to be generated in every hit simulation loop. Regeneration of random numbers is only triggered when they are all consumed. In our benchmark test, we generate 200,000 random numbers each time. The infrequent generation of these random numbers places a very small burden on the GPU, and does not significantly affect performance.

We noticed that for some combinations of (particle, energy, $\eta$),  the numbers of hits were very small. Since we only offloaded the simulation to the GPU when there were a minimum of 500 hits, most  of the calculation could still be done on the host.  To solve this problem, we explored the method of grouping hit simulations. We combined  different layers, particles, events together so that each GPU simulation  kernel launch can process more hits, with tasks of multiple entries to the simulation function as shown in Listing \ref{hitsiml}.  In order to combine these tasks, every time the hit simulation part is executed, instead of directly  doing the calculations on either the GPU or the CPU, we only collect the parameters needed for simulation, such as particle type, position, layer, number of hits, energy etc.  The  simulation calculation is delayed until enough hits have been accumulated. This information is then loaded onto the GPU and the simulation kernel is launched with more threads. Each thread can discover its own parameters needed for the hit it simulates and perform proper simulation. Our preliminary results in Section~\ref{sec:cuda_bench} show that this technique can lead to positive gains.

\section{Performance Benchmarks}\label{sec:cuda_bench}

To perform the benchmarks, we used the GPU nodes on the Brookhaven National Lab Institutional Cluster (BNL-IC) and the NERSC Cori GPU cluster.  The BNL-IC node has 2 Intel\textsuperscript{\textregistered} Xeon\textsuperscript{\textregistered} CPU E5-2695 processors, each with 18 cores.  It  also  has 2 NVIDIA P100 GPUs.  Only one GPU was used when running a single instance of the code.
Two GPUs were used only when we tested multiple instances of the program running  with  \texttt{nvidia-cuda-mps-server}.

Cori GPU nodes  have 2 Intel\textsuperscript{\textregistered} Xeon\textsuperscript{\textregistered} Gold 6148 CPU and 8 NVIDIA V100 GPUs.  Our benchmark job used only one CPU core and one GPU.

The program starts by reading parametrization data files, unpacking them, and setting up the simulation chains. It then goes through an event loop to perform simulation for each event. The accumulated results  are saved  after event loop is completed.  Other than the event loop, the majority of the time was spent on unpacking data which doesn't change much with different  input sample files.

Our benchmarks collect time  spent in the event loop  and for the complete run by wrapping the code sections with calls to \texttt{std::chrono::system\_clock::now()}.

Table \ref{tab:allp} shows the result from running on the BNL-IC, comparing the original CPU  code and GPU accelerated code run time for a few different  particles and pseudorapidities with an energy of less than 65.5~GeV.
Multi-particle  input data  has 500 events  while all others have 10,000 events.
For multi-particle data which simulate $t\bar{t}$ events with center of mass energy of 13~TeV, each event has about 600-800 particles with different energies and angles.
All  timings were obtained  by running on the same node.  Numbers were  averaged over 10 repeated runs.

  \begin{table}[h]
\renewcommand{\arraystretch}{1.3}
\begin{center}
    \begin{tabular}{| r | c | c || c | c| c || c |c|c||}
    \hline
    &&& \multicolumn{3}{|c||}{Total Time(seconds)} & \multicolumn{3}{c||}{Event Loop time (seconds)} \\
    \cline{4-9}
    \multicolumn{1}{|c|}{particle} & E(GeV)& $\eta_{min}$& CPU & GPU & GPU$_g$& CPU & GPU & GPU$_g$ \\
    \hline
    Electron         & 65.5 & 0.0  & 34.62 & 19.26 &17.54 & 18.30 &2.71 & 0.90  \\
    Electron          & 65.5 & 0.2 & 34.97 & 18.98 & 17.61 & 18.67 &2.37 &0.97  \\
    Electron        & 65.5 & 2.2 & 36.09 & 19.13 &17.83  & 19.69  &2.56&1.19\\
    Electron        & 65.5 & 3.4 &17.95 &18.22  & 18.14 & 1.60 &1.64&1.54 \\
    \hline
     Photon         &32.8 & 0.2 &26.15 &18.60 &18.29 &9.71  &2.06 &0.80 \\
     Photon         &65.5 & 0.2 &34.94 &18.91 &17.83  &18.60  &2.31 &0.97 \\
     \hline
     Pion         &16.4 & 0.2 &19.02 &18.41 &17.46  &2.67  &1.76 &0.89 \\
     Pion         &32.8 & 0.2 &21.06 &18.83 &17.67  &4.57  &2.17 &1.08 \\
     Pion         &65.5 & 0.2 &24.13 &19.28 & 17.90 &7.68  &2.70 &1.30 \\
     \hline
     \hline
     Multi-particle &&& 52.69& 46.39&36.17&35.87&29.30&19.05\\

     \hline
    \end{tabular}
\end{center}
    \caption{ Execution time for the complete job and for only the event loop part of the job, for various input files.
    GPU$_g$ is for GPU group simulation code.}
    \label{tab:allp}
\end{table}

In Table \ref{tab:allp}, the CPU code event loop time roughly reflects the total number of hits. The fact that the GPU code  does not offload simulation of forward detector layers is  shown in the case of \texttt{$\eta_{min}$=3.4} where  most of the  simulation calculation was done on the host because they not only have small numbers of hits but are also  mostly on the forward detector layers.  Therefore in this case the GPU codes are even  slower due to overheads of loading geometry etc. We can see that for the data file with more hits, such as electrons and photons with higher energy  and smaller $\eta$, both GPU and group simulation (Column GPU$_g$) can reduce total run time by 40-50\% compare to CPU code. For the event loop part of the run, GPU code can run faster than CPU code by a factor of 6-8 while group simulation can improve this by another factor of 2-3.  For  input data with fewer hits, such as Pions, the gain is less.  For  the multi-particle data we found that in the simulation, half of the layers only received 1 hit and many  layers have small numbers of hits, with the average hit count to be about 50. This is due to many low energy particles in each event, resulting in a low GPU simulation rate.  Furthermore, a lot more parametrization data need to be loaded onto  the  GPU  due to the variety of particles, energies and angles.   Hence in this case  the speedup on event loop is only a factor of 1.2 for GPU code and 1.9 for GPU$_g$ code.

Table \ref{tab:nsc} shows the results of running on the NERSC Cori GPU node which are consistent with the BNL-IC results. Note that Cori has slightly faster CPUs and GPUs.

\begin{table}[h]
\renewcommand{\arraystretch}{1.3}
\begin{center}
    \begin{tabular}{| r | c | c || c | c| c || c |c|c||}
    \hline
    &&& \multicolumn{3}{|c||}{Total Time(seconds)} & \multicolumn{3}{c||}{Event Loop time (seconds)} \\
    \cline{4-9}
    \multicolumn{1}{|c|}{particle} & E(GeV)& $\eta_{min}$ & CPU & GPU & GPU$_g$& CPU & GPU & GPU$_g$ \\
    \hline
    Electron          & 65.5 & 0.2 & 34.13 &18.12&  16.87 & 17.98 &2.02&0.88  \\
    \hline
    Multi-particle   &&& 51.5&42.4&34.7&  35.1& 26.2 &18.1 \\
     \hline
    \end{tabular}
\end{center}
    \caption{ Execution time of complete run and event loop on Cori GPU node.}
    \label{tab:nsc}
\end{table}

Simulations with higher energy (1~TeV - 4~TeV) for fewer particles were also benchmarked  on the BNL-IC. Results are shown in Table \ref{tab:high}.
We can clearly see that the GPU code speed-up is much better than for the lower energy simulations. Speed-ups of up to a factor of 60 can be achieved for the event loop, and 4.6 for the total run.
It should be noted that group simulation code does not gain as much over normal GPU code for electrons and photons compared to lower energy simulations.
In the case of higher energy particle simulations, such as with the 4~TeV electrons, performance of the group simulation can be worse than the regular one, due to the fact that each layer already has many simulated hits which make normal GPU code calculation already efficient, while the group simulation code needs to perform extra steps collecting simulation parameters.

\begin{table}[h]
\renewcommand{\arraystretch}{1.3}
\begin{center}
\resizebox{\textwidth}{!}{%

    \begin{tabular}{| r | c | c |c|c|c|c| c || c|c | c |c|c||}
    \hline
    &&& \multicolumn{3}{|c|}{Total Time(s)} &\multicolumn{2}{|c||}{Speed Up}& \multicolumn{3}{|c|}{Event Loop time (s)}&\multicolumn{2}{|c||}{Speed Up} \\
    \cline{4-13}
    \multicolumn{1}{|c|}{particle} & E (TeV)& Events & CPU & GPU & GPU$_g$&GPU&GPU$_g$& CPU & GPU& GPU$_g$&GPU&GPU$_g$ \\
    \hline
Electron&	1.0&	3000&81.4&18.02&	17.60&	4.5&	4.6&	65.12&	1.58&	1.10&	41&	59\\
Electron&	2.0&	1000	&51.0&	17.12&	17.10&	3.0&	3.0&	34.62&	0.70&	0.63&	50&55\\
Electron&	4.0&	900&	60.5&	17.18&	17.33&	3.5&	3.5&	44.23&	0.71&	0.84&	62&	53\\
Photon&	1.0& 	3000&	80.5&	18.15&	17.54&	4.4&	4.6&	63.83&	1.59&	1.06&	40&	60\\
Photon&	2.0&	1000	&50.4&	17.29&	17.14&	2.9&	2.9&	34.14&	0.78&	0.61&	44&	56\\
Photon&	4.0& 	1000	&64.5&	17.38&	17.35&	3.7&	3.7&	48.21&	0.92& 	0.84&	53&	57\\
Pion&	1.0 & 3000&	37.4&	18.47&	17.54&	2.0&	2.1&	21.17&	1.99&	1.08&	11&	20\\
Pion&	2.0	& 1000&	30.8& 	17.16&	16.88&	1.8&	1.8&	14.52&	0.67&	0.37&	22&	39\\
\hline
    \end{tabular}
    }
\end{center}
    \caption{ Execution time of complete run,  event loop  and speed up over CPU only code for high energy input files.
     All simulations are with $\eta_{min}= 0.2$.  Speed up is over CPU code.}
    \label{tab:high}

\end{table}

In all cases  above, a single instance of the job will only use  a small portion of the available GPU resources.  Current high energy physics production calculations are performed  in the High Throughput Computing (HTC) environments  where many jobs are run on each multi-core node at the same time, usually with one thread or one process per core.  To emulate this environment, we started multiple concurrent instances of the same job, and shared the GPUs on the node between them by means of \texttt{nvidia-cuda-mps-server}\cite{cuda-mps}.

Figure \ref{fig:nproc} shows the result of running the benchmark on the same BNL-IC node, launching 2-32 processes of the program with CPU, GPU and group simulation codes. Input sample data is Electron with \texttt{E=65.5~GeV, $\eta_{min}$=0.2} and 10,000 events. Since \texttt{nvidia-cuda-mps-server} for the NVIDIA P100 GPUs only supports a maximum 16 processes, we used 2 GPUs for a  maximum of 32 processes for our benchmark.  We can see that the CPU code slows down and the curve flattens as the number of concurrent jobs increases. This is due to core clock speed throttling from turbo boost when more cores are busy.
The event loop time for per-event based simulation GPU code increased by a factor of more than 2 when the number of processes increased to 32, while the performance of grouped simulation codes is not significantly degraded as the number of concurrent jobs increase.
While both codes have same amount of calculations, the per-event GPU code has significantly more small kernel launches and small data copy API calls than the grouped simulation. For many of them, the kernel launch latency overheads are greater than the actual kernel run times and data transfer times. The shift of the per-event GPU code (green triangles) indicates that the cuda-mps overhead is fairly significant compared to the GPU run time, while such overhead is negligible in the other cases where run times are significantly longer. The fact that we can run multiple concurrent processes without increasing much the execution time demonstrates that the computing units of GPUs are not fully utilized by any single job instance, even when shared among many jobs. However the overheads of the large number of small kernel launches and small data transfers start to saturate the cuda-mps managed GPUs when the number of processes reaches around 20 for the per-event based simulation. For that simulation we pinned the \texttt{nvidia-cuda-mps-server} processes as well as the simulation processes to the NUMA domains corresponding to  the GPU devices that they are associated with.

\begin{figure}[h!]
\begin{center}
\includegraphics[width=0.8\textwidth]{figs/cuda_mps_2.eps
\end{center}
\caption{ Normalized timing of multiple concurrent jobs running on the same node sharing GPUs with \texttt{nvidia-cuda-mps-server}}\label{fig:nproc}
\end{figure}

\section{Validation of the CUDA Implementation}
Since GPU ported code uses the \texttt{cuRAND} random number generator, which creates a different random number set from the CPU  simulation, we do not expect that the simulation results should be exactly comparable between the CPU and GPU. We should, however, be able to produce statistically consistent results, which we have checked and confirmed.
In order to further validate the GPU code, we used the option of copying the same  set of CPU generated random numbers to the GPU, performing the simulation on the GPU, and comparing the results with the simulation performed solely on the CPU.

We compared the results of around 62 million hits in the Electron 65.5~GeV run,  and found that  in the GPU simulation, there were 2 hits whose calculations ended up in different calorimeter cells compared to the CPU calculation. Further investigation showed that this was  due to floating point rounding error differences between CPU and GPU. The 2 mismatches can be eliminated by using double precision (\texttt{double}) variables for certain calculations instead of single precision \texttt{float}.  We can thus conclude that the GPU ported code is correctly performing the same calculations as the original CPU code.

\section{Portable Implementation with Kokkos}

Kokkos \cite{kokkos} is a  C++ library, developed at Sandia National Laboratory, with a rich history
and a strong user base. It is a
portable, performant, single source shared-memory programming model that abstracts away the hardware, allowing developers to write algorithms once and run on
any supported backend architectures. Currently supported backends are serial processing on a traditional
CPU, parallel processing on multi and many-core CPUs via pThreads and OpenMP, and NVIDIA and
AMD GPUs via CUDA and HIP. An Intel GPU backend is currently under development as well.
A common API that hides the architecture-specific implementation
details greatly reduces the intimate device knowledge that would otherwise be needed by developers.
At compile time, Kokkos permits the selection of one host serial
backend (e.g. Intel Haswell architecture), one
host parallel backend (e.g. OpenMP), and one device parallel backend (e.g. NVIDIA V100). If different hardware
backends are desired, both the user code and Kokkos runtime library must be recompiled.

Kokkos has a number of different parallel abstractions available,
such as \texttt{parallel\_for}, \texttt{parallel\_reduce}, and \texttt{parallel\_scan}. It also provides
utilities such as random number generators, and supports
atomic operations. Data access is accomplished via multi-dimensional structures that are templated in the
type of the data called
\texttt{Kokkos::Views}. The sizes of the structures can be specified at either run or compile time.
Kokkos' usage of C++ language features in kernels are only limited by the compilers necessitated by
the selected backends - e.g. if the CUDA backend is chosen for NVIDIA hardware, and the available version
of the nvcc compiler is 10.2, then \texttt{C++14} is the latest standard that can be used in a kernel.

\subsection{Build Infrastructure}

FastCaloSim uses cmake as its build tool. Modifying the cmake configuration files to pickup the Kokkos
installation was relatively simple, as Kokkos provides the infrastructure support for cmake, with
\texttt{find\_package(Kokkos)}.
Configuring Kokkos was more complicated. There are three build parameters of Kokkos that must be
selected when installing the library, which can affect how the FastCaloSim code is organized:
\begin{itemize}
    \item The hardware backends (host serial, host parallel, device parallel)
    \item The use of shared libraries
    \item The use of relocatable device code
\end{itemize}
For testing purposes, the choice of the hardware backends is not problematic, as the developer knows
where the code will run. It is more of an issue in a production environment where binaries might be
distributed to a broad range of hardware architectures, requiring multiple different configurations to be
built. The choice of using shared vs static libraries is usually dependent on what already exists in the
code base that will be using Kokkos. In this case, FastCaloSim uses shared libraries, and modifying
it to use static libraries would have made re-integration to the original code base much more
challenging, so we decided to stay with shared libraries. Unfortunately, enabling shared libraries and
relocatable device code are mutually exclusive in the Kokkos build. Disabling relocatable device
code has significant implications for the structure of the kernel sources - all symbols in a kernel
must be visible and resolvable within a single compilation unit.

In order to keep all source files as synchronized as possible, we chose to keep the serial, CUDA and
Kokkos codes in the same git branch, selecting which feature to build at cmake configuration time via
command line switches, such as \texttt{-DENABLE\_GPU=on} and \texttt{-DUSE\_KOKKOS=on}.

\subsection{Porting CUDA Kernels to Kokkos}

If the CUDA backend of Kokkos is selected as the device parallel
backend, CUDA and Kokkos code can freely interoperate, with Kokkos kernels having the ability to call
CUDA, and access data allocated via \texttt{cudaMalloc}. This feature allows for
the incremental migration of CUDA kernels to Kokkos, simplifying testing and reproducibility. One
by one, we rewrote the CUDA kernels as Kokkos lambda expressions and inlined functions,
to be executed as \texttt{parallel\_for}
loops, but maintaining the CUDA memory allocations, and accessing the bare device pointers in
the kernels. At each step, we validated the outputs to make sure that the results were consistent with the original CUDA code.

Once all kernels had been migrated to Kokkos, we tackled the data structures. Most are simple scalars,
1D and 2D arrays. These were easy to translate to \texttt{Kokkos::Views}. A couple of data structures,
however, are jagged 2D arrays, i.e. arrays of arrays where the sizes of the inner arrays are not all the
same. While it is technically possible to create \texttt{Kokkos::View< Kokkos::View< T >>}, this is not
advised by the Kokkos developers, and requires a significant amount of extra boiler-plate code to create
and access. Instead we chose to flatten the jagged arrays, either by adding a new array which indexed
the starting locations of the inner arrays in a 1D structure when there was significant variation in the
sizes of the inner arrays, or by padding to the largest size of the inner arrays, when they were close
in size.

\subsection{Source File Organization}
Since Kokkos' use of relocatable device code was disabled, we had to structure the source files
so that all symbols were visible within a single compilation unit. This was accomplished by
creating a wrapper file that merely included the various individual kernel source files. Since
the majority of the kernel device functions were identical for the Kokkos and CUDA implementations,
in order to share as much of the code as possible, the shared code was in a single file, with
CUDA and Kokkos specific features triggered via preprocessor macros (see listing \ref{devsel}).

\begin{lstlisting}[language=C++, frame=single, caption=Language functionality selection, basicstyle=\small, label=devsel]
#ifdef USE_KOKKOS
#  include <Kokkos_Core.hpp>
#  include <Kokkos_Random.hpp>
#  define __DEVICE__ KOKKOS_INLINE_FUNCTION
#else
#  define __DEVICE__ __device__
#endif

__DEVICE__ void HitCellMapping_d( Hit& hit, Chain0_Args args ) {
  long long cellele = getDDE( args.geo, args.cs, hit.eta(), hit.phi() );

#ifdef USE_KOKKOS
  Kokkos::View<float*> cellE_v( args.cells_energy, args.ncells );
  Kokkos::atomic_fetch_add( &cellE_v( cellele ), hit.E() );
#else
  atomicAdd( &args.cells_energy[cellele], hit.E() );
#endif
   ...
}
\end{lstlisting}

Implementations that were very different between the Kokkos and CUDA versions were kept in
separate files, and included by wrappers appropriately.

Another complication is that Kokkos will compile source files differently depending on the
suffix of the file name. If the suffix is \texttt{.cu}, the NVIDIA compiler \texttt{nvcc} will be
used for the compilation with a specific set of command line switches that is different than if
the same file had a \texttt{.cpp} or \texttt{.cxx} extension. Compilation errors will occur
if Kokkos core headers are exposed to \texttt{nvcc} in \texttt{.cu} files, which required
some careful organization
of file names and wrappers.

\subsection{Random Numbers}
A large number of random numbers are required during the execution of FastCaloSim. They
are used in code that is executed on the device, and are not needed on the host. Thus it would
make sense to generate them on the device, to avoid large memory transfers. Both CUDA and
Kokkos provide libraries to generate high quality random numbers. However, in order to verify
that different implementations produce the same results, it is very useful to be able to always
use the same set of random numbers. We thus provided the ability to generate random numbers
either on the host, and transferred to the device as needed, or directly on the device using
either the CUDA or Kokkos RNGs.

\subsection{Performance Studies}

Once we had fully migrated all kernels and data structures to Kokkos, we were able to test the
performance of the code on multiple different parallel devices. The hardware platform was
a dual Intel Xeon Gold 5220 CPU with 18 physical cores per CPU, for a total of 72 hyperthreaded
cores. The host OS was Centos~7, and packages were compiled with gcc8.2. Also installed
on the server were a NVIDIA RTX 2080 Super GPU, and an AMD Vega 56 GPU. We used Kokkos version 3.1
for testing purposes, which was configured to target the CPU with \texttt{Kokkos\_ARCH\_SKX}, and the GPUs
with \texttt{Kokkos\_ARCH\_TURING75} and \texttt{Kokkos\_ARCH\_VEGA900}. The installed version of CUDA
was 10.2.89, and of the ROCm libraries was 3.5.1. We also converted the CUDA code directly to HIP using
\texttt{hipify-perl}. For the sake of brevity, when we refer to HIP
in the subsequent sections, we mean the hcc backend of HIP that targets AMD devices, as opposed
to the nvcc backend that targets NVIDIA devices.

In Table \ref{tab:kokkos_perf} we show the relative execution times of various kernels with
different Kokkos backends, using the
CUDA implementation as a baseline, and reporting times as a factor \textbf{slower} than that of
the CUDA implementation.

The aspects that we measured were the time it takes to re-initialize a large array of approximately
200000 floats (\texttt{clean})
at the start of each event, the time to execute the main kernel (\texttt{sim\_a}), the time to execute the
reduction kernel (\texttt{sim\_ct}), the time it takes to copy the results back from the device to the host
(\texttt{copy d$\rightarrow$h}), and finally the time for the full event loop of 10000 events. This was
done by placing synchronization barriers around the kernels, and using host based timers to
measure the execution times. The best
performance using the pThread and OpenMP backends were achieved using 15 threads/processes.

\begin{table}[h]
\renewcommand{\arraystretch}{1.3}
\begin{center}
    \begin{tabular}{| r | c | c | c | c | c || c |}
    \hline
    & \multicolumn{5}{|c||}{Kokkos} & \\
    \cline{2-7}
    \multicolumn{1}{|c|}{kernel} & CUDA & HIP$^1$ & Serial & pThread & OpenMP & HIP$^2$\\
    \hline
    clean           & 2.83 & 20.1 & 7.67 & 1.61 & 2.26 & 5.02 \\
    sim\_a          & 1.17 & 37.9 & 36.9 & 11.2 & 11.5 & 37.0 \\
    sim\_ct         & 1.35 & 8.93 & 12.9 & 1.66 & 1.99 & 3.99 \\
    copy d$\rightarrow$h & 1.30 & 3.92 & 0.02 & 0.06 & 0.04 & 2.98 \\
    event loop      & 1.16 & 11.9 & 7.26 & 2.61 & 2.79 & 6.83 \\
    \hline
    \end{tabular}
\end{center}
    \caption{Relative kernel average execution times with various Kokkos backends, as a factor of the pure CUDA implementation execution time.
    HIP\textsuperscript{1} is for Kokkos with the HIP backend, whereas HIP\textsuperscript{2} is a direct port from CUDA to HIP}
    \label{tab:kokkos_perf}
\end{table}

We also directly measured the performance of the kernels using \texttt{nvprof} and \texttt{rocprof}, as
show in Table \ref{tab:gpu_perf}.

\begin{table}[h]
\renewcommand{\arraystretch}{1.3}
\begin{center}
    \begin{tabular}{| r | c | c || c |}
    \hline
    & \multicolumn{2}{|c||}{Kokkos} & \\
    \cline{2-4}
    \multicolumn{1}{|c|}{kernel} & CUDA & HIP$^1$ & HIP$^2$ \\
    \hline
    clean           & 1.60 & 4.07 & 2.62 \\
    sim\_a          & 0.93 & 43.6 & 44.3 \\
    sim\_ct         & 1.43 & 3.58 & 2.41 \\
    copy d$\rightarrow$h & 1.02 & 1.65 & 1.32 \\
    \hline
    \end{tabular}
\end{center}
    \caption{Relative kernel average execution times with CUDA and HIP\textsuperscript{1} Kokkos backends, and a direct HIP\textsuperscript{2} port, as a factor of the pure CUDA implementation execution time, as measured by GPU profiling tools \texttt{nvprof} and \texttt{rocprof}}
    \label{tab:gpu_perf}
\end{table}

From these measurements, we can extract several important conclusions. First, using Kokkos to
initialize memory on the device does impose some significant penalties as compared to a simple
CUDA operation. This is due to extra steps required to create the host views before the
device memory is reset with a \texttt{Kokkos::deep\_copy}. In general though, while we do see some overhead from using Views and
some small penalties from launching Kokkos kernels, the
performance of the CUDA backend of Kokkos is reasonably close to that of the pure CUDA implementation.
As the kernels get smaller and simpler, e.g. comparing the performance of \texttt{sim\_ct} to that
of \texttt{sim\_a}, the Kokkos overhead become more noticeable.

Second, the performance of AMD GPU is noticeably worse than that of the NVIDIA GPU. While it is
hard to do a direct comparison as the architectures are very different, the clock speeds and
number of computation units are within 25\% of each other, yet we see timing differences in
some cases of one and a half orders of magnitudes. Further tests of simple kernels seem to
suggest that the kernel launch latencies on the AMD GPU are significantly higher than that
of the NVIDIA GPUs (see Table \ref{tab:launch_latency}). The clock frequency of the CPU on the
host also has a significant impact on the launch latency.

\begin{table}[h]
\renewcommand{\arraystretch}{1.4}
\begin{center}
    \begin{tabular}{| c | c || c | c || c |}
    \hline
    & &\multicolumn{2}{|c||}{Kokkos} & \\
    \cline{3-4}
     CPU freq & CUDA & CUDA & HIP & HIP$^2$ \\
     \hline
     2200 MHz & 5.6 & 9.4 & 152 & 88 \\
     3700 MHz& 3.4 & 5.3 & 60 & 30 \\
    \hline
    \end{tabular}
\end{center}
    \caption{Kernel launch latencies in microseconds as a function of CPU clock frequency}
    \label{tab:launch_latency}
\end{table}

It is also worth noting that when the job was pinned to a single CPU core, the timings for the CUDA and Kokkos/CUDA did not change appreciably, however the time to execute the full event loop when utilizing the HIP backed of Kokkos increased by 30\%, suggesting that scheduling and launching kernels
on AMD devices places significant burden on the CPU.

\subsection{Porting vs Rewriting}

The method that we chose to translate the CUDA code to Kokkos was obviously non-optimal, and in
some ways
deliberately so. Part of the exercise was to understand the issues involved in porting a
project which may have a large amount of existing code from serial C++ to CUDA, and from CUDA to Kokkos,
while attempting
to leave existing serial C++ and CUDA code unmodified as much as possible. Given the
relatively small size of the FastCaloSim code base, it would have been easier to completely rewrite
the entire code base, refactoring to optimize for Kokkos data structures and kernels. This would
no doubt have improved the performance of the Kokkos version. However, ultimate performance is
not the end goal of our study, but rather performant portability, reproducibility, and the burden
of code maintenance are more important metrics. With a large code base, it is very important
to make sure that the it is easy to maintain compatibility between all implementation flavours,
and also to ease migration of code from serial C++$\rightarrow$CUDA$\rightarrow$Portability Layer.

\section{Conclusions}
We have successfully ported the ATLAS fast calorimeter simulation code, FastCaloSim, to use GPUs. The code was first ported to NVIDIA GPUs using the proprietary CUDA API. As with many HTC applications, the compute-to-I/O ratio of FastCaloSim is not very large, limiting the potential gains of using GPUs. Several techniques have been incorporated in our implementations to increase the speed and efficiency of the computations on the GPUs, as well as the GPU's overall utilization. The performance gains on the GPUs are greatly dependant on the types of simulations (particle type, energy, angle etc.). For the test cases we ran, we observed overall speedups ranging from none to about 5 times. But if we compare only the timing of the simulation part, the speedups on the GPUs were much more significant, as much as 60 times relative to executing with a single thread on a CPU core. The GPU was found to be underutilized, and a single GPU may be shared among many concurrent CPU jobs of FastCaloSim via \texttt{nvidia\_cuda\_mps\_server} with very little loss of performance.

We have also shown a preliminary Kokkos implementation that supports running on multiple architectures, including multicore CPUs, NVIDIA GPUs and AMD GPUs. Kokkos has proven to be a powerful tool for performant portability on a number of different architectures. While there is some overhead with utilizing its data structures as compared with raw memory allocation, these are relatively small and offset by the gains in portability. Direct comparisons of kernel execution time running on both NVIDIA and AMD GPUs show reasonably similar measurements between Kokkos and device specific CUDA/HIP implementations. Since the Kokkos implementation was designed as a port from CUDA, and not rewritten from the ground up, it is very likely that further optimization of the Kokkos code would be possible. The ability to run on both NVIDIA and AMD GPUs, as well as multicore CPUs via OpenMP or pThreads, with minimal changes to user code, vastly outweighs any small performance losses.

\section*{Acknowledgments}
This work was supported by the DOE HEP Center for Computational Excellence at Brookhaven National Laboratory, and Lawrence Berkeley National Laboratory under B\&R KA2401045.
 This work was also done as part of the offline software research and development programme of the ATLAS Collaboration, and we thank the collaboration for its support and cooperation.
 This research used resources of the National Energy Research Scientific Computing Center (NERSC), a U.S. Department of Energy Office of Science User Facility located at Lawrence Berkeley National Laboratory, operated under Contract No. DE-AC02-05CH11231.
 We gratefully acknowledge the use of the computing resources provided by Brookhaven National Laboratory.

redpajama_set_name: RedPajamaArXiv


Chocolate chip zucchini brownies are rich, decadent and delicious – no one will realize they are secretly healthy! Whole wheat and loaded with zucchini, these brownies make a perfect treat!
I hope you all have some fun celebrations to look forward to today, or at least later this week!
That's right, these gorgeous brownies aren't what they seem at first glance. They are loaded up with veggies! What?!
So, I've been making these chocolate chip zucchini brownies for years.
Normally, it was always to help use up some of the zucchini from my garden at our old house – which tended to produce A LOT of zucchini! You know how it is… the plant goes crazy and you are constantly uncovering more and more and more zucchini.
So I'd be making chocolate chip zucchini bread, vegetarian stuffed zucchini, harvest zucchini muffins, almond butter zucchini muffins, zucchini turkey burgers and even zucchini cheddar whole wheat waffles.
They feel festive with their little green bits tucked away in all that chocolate.
Soft, melty, festive chocolate — yup, definitely how I want to celebrate!
I don't bother peeling the zucchini. Just use the fine side of your grater.
The 2 cups of grated zucchini is measured before you squeeze the extra moisture out. Grate 2 cups then squeeze with paper towels or a kitchen towel to get the extra water out. You'll have a compact little ball of zucchini. Break it up a bit as you add it to the batter so the zucchini gets incorporated.
This size batch yields 9-16 brownies, depending on how you want to cut them. (We like big, thick brownies here!) I've also baked these in an 8×4 bread pan to have a mini batch. Alternatively, you could double the recipe and use a 9×13 baking dish.
Make these chocolate chip zucchini brownies. Soon.
You'll be getting wafts of chocolate-y goodness in no time and they'll come out warm and melty and thick and beautiful.
You'll be in chocolate heaven!
Decadent and delicious (and secretly healthy) chocolate chip zucchini brownies!
Preheat the oven to 350 and spray an 8x8 pan with cooking spray.
In a large bowl, mix the eggs, sugar, applesauce and vanilla until well combined.
In a separate bowl, mix the flour, cocoa powder, baking soda and salt until well combined.
Add the dry mix to the wet and mix until just combined. Don't over mix.
Add the zucchini and chocolate chips and gently fold in to the batter so everything is evenly distributed.
Pour into your prepared pan and level out the mix with the back of a spoon.
Bake for 30-35 minutes, until done. Check with a toothpick - the brownies are gooey but shouldn't be raw.
Once cooled, store in the refrigerator. (They keep better this way because of all the moisture from the zucchini.) I kind-of prefer them cold but you can heat in the microwave if you want them warmed up.
You can substitute regular all-purpose flour here if you prefer.
I don't bother peeling the zucchini. Just use the fine side of your grater.
The 2 cups of grated zucchini is measured before you squeeze the extra moisture out. Grate 2 cups then squeeze with paper towels or a kitchen towel to get the extra water out. You'll have a compact little ball of zucchini. Break it up a bit when you add it to the batter so the zucchini gets incorporated.
Adapted from Texan Erin. She's got tips on making this gluten-free and dairy-free too.
Zucchini really does do amazing things in baked goods – especially when paired with chocolate!!
Sneaky veg in a sweet treat – what's not to love!
These are very good. My family really liked them.
I'm SO happy to hear that Ann! Thanks so much for sharing!
Was the 2 cups of zucchini measured before or after you squeezed it out? I plan to make these soon!
Great question Cheryl! I grate a full two cups then squeeze it out (it will be a lot less and very compact). Make sure to fluff it back up a bit when you're adding it so the zucchini gets incorporated. I hope you enjoy them!
Kathryn, thank you for the quick reply – making soon!
Just put them in the oven! Double batch so I can take one to a friends house and leave some for my family. I used dark chocolate cocoa powder and they smell so good!
Aren't you a sweet friend?! Thanks so much for leaving a comment – I hope you all enjoy them!


Lester Holt received a great deal of criticism for his performance as debate moderator during Monday's first presidential debate, so Seth Meyers took him to task on Wednesday's Late Night by showing him how it's done.
Meyers used cleverly edited debate footage to get Hillary Clinton and Donald Trump to give absurd answers to his absurd questions.
Watch the video for more wacky proof that Seth Meyers should be the next debate moderator.


Steve Morse has a unique approach to playing guitar, and we designed his pickups to match his technique. Steve commonly switches between pickups mid-solo — playing the neck pickup on the high notes and the bridge pickup on the lows — so both pickups must be evenly-balanced over the entire fingerboard. We engineered the Steve Morse Model™ Neck specifically for this purpose, with a dark, smooth sound that is clean and uncompressed.
Steve's neck model has a very specific sound, which is almost exactly the opposite of his bridge pickup in every way — it has low output and a warm, very smooth, almost flat attack. The DC resistance spec is very high, but the magnet is not strong, and this combination accounts for the unusual mix of dark tone and low output.


text:
School of Languages and Linguistics
School of Languages and Linguistics - Research Publications
Pair versus individual writing: Effects on fluency, complexity and accuracy
Wigglesworth, G; Storch, N
Storch, Neomy; Wigglesworth, Gillian
Languages and Linguistics
Wigglesworth, G. & Storch, N. (2009). Pair versus individual writing: Effects on fluency, complexity and accuracy. LANGUAGE TESTING, 26 (3), pp.445-466. https://doi.org/10.1177/0265532209104670.
<jats:p> The assessment of oral language is now quite commonly done in pairs or groups, and there is a growing body of research which investigates the related issues (e.g. May, 2007). Writing generally tends to be thought of as an individual activity, although a small number of studies have documented the advantages of collaboration in writing in the second language classroom (e.g. DiCamilla & Anton, 1997; Storch, 2005; Swain & Lapkin, 1998). Particularly in university contexts, group or pair assignments are widely used in many disciplines. In addition, collaborative writing could be used in second language classroom assessment contexts as formative assessment. However, research which compares texts produced by learners collaboratively to texts produced individually, and the implications of this for assessment practices, is rare. </jats:p><jats:p> This study is a first step in the investigation of using collaborative writing in second language contexts and comparing the performance of two groups of second language learners: one group worked individually, and the other group worked in pairs. When writing in pairs, each pair produced a single text. All participants completed one writing task: an argumentative essay. The performances of the individuals (N = 48) and the pairs (N = 48) were compared on detailed discourse analytic measures of fluency, complexity and accuracy. This comparison revealed that collaboration impacted positively on accuracy, but did not affect fluency and complexity. A detailed analysis of the pair transcripts recorded during the writing activity provides insights into the ways in which pairs work together, and the foci of their endeavour. The implications of these findings for in-class assessment of second language writing are discussed. </jats:p>
School of Languages and Linguistics - Research Publications [842]

redpajama_set_name: RedPajamaCommonCrawl


"Most of our politicians do not have the details of the Indo-US nuclear deal. However, they are taking a stand.
"When I ask them if you have read the 123 agreement, please tell me what areas which you see are bad, very few of them were in a position to give facts. However, they have already taken a position. Please avoid taking decisions based on ego, based on perceptions. Go for an analysis based on data and facts. In the end, everybody is happy.
This entry was posted in Issues and Ideas, IT-BT, IIT-IIM, Not Bad News and tagged Infosys, N.R. Narayana Murthy, Nuclear Deal. Bookmark the permalink.
It was wonderful to listen to the discussion.
I normally donot see these channels, but this discussion really didnt let me move from my place for more than an hour.
Actually, I was a supporter of the Nuclear deal, but after listening to Mr.Shourie's arguments, I have a doubt about the deal itself.
The sage of Electronics City is back at doing what he is best at….carping and complaining. NRN has hardly ever been original, his comments on politics and politicians have been cliches that are repeated ad nauseum in any upper class drawing room.
But back to his point, there is nothing farther away from thr truth as his generalisation about the Nuclear Deal. I wanted to cite an example of a debate on Quality of Mineral Water Bottles supplied in Railway Stations that came up for debate a while ago in Parliament. I was surprised at the very high levels of knowledge and insight displayed by the members on such a mundane topic. I am sure on more critical issues, the preparation and knowledge is even higher. At the same time, it would be downright silly to expect ALL the MPs to have read the deal. There is always a core group in any political grouping that does the background work on specific issues.
Despite our persistent stereotyping, it is a fact that legislatures, particularly at the state level have a significant proportion of smart, intelligent & successful people who graduate to politics. There have been numerous doctors, lawyers, farm sector leaders, trade union leaders who have won elections (They may not have the normal definitions of Smart – US Degrees, Corporate style clothing, but nevertheless have high levels of perception and intelligence). That they choose money over service once in power is another matter.
I agree with NRN. Barring a few, most politicians simply nod their heads in agreement to whatever their leaders say without reading or understanding what is in the 123 agreement.
A bipartisan group of highly distinguished people have appealed to these leaders to accept the agreement and sure no agreement will be 100 % perfect.
Our politicians should be more pragmatic without being too much dogmatic.
I feel nuclear agreement is in Country's best interest.
NRN…Please stop criticizing politicians. Because you are also like one of them – Hypocrite, Non-Original.


text:
Дубровачки архив, или Државни архив у Дубровнику, државни је архив смјештен у Дубровнику у Хрватској. У свом саставу има материјал настао у државној служби Дубровачке републике, тј. нотарске и секретарске услуге из 13. вијека и након пада Републике 1808. године документа настала у службама и институцијама у граду Дубровника током француске, хабзбуршке, југословенске и хрватске власти.

Архив је важан јер је Дубровачка република имала трговинске и политичке везе са свим земљама на Средоземном мору и унутрашњости југоисточног дијела Европе; штавише, документи и записи из тог периода представља врло драгоцену историјску грађу за проучавање дубровачке историје, као и историје Хрватске и сусједних земаља.

Историја
Дубровачки архив је настао као резултат нотарских и секретарских услуга из 13. вијека, са архивским материјалом смјештеним на различитим локацијама, као што су канцеларије цивилне службе, магистрата и канцеларије изван зидина града. С обзиром да су материјали из средине 15. вијека чувани у украшеним орманима у Кнежевом двору, који и даље имају исту сврху. Други дио архивираног материјала састоји се од најважнијих привилегија и уговора склопљених између Републике и околних владара, вјерских великодостојника и градова са којима је Дубровник одржавао политичке и трговинске односе, чувани су у ризници катедрале Узнесења Блажене Дјевице Марије на небо. Поред ове двије архивске локације, Дубровачка република је имала још неколико, али су сви уништени, као што је и архив у Надбискупском двору уништен током пожара 1667. године.

Највећи дио архивираних докумената настао је радом канцеларија градске власти, чија је организација завршена током 15. вијека. Постојећа државна архива је основана као засебна установа, разним извршним одлукама највиших органа Дубровачке републике. Термин архив је први пут потврђен 1599. године, а 1760. године архивирани документи са разних локација су били смјештени у Кнежевом двору. Потврда термин државни архив датира из 1783. године. Те године је забиљежена структура архиве подијељена на 14 серија (861 томова), док су остали неувезани документи били распоређени у 10 столова, од слова А до слова К. Током 19. и 20. вијека највећи дио архивске грађе је одузет. Архив је 1920. године почео са радом као независна институција, а 1952. године је премјештен у палату Спонза, гдје се и данас налази.

Архив данас
Данашња збирка садржи 300 фондова и колекција, са више од 2,7 милиона страница писаних докумената, исправа, докумената и осталих. Најстарији документи су була папе Бенедикта VIII из 1022. године и документи о оснивању бенедиктинског манастира на острву Локруму из 1023. године, сачувани у препискама из каснијих вијекова. Главни фондови архивске грађе су из времена Републике и француске управе. У архиву се налазе записи канцеларија и установа које су постојале током 19. и 20. вијека у граду и његовој околини. Осим њих, у архиву се налазе и књиге из Дубровника и острва Корчула до 1860. године, Збирка географских карата, као и приватни породични архиви. Данас, архив складишти материјал из Дубровачко-неретванске жупаније, функциониште као спољна архивска служба. Архив је данас подијељен на сљедеће организационе јединице:
 Одјељење за заштиту и обраду архивске грађе од националног значаја — архивни фондови Дубровачке републике, француске управе и старије архивске грађе Корчуле до 1797;
 Одјељење за заштиту и обраду архивске грађе управе до 1814. до 1945;
 Одјељење за заштиту и обраду архивске грађе управе од 1945;
 Одјељење за заштиту и обраду архивске грађе правосуђа, нотаријата и тужилаштва;
 Одјељење за заштиту и обраду архивске грађе приватних и породичних фондова и збирки, културних, просвјетних, социјалних и здравствених установа, друштевних и политичких организација друштава и удружења;
 Одјељење за заштиту и обраду архивске грађе привредних, новчарских и поморских фондова;
 Одјељење за заштиту и обраду архивске и регистратурне грађе изван архива;
 Одјељење за фотографију, микрографију, репрографију и дигитализацију архивске грађе, документацијско-информацијске послове с библиотеком;
 Лабораторија за конзервацију и рестаурацију архивске грађе и
 Одјељење за правске, управне, рачуноводствене и опште послове Архива.

Види још
 Дубровачка република

Референце

Спољашње везе

Архиви
Дубровник

redpajama_set_name: RedPajamaWikipedia


text:
DESTINY: New Single Delayed
December 15, 2012 0 Comment Destiny (Band)
The new digital single from long-running Gothenburg, Sweden-based metallers DESTINY, which was originally scheduled for release this month, has been delayed until January. The single will contain material from the band's anniversary album, "The Incompatibility Of Philosophical Terminology", which is due in the spring. It will mark the recording debut of new DESTINY singer Jonas Heidgert (DRAGONLAND) and will also be the first release featuring Michael Åberg (NOSTRADAMEUS) on guitar, even though he has been with the band since 2005. Guesting on drums is Adde Larsson (M.A.N., BLACK CANDY STORE) who recorded the drums about a year ago well before current drummer Roger Christiansson joined DESTINY a couple of weeks ago. The bass will, as always, be handled by DESTINY mainman Stefan Björnshög.
Everything but the vocals has already recorded for DESTINY's next official full-length CD, "Global Warming".
Way back in 1982, bassist Stefan Björnshög founded one of Sweden's first true heavy metal bands, DESTINY. Unlike most other Swedish acts that were active around the same period, DESTINY never broke up. In 2005, DESTINY was joined by guitarist Michael Åberg, who previously played with NOSTRADAMEUS.
Following a tour opening for former BLACK SABBATH vocalist Tony Martin in 2006, DESTINY kept an unusually low profile, partly because the band was unable to find a suitable vocalist and also because of illness and various personal tragedies.
← AC/DC's Music To Accompany Drones Via Laser Beam Signal
RUDY SARZO, ZAKK WYLDE Attend 'Randy Rhoads: The Quiet Riot Years' Screening In Los Angeles →
Destiny Posts 30th Anniversary Album Teaser
Former SAXON Drummer NIGEL DURHAM Joins MORPHEUS RISING
ASKING ALEXANDRIA's 'From Death To Destiny' Lands In U.S. Top 10


You are here: Home / adventures / Detroit / Life around here.
It's time for us to catch up! :) January and February went by equally fast (so much excitement over living in Michigan again) and slow (missing Zack so, so terribly while he was finishing work in California), and suddenly we've welcomed March into our lives. Shortly after Zack arrived in Detroit (Feb. 20) our entire family was hit with the stomach flu — my second time this pregnancy (what?!), and so for an entire week we were wiped out. Like, we powered through an entire jug of laundry detergent in February ;) As soon as the weather was mild and we had our health back, I opened up the windows and cleaned our house just like a 34 week pregnant woman! It was glorious. I seriously loved being able to clean our home (and have it stay clean).
Yes, today marks 34 weeks with this pregnancy. I'm in awe over how fast it's going, and I can't wait to meet our third babe (and see if baby resembles Natalie or Max at all). While we anticipate my due date, we have been busy searching for and buying furniture on Craigslist, hanging some of our art on the walls, moving things around, and shoveling snow (just the boys). For the most part I've felt amazingly great. My legs get so sore and restless come evening time, but this isn't a new symptom. And I finally bought a body pillow — what was I waiting for? I've had terrible back pain, but it suddenly vanished with the right kind of pillow support. Max wants me to pick him up all the time (which I can't), so we've compromised by snuggling a lot on the couch and before nap/bed time. He thinks he's getting a sweet deal, but we all know mama is the one getting the reward ;) snuggling with my firstborn, it NEVER gets old!
Even though the weather (and our neighborhood) is totally different from Los Angeles, it doesn't feel like a hard adjustment. Many have asked us about this, and I'm happy to share it's been quite easy :) Zack and I both grew up in Michigan, and we lived in the Detroit metro area about eight years ago, just never with children. And, the amount of living space we now have is just enough to help us feel like we can fully create and play and rest and cook and just be.
Life around here has been full and exciting and tiring and challenging and the color of fresh snow. I'm definitely amazed at how far we've come since we last lived in Detroit, and so grateful for all the adventures we've had throughout that span of time. And let's not forget: ZACK IS HOME, ZACK IS HOME, ZACK IS HOME! :) it's the best thing ever seeing him every single day!
« Our 2016 Valentine's Day cards!
Yay your home is looking so cute, love your sons toddler bed by the way!! We need to start looking for one for our second who just turned two! I think she will be our last so it will be SO strange to not have a crib anymore! So happy things are going well for you all, and your family is reunited! Its the best having your husband to chat with at the end of the day Take care!!
I love our family and your blog so much. Stay perfect, little Queen.
OH goodness, thank you :) I suppose when I look at the big picture of our home I've gotten more done than I realize.
Yea! Hooray! Loved this little update and seeing pictures of your home! Hope you have a marvelous weekend!
Just caught up with your blog and I see how very sweet you are to mention my little kits! I'm so happy to know you are enjoying the bicycle – seriously seriously happy! And happy for you to be in enjoying a bit more space as you prepare to welcome #3. You're doing an amazing job at life right now!


Order Sildenafil Medana Low Price! Buy SILDENAFIL MEDANA Now.
sildenafil medana cost 120 mg Sildenafil Medana vigora What's more consumers can choose from a variety of products and systems. AnnA is an author and speaker on health and personal development with a particular interest in hormonal health. 7. As soon as you reach the head, bring it back again to the base. Since you're reading this article I'm guessing that you really WANT to grow yourself a more powerful package. How it WorksFirst thing that comes to mind when finding out bout enlargement product is whether they work. sildenafil medana You need to know what these are so that you can get rid or take control of them. generic sildenafil medana canada price So as you can see, it's pretty clear why many men (including myself) have chosen to go the natural route and exercise their manhood. The second commonly used method (that works) is the enlargement exercise. This will increase the level of testosterone and promote cell regeneration.


text:
Die Börse Shenzhen () ist eine der beiden Börsen der Volksrepublik China und wurde 1990 gegründet.

Nach der Börse Shanghai ist die Börse in Shenzhen die zweitgrößte Festlandchinas.

Entwicklung
Im Jahr 2014 listete die Börse Shenzhen 1581 Unternehmen, welche zusammen eine Marktkapitalisierung von 3.932 Milliarden US-Dollar aufwiesen. Bis Juli 2013 war die Zahl der Aktiengesellschaften auf 1537 gestiegen. Im August 2017 betrug die Marktkapitalisierung 3.562 Milliarden US-Dollar, womit sie die achtgrößte Börse der Welt nach Marktkapitalisierung ist. Inzwischen sind über 2000 Unternehmen gelistet.

IPOs an der Börse in Shenzhen wurden 2000 eingestellt und werden nur noch an der Börse in Shanghai gehandelt. In Shenzhen wurde ein der NASDAQ ähnlicher zweiter Markt geschaffen (ChiNext), welcher für Privat- und Technologieunternehmen gedacht ist.

Aktienindizes
Der Hauptindex SSE wurde im Mai 2004 eingeführt. Erst einige Jahre später wurden spezielle Indizes dem Listing hinzugefügt. Der ChiNext Market wurde im Oktober 2009 gegründet. Seit Januar 2006 wurde der OTC-Market-Index gelistet.

Siehe auch
 Hong Kong Stock Exchange
 Shanghai Stock Exchange
 CSI 300

Weblinks
 Offizielle Webseite der Börse Shenzhen (chinesisch)
 Offizielle Webseite der Börse Shenzhen (englisch)
  In: handelsblatt.com, 25. Januar 2006

Einzelnachweise

Wertpapierbörse
Finanzdienstleister (Volksrepublik China)
Unternehmen (Shenzhen)
!
Rem Koolhaas


We attract and retain highly motivated individuals – the best of the best in our industry.
Pharos Offshore are committed to supporting our offshore personnel and by delivering engineering solutions which will provide an operational, commercial and technical advantage to our clients, through this activity we develop new fields of business in which our offshore personnel can provide support, developing the growing needs of our client base and our internal personnel commitment.
We offer an attractive employee package with clear progression and development opportunities, health cover and flexible working hours.
We offer competitive and attractive rates of pay, in line and above industry standards.
We place emphasis on clear progression and development opportunities for staff.
We aim to create a positive and friendly working environment.
Pharos Offshore offers private health care to all staff.


text:
Charles Clary
Charles Clary (March 24, 1873 – March 24, 1931) was an American actor of the silent film era.[1] Clary appeared in 206 films between 1910 and 1930. He was born in Charleston, Illinois and died on his 58th birthday in Los Angeles, California. He worked for Selig and the Fine Arts Film Company.[2]
Charleston, Illinois, United States
March 24, 1931(1931-03-24) (aged 58)
Selected filmography
Brown of Harvard (1911)
The Carpet from Bagdad (1915)
The Penitentes (1915)
Joan the Woman (1917)
The Innocent Sinner (1917)
The Spy (1917)
The Honor System (1917)
A Tale of Two Cities (1917)
Madame Du Barry (1917)
The Silent Lie (1917)
The Soul of Satan (1917)
The Conqueror (1917)
The Rose of Blood (1917)
The Man Hunter (1919)
The Lone Star Ranger (1919)
A Girl Named Mary (1919)
The Day She Paid (1919)
The Woman in Room 13 (1920)
The Penalty (1920)
A Connecticut Yankee in King Arthur's Court (1921)
The Sea Lion (1921)
Sunset Jones (1921)
The Hole in the Wall (1921)
Two Kinds of Women (1922)
Heroes and Husbands (1922)
Money! Money! Money! (1923)
Prodigal Daughters (1923)
Thundering Dawn (1923)
On Time (1924)
In Fast Company (1924)
Empty Hands (1924)
Seven Days (1925)
The Golden Bed (1925)
Jimmie's Millions (1925)
Super Speed (1925)
An Enemy Of Men (1925)
The Unwritten Law (1925)
The Coast of Folly (1925)
The Volga Boatman (1926)
The Blind Goddess (1926)
The Blue Streak (1926)
Thrilling Youth (1926)
Beverly of Graustark (1926)
Mighty Like a Moose (1926 short)
Satan Town (1926)
Smile, Brother, Smile (1927)
The Magic Garden (1927)
See You in Jail (1927)
Pretty Clothes (1927)
A Woman Against the World (1928)
Jazz Mad (1928)
Sailor's Holiday (1929)
"Charles Clary". The New York Times. Archived from the original on March 26, 2016. Retrieved March 4, 2017.
"fine+arts+film+company"#v=onepage&q="fine%20arts%20film%20company"&f=false "The Movie Magazine: A National Motion Picture Magazine". 1915.
Wikimedia Commons has media related to Charles Clary.
Charles Clary on IMDb
Charles Clary at the Internet Broadway Database
BNF: cb14165024b (data)