Abstract-The main objective of this paper is to speed up the simulation performance of SystemC designs at the RTL abstraction level by exploiting the high degree of parallelism afforded by today's general purpose graphics processors (GPGPUs). Our approach parallelizes SystemC's discrete-event simulation (DES) on GPGPUs by transforming the model of computation of DES into a model of concurrent threads that synchronize as and when necessary. Unlike the cooperative threading model employed in the SystemC reference implementation, our threading model is capable of executing in parallel on the large number of simple processing units available on GPUs. Our simulation infrastructure is called SCGPSim 1 and it includes a source-to-source (S2S) translator to transform synthesizable SystemC models into parallelly executable programs targeting an NVIDIA GPU. The translator retains the simulation semantics of the original designs by applying semantics preserving transformations. The resulting transformed models mapped onto the massively parallel architecture of GPUs improve simulation efficiency quite substantially. Preliminary experiments with varying-sized examples such as AES, ALU, and FIR have shown simulation speed-ups ranging from 30x to 100x. Considering that our transformations are not yet optimized, we believe that optimizing them will improve the simulation performance even further. SystemC, a modeling and simulation language, has been used in making early trade-off analysis and design-space exploration. It has been shown that SystemC methodologies shorten design cycles allowing designers to meet the ever tightening time-to-market constraints. A key contributing factor to SystemC's success is its ability to provide mechanisms for fast simulation. However, SystemC's reference implementation makes no attempts to leverage the parallel architectures that are becoming increasing popular. The OSCI implementation of SystemC is built using application-level threading libraries (co-operative threads) that results in an unparallelizable simulation kernel, as co-operative threads cannot be scheduled by the operating system and hence cannot be dispatched to different processing elements. Clearly, this implementation is not designed to take advantage of multi-core processing systems. A significant amount of effort has gone into improving the performance of SystemC simulations [9, 7, 6 ], but only a handful of attempts into parallelizing them [2, 8] . This paper presents a source-to-source (S2S) translator that transforms synthesizable SystemC designs into parallelly exe-1 This work was partially funded by an Air Force Labs, an AFOSR funding and NSF-CPA funding. cutable CUDA programs that run on NVIDIA Graphics Processing Units (GPUs). To convert the SystemC designs into CUDA designs, we use a programming model where cooperative threads can be mapped to parallel threads, and they can be efficiently simulated providing an order of magnitude decrease in simulation times under the constraint of preserving the existing SystemC simulation semantics. Note that we use semantic-preserving transformations to obtain the exact same behaviors as when using SystemC's discrete-event scheduler. Our translator accepts synthesizable SystemC designs; however, we require modules in the designs to have registered outputs. As a result, we do not currently support combinatorial paths between modules. We use NVIDIA TESLA D870 GPU as an example target platform. This paper is organized as follows, Section 2 presents related work. In Section 3 we present a brief introduction on the SystemC simulator. Section 4 presents the GPU architecture and the CUDA programming model. In Section 5 we explain the conversion procedure from SystemC to CUDA and the translator that does this automatic conversion. We present our results in Section 6 and conclude with future work in the last section. Several attempts have been made in the past on parallelizing a SystemC simulator, though none of them are targeted at a massively parallel platform such as a GPGPU. In [9] Savoiu et al. show a static transformation algorithm for SystemC models. This work tries to improve simulation efficiency by reducing synchronization overheads by changing the structure of the model completely to fuse together the various modules and cooperative threads into a much higher granularity thread. In [2], Kaouane et al. have proposed a design flow for a CELL processor based SystemC simulation. With limited number of synergistic processing elements, it requires a lot of context switches that reduces the efficiency. This implementation also has several constraints; for example, the processes are sensitive to a single signal, which is very rare in hardware modeling. In addition, they manually divided SystemC code to extract compute intensive and control intensive parts. Although much work has been done in parallel discrete-event simulation, to the best of our knowledge no true parallel implementation of SystemC exists. In [8] Perumalla evaluated discrete-event simulation on GPUs, but it yielded no real speed up. In fact, the simulation time was more than what CPU would take without the GPUs. Naguib et al. [3] showed that by means of process splitting, one can speed up the SystemC simulations by automatically optimizing the model for simulation with results that showed 978-1-4244-5767-0/10/$26.00 2010 IEEE 2C-2 149 15% improvement. Our approach shows a significant simulation performance improvement over existing approaches. We implement a GPU based cycle-accurate simulator by automatically carrying out a source-to-source transformation on SystemC designs to CUDA and executing it efficiently on the GPU. SystemC is an open-source system-level design language based on C++ that has its own simulation kernel. The SystemC simulation kernel is a run-time scheduler that handles both the synchronization and scheduling of concurrent processes. It implements a discrete-event scheduler that executes processes in response to occurrence of events. When an event occurs, the corresponding process is pushed into the runnable processes queue for execution. The scheduler maintains a queue of runnable processes, and it keeps executing them until they complete execution or reach a suspension wait() statement. Execute all processes in random way Figure 1 shows the flow of SystemC simulation kernel. In the Initialization phase all processes are executed in an unspecified order. In the Evaluation phase, runnable processes are selected and executed until no runnable processes exist. This can cause additional processes to become ready to run in this phase due to immediate event notifications. When no such ready to run processes exist, the simulator enters the Update phase, where signal values are updated to the values computed in the evaluate phase. At this point, if there are any pending delayed notifications, the simulator enters Delayed notification phase that determines which processes are ready to run due to the delayed notifications and returns to the Evaluate phase. Otherwise, if there are timed notifications the simulator enters the Timed notification phase where it advances the current simulation time to the earliest pending timed notification and reenters the Evaluate phase. If there are no timed notifications, the simulation is finished. While this approach works perfectly fine for a single-core processor, it does not scale and take advantage of multi-core/multi-processor environment. IV THE CUDA PROGRAMMING MODEL CUDA (Compute Unified Device Architecture) is an extension to the C language that exploits the processing power of GPUs to solve complex compute-intensive problems efficiently. High performance is achieved by launching a number of threads and making each thread execute a part of the application in parallel. These threads are grouped and arranged as blocks. Each block contains maximum of 512 threads. These blocks are arranged in 2 dimensional matrix structures to form a grid. Each grid consists of a maximum of 65536 blocks. Figure 2 illustrates the structure. Threads from the same block access fast shared on-chip memory and can be synchronized using built-in functions. Threads from different blocks can communicate data using the slower global memory. CUDA thread execution differs from that of CPU thread execution in terms of scheduling of threads. In CUDA, threads are grouped to form warps and these warps are executed in single-instruction multiple-thread (SIMT) way. All threads in a warp execute one instruction at a time. If threads that belong to a warp diverge at a data-dependent conditional branch, then they are executed sequentially and they all converge on to the same execution path in later instructions. Branch divergence applies only for threads of a particular warp; threads from different warps execute independently on different multiprocessors. We have to create threads that belong to different warps to get the maximum speed up. HOST DEVICE The bottleneck in the CUDA programming model is the data transfer rate between the Central Processing Unit (CPU) and the Graphical Processing Unit (GPU). Currently this is limited by the maximum speed of the PCIe bus bandwidth. The overall performance of the applications depend significantly on the strategy of use of the on-chip GPU memory. Frequent memory transfers between CPU and GPU deteriorates the performance. This does not hamper efficiency in our case as we do not communicate with GPU frequently. V SYSTEMC TO CUDA A general SystemC design consists of producer/stimulus, consumer/response and channel/computation modules. The channel consists of numerous modules which model the behavior of the hardware. There may be one or many producers and consumers. Figure 3 shows the execution model of "N" runnable processes with one producer and one consumer process. Simulation of the design continues until all the processes get suspended, i.e. the end of simulation condition has been met. Threads created during SystemC simulation are cooperative threads and they effectively execute sequentially even in a multiprocessor environment. Figure 4 shows the proposed 2C-2 150 parallel execution model of the SystemC design. Individual threads are executed in parallel and barriers are used for synchronization between all the threads. To prevent the race condition on shared variables we use "Double buffering" on such variables. In Figure 4 , the threads execute in the vertical direction and the data flows in the horizontal direction. Threads are colored grey for certain part during the execution cycle. In the grey colored region, that particular thread does meaningful computation, otherwise it just waits at the barrier for the other threads to complete their job. At the start of the simulation cycle, all the threads start executing in parallel, but only Stimulus thread does meaningful work whereas the rest of the threads hit Barrier-1 and wait for synchronization with other threads. This is Barrier Synchronization. After Barrier-1, Stimulus thread just waits at Barrier-2, while the threads 2 to N-1 execute and reach Barrier-2. The Consumer thread which was waiting at Barrier-2, now does meaningful work by saving the current cycle output. All the shared variables have 2 copies (Double buffering technique) of themselves. At the end of Barrier-2 we make the contents of both the copies of variables same, so that next new data is read at the next cycle. Barrier-1 ensures that input is available before the channel operation is started, whereas Barrier-2 ensures the outputs of current cycle are stored and the double buffers have same content in them. Concurrency is extracted by executing the threads 2 to N-1 in parallel. In this section we describe the structure of the CUDA program used in our implementation. In CUDA architecture all the threads can see the whole KERNEL code. The only way to make each thread execute a part of the whole code is by distinguishing that part of code with a threadId. Listing-1 shows a CUDA KERNEL CODE structure with four threads executing different parts of the whole code for MAX CYCLES number of times. The in-built function syncthreads() is used for synchronization between the various threads. Figure 5 shows the overall design flow of SCGPSim. We divide the entire process of SystemC simulation on CUDA into 5 different phases as explained below in detail.  / Code t h a t has t o be e x e c u t e d by t h r e a d 3 } __syncthreads ( ) ; / / B a r r i e r S y n c h r o n i z a t i o n } } C Design Flow SystemC design
The authors of REF use general purpose graphics processors (GPGPUs) for kernel parallelization.
SCGPSim: A fast SystemC simulator on GPUs
2010 15th Asia and South Pacific Design Automation Conference (ASP-DAC)
2010 15th Asia and South Pacific Design Automation Conference (ASP-DAC)
Computer Science