2010 10th IEEE/ACM International Conference on Cluster, Cloud and Grid Computing

Cooperative Multitasking for GPU-Accelerated Grid Systems
Fumihiko Ino, Akihiro Ogita, Kentaro Oita and Kenichi Hagihara
Graduate School of Information Science and Technology
Osaka University
1-5 Yamadaoka, Suita, Osaka 565-0871, Japan
Email: ino@ist.osaka-u.ac.jp

a single kernel to run at a time and (2) switches the runningkernel only when it completes its execution. Therefore, theframe rate will signiﬁcantly drop if the guest kernel occupies
the resources for long time.
To avoid this problem, current systems [4], [5] use screensavers capable of detecting fully idle GPUs for accelerationof guest applications. In other words, both guest and hostapplications are exclusively executed in the systems, where
idle resources are typically dedicated to guest applications.
In contrast to these dedicated systems, we are focusing on
non-dedicated systems, which have a true resource sharing
mechanism to harness the power of GPUs in the homeand ofﬁce. For example, such mechanisms allow us torun guest applications on lightly loaded GPUs, where hostusers operate their machines for ofﬁce work with almost noworkload on the GPU.
Our main goal is to achieve true resource sharing between
hosts and guests for GPU-accelerated grid systems. To
achieve this, we are developing a cooperative multitasking
method capable of simultaneous execution of a graphics
application and a CUDA-based scientiﬁc application. Using
this method, screensavers are not needed to ensure exclusive
execution of guest and host applications. Instead, hosts
who accept scientiﬁc jobs from guests should specify theminimum frame rate they need. According to this requirement, the proposed method divides the workload of scientiﬁcapplications such that each kernel can complete its execution
within the desired period. The method assumes that guest
applications are implemented using CUDA while host applications are implemented using a periodical rendering modelwith graphics libraries such as DirectX [9] and OpenGL[10].
The rest of this paper is organized as follows. Section II
presents preliminaries including an overview of CUDA andthe periodical rendering model. Section III then describes thedetails of our multitasking method and Section IV shows
experimental results. Finally, Section V summarizes thepaper with future work.

Abstract—Exploiting the graphics processing unit (GPU) is
useful to obtain higher performance with a less number of hostmachines in grid systems. One problem in GPU-accelerated
grid systems is the lack of efﬁcient multitasking mechanisms.
In this paper, we propose a cooperative multitasking method
capable of simultaneous execution of a graphics application
and a CUDA-based scientiﬁc application on a single GPU. To
prevent signiﬁcant performance drop in frame rate, our method(1) divides scientiﬁc tasks into smaller subtasks and (2) serially executes them at the appropriate intervals. Experimental
results show that the proposed method is useful to control theframe rate of the graphics application and the throughput of
the scientiﬁc application. For example, matrix multiplication
can be processed at 50% of the dedicated throughput whileachieving interactive rendering at 54 frames per second.
Keywords-multitasking; GPU; CUDA; grid;

I. I NTRODUCTION
The graphics processing unit (GPU) [1], [2] is an accelerator originally designed for graphics applications. It provides
us high memory bandwidth and ﬂoating-point performancewith a single-instruction, multiple-data (SIMD) capability.
Furthermore, it now has a ﬂexible development framework,
called compute uniﬁed device architecture (CUDA) [3],demonstrating many acceleration results typically with a 10fold speedup over CPU-based implementations. In CUDA,
the compute-intensive code is usually implemented as akernel, namely a function that runs on the GPU.
The GPU is also emerging as a powerful computational resource in grid environments, where distributed resources are
virtually collected into an integrated system. For example,
the Folding@home project [4] demonstrates that 70% of theentire performance is provided by idle GPUs, which accountfor only 10% of all resources available in the system. Thus,
exploiting the GPU is useful to obtain higher performancewith a less number of host machines in grid systems. In thispaper, we denote hosts as users who donate their resources
to the grid system. On the other hand, guests are grid users
who submit jobs to the system.
One problem in GPU-accelerated grid systems [5], [6] isthe lack of efﬁcient multitasking mechanisms. For example,
the frame rate of graphics applications and the throughput of
scientiﬁc applications can signiﬁcantly drop when they are
simultaneously executed on the same GPU [7], [8]. This isdue to the current GPU architecture, which (1) allows only
978-0-7695-4039-9/10 $26.00 © 2010 IEEE
DOI 10.1109/CCGRID.2010.18

II. P RELIMINARIESThis section presents preliminaries needed to understand
our method.
774

1/ F

Grid
Generation

TB

TB

……

TB

TB

TB

……

TB

1/ F

t

t

t

Frame #1

Frame #2

Frame #3

KernelinvocationTime
Thread block (TB)
Thread

TB

TB

……

TB

TB

TB

……

TB

ThreadFigure 2. Periodical rendering model. Rendering tasks are executed atregular intervals to produce a series of frames.

…

Generation
Kernelinvocation

…

Grid

… Thread… Thread

PC games and movie players are required to ensure thesame rate to avoid extremely high frame rates per second(fps). In contrast, the latter executes rendering tasks on the
GPU as fast as possible. A rendering task here corresponds
to the drawing code ﬂushed by the glFlush() function,
which produces a single frame in graphics applications.
On the other hand, a CUDA task in scientiﬁc applications
corresponds to a kernel invocation. As we mentioned in
Section I, our method assumes that host applications are
implemented using the periodical model.
Figure 2 shows how the periodical model produces aseries of frames. In this model, rendering tasks are executedat regular intervals. The intervals are given by 1/F , where F
denotes the frame rate needed for the graphics application.
Let t be the time spent for producing a frame. The idle
period W between frames then can be represented by

Time
Figure 1. Hierarchical thread structure in CUDA. The CPU code invokes
kernels, which generate thousands of threads to accelerate computationon the GPU. An implicit global synchronization is performed between
successive invocations.

A. Compute Uniﬁed Device Architecture (CUDA)
CUDA [3] is a ﬂexible development framework for the
NVIDIA GPU. Since it is based on an extension of the C
language, it allows us to easily implement GPU-accelerated
applications, which consist of kernels and CPU code.
CUDA-based kernels typically generate millions of threads
on the GPU, which then accelerate heavy computation by
SIMD instructions on multiprocessors (MPs). Currently, thenumber M of MPs ranges from 16 to 30 depending on theGPU.
Figure 1 illustrates an overview of the hierarchical thread
structure in CUDA. In CUDA programs, threads are classiﬁed into thousands of groups, each called as thread blocks(TBs). Threads belonging to the same TB are allowed tosynchronize each other, so that such threads can share datausing fast on-chip memory, called shared memory. On theother hand, data dependencies between different TBs are notallowed in the kernel. Therefore, we have to separate thekernel into multiple pieces to deal with such dependencies.
Separated kernels are then serially executed with global
synchronization.
TB is the minimum unit allocated to MPs. Therefore,
the number n of TBs should be a multiplier of M forload balancing between MPs. It also should be noted herethat the GPU architecture is designed to hide the memory
latency. This is achieved by concurrently running multiple
TBs on each MP. Since there is no data dependence between
different TBs, each MP is allowed to switch them to perform
computation during a memory fetch operation. To maximize
the effects of this latency hiding, MPs run more TBs as longas registers and shared memory are available.

W = 1/F − t.

(1)

According to the model mentioned above, the frame rate
F can be decreased if W < 0. One problem here is thatthe time t cannot be directly measured in grid environments
because it is not realistic to obtain and modify the sourcecode of arbitrary host applications. It is important to take thisconstraint into account when developing grid middleware.
III. P ROPOSED M ULTITASKING M ETHOD
Figure 3 illustrates how guest and host applications are
executed under the periodical rendering model. According
to our preliminary experiments, we ﬁnd that rendering tasks
and CUDA tasks are always executed in exclusive mode.
That is, the GPU switches tasks when the current kernel
completes its execution. Therefore, rendering tasks can bedelayed or cancelled if a scientiﬁc task occupies the GPU
for long time, as shown in Fig. 3(a).
To solve this problem, we have to control the execution
time of guest kernels such that the host application canprocess a rendering task at every time 1/F . Figure 3(b)illustrates how our cooperative method realize efﬁcient multitasking on the GPU. The proposed method consists of two
strategies as follows.
1) Task division. The method divides guest tasks into
smaller subtasks such that each kernel completes its
execution within the idle period W .
2) Alternative execution. The method invokes the scientiﬁc kernel at almost the same intervals as the host

B. Periodical Rendering Model
There are two typical rendering models that can be
employed for graphics applications: a periodical model anda non-periodical model. The former is intended to provide
the same frame rate on any graphics card. For example,
775

1/ F
Host
process

Frame #1

OriginalMatrixMultiplication()
1: // setup execution parameters
2: dim3 threads(BLOCK SIZE, BLOCK SIZE);
3: dim3 grid(WC / threads.x, HC / threads.y);
4: // execute the kernel
5: matrixMul<<< grid, threads >>>(d C, d A, d B, WA, WB);

Frame #2

GPU
Time
Guest
process

global voidmarixMul(ﬂoat *C, ﬂoat *A, ﬂoat *B, int wA, int wB) {6:// Block index
7:
int bx = blockIdx.x;
8:
int by = blockIdx.y;
9:... // omitted
10: }Original task

(a)
1/ F
Host
process

Frame #2

Frame #1(a)GPU

ModiﬁedMatrixMultiplication(GRID YSIZE)
1: // setup execution parameters
2: dim3 threads(BLOCK SIZE, BLOCK SIZE);
3: dim3 grid(WC / threads.x, GRID YSIZE);
4: // execute the kernel
5: for (int i=0; i<HC / threads.y / GRID YSIZE; i++) {6:
matrixMul<<< grid, threads >>>(d C, d A, d B, WA, WB, i*GRID YSIZE);
7:
cudaThreadSynchronize();8:
Sleep(1/F );
9: }

Time
Guest
process

Subtask #1

Subtask #2

1/ F
Sleep()(b)
Figure 3.
Timeline view of multi-task execution on the GPU. (a) A
naive method will drop the frame rate but (b) our method prevents thisperformance degradation. The proposed method divides each of scientiﬁctasks into smaller subtasks, which are then serially processed at regularintervals.

global voidmarixMul(ﬂoat *C, ﬂoat *A, ﬂoat *B, int wA, int wB,
int offset) {10:// Block index
11:
int bx = blockIdx.x;
12:
int by = offset + blockIdx.y;
13:... // omitted
14: }application. This prevents resource starvation because
host and guest applications are given equal chances to
use the GPU.
Note that our method requires code modiﬁcations of guest
applications. In contrast, there is no need to modify thecode of graphics applications that run as host applications
on grid resources. This is essential to run the method in
grid environments, where there can be a large number ofhost applications and their code is not allowed to edit. In
this sense, the method provides a realistic solution to the
problem of multitasking in grid environments.

(b)
Figure 4. Example of guest code modiﬁcations. (a) The original code[3] for matrix multiplication C = AB and (b) the modiﬁed code forcooperative multitasking. WC and HC represent the width and the heightof matrix C. The input parameter GRID YSIZE determines the number ofTBs.

almost all of the original kernel code can be reused aftertask division. The modiﬁcation only needed is that we haveto add an offset as a kernel argument and have to specify
the appropriate address of output data by using the offset.
Figure 4 shows an example of code modiﬁcations. It
explains how matrix multiplication C = AB can beadapted to cooperative multitasking. In this example, theoriginal code [3] in Fig. 4(a) generates (WC/BLOCK SIZE)× (HC/BLOCK SIZE) TBs, where BLOCK SIZE is thesize of TBs, and WC and HC represent the width
and the height of matrix C. On the other hand, themodiﬁed code in Fig. 4(b) reduces this number to(WC/BLOCK SIZE) × GRID YSIZE TBs, as shown inline 3, where 1 ≤ GRID YSIZE ≤ HC/BLOCK SIZE.
Parameter GRID YSIZE here represents the number of TBs
in y direction. This parameter can be given to the CPU code
to change the granularity of subtasks at run-time.
Let K (> 0) and k (≤ K) be the execution time of theoriginal task and that of a divided subtask, respectively. In

A. Task Division
We now explain how task division can be done for CUDAbased applications. As we mentioned in Section II-A, there isno data dependence between different TBs. Therefore, ourmethod divides the original task into smaller subtasks bysimply reducing the number n of TBs given to the CUDA
kernel. In addition to this task division, it serially invokes
the kernel with changing TBs to obtain the same results as
before.
Our division strategy has an advantage in code modiﬁcation. Firstly, we use the same TB size as the originalcode. This means that kernel optimization inherent in TBs
is kept as the original. For example, we do not have toconcern about reducing the degree of parallelism available
in each TB. One exception is that the memory latency
hiding mentioned in Section II-A can be cancelled due tothe reduced number of TBs. The second advantage is that776

general, the kernel workload is proportional to the number
n of TBs. Therefore, our method assumes that the execution
time K can be represented by
K = Bdn/M e,

Table IS PECIFICATION OF EXPERIMENTAL MACHINES .
ItemOperating system
CPU
GPU
CUDA
DriverDesktop resolution(2)

where B represents the time needed for processing a single
TB on an MP. In optimized kernels, we can assume thatn À M and n ≡ 0 (mod M ). Suppose that the original
kernel takes long time such that K > W . In this case, we
divide the original task into dK/W e subtasks in order tosatisfy k ≤ W . Notice that this estimation is not precise
because MPs can run multiple TBs at a time.

Machine #1
Machine #2
Windows XP
Windows 7
Core 2 Duo
Xeon W3520
GeForce 8800 GTS (G80) GeForce GTX 280
1.1
2.3
169.21
191.071280 × 1024 pixels
1920 × 1080 pixels

IV. E XPERIMENTAL R ESULTS
We now show experimental results to understand theeffects of the proposed method. For experiments, we used
two desktop PCs, as shown in Table I. One is equippedwith an Intel Core 2 Duo CPU running at 1.86 GHz. Thismachine has an NVIDIA GeForce 8800 GTS (G80) cardwith M = 12. We have installed Windows XP, CUDA 1.1,
and graphics driver 169.21. The other one has an Intel Xeon
W3520 CPU running at 2.66 GHz. This machine has an
NVIDIA GeForce GTX 280 card with M = 30. We haveinstalled Windows 7, CUDA 2.3, and graphics driver 191.07.
With respect to guest applications, we used two applications. One is matrix multiplication [3] and the otheris biological sequence alignment [12]. The former solves
the problem with the matrix size of 2048 × 2048. Duringexecution, the kernel generates n = 16, 384 TBs, each consisting of 16 × 16 threads. The latter implements the SmithWaterman algorithm [13] to perform sequence alignment
between a database of 250,143 entries and a query sequenceof length 512. The kernel generates n = 250, 143 TBs withthread block size 128. Both guest applications are manually
modiﬁed for the proposed method.
On the other hand, a phong shader [14] is employed asa host application. This shader is implemented using the
OpenGL library [10] with the periodical rendering model.
Since the shader runs at F = 60 fps, the sleeping time 1/F
is set as 17 ms in experiments.

B. Alternative Execution
Although each of subtasks completes its execution within
the idle period W , the frame rate of the graphics application
can drop if guest subtasks are continuously executed on theGPU. To prevent guest subtasks from occupying the GPU,
our method tries to ensure that at least a rendering task isprocessed between successive guest subtasks.
Such alternative execution requires synchronization between host and guest applications, because they are independently executed in grid environments. However, it is
not realistic to develop a synchronization mechanism forarbitrary combinations of graphics applications and scientiﬁcapplications. Therefore, our method invokes the guest kernel
at almost the same intervals as the graphics application. Thiscan be simply realized by calling a sleep function between
guest kernel calls, as shown at line 8 in Fig. 4(b). The sleep
function then sleeps time 1/F , so that at least a frame will be
produced before the next call of the guest kernel, as shown in
Fig. 3(b). Note that we must call cudaThreadSynchronize()
before calling the sleep function because CUDA kernels are
currently launched in an asynchronous, non-blocking mode[3]. Otherwise, CUDA kernels can be continuously executed
between successive frames.
For the sleep function, we currently use Sleep() provided
by Windows API [11]. This has an advantage over a naive
implementation that enters a busy loop because Sleep()
allows the guest process to move to the waiting state.
However, we need an accurate sleep mechanism with 1-ms
resolution to deal with graphics applications with a higher
frame rate F ranging from 30 fps to 60 fps. On the other
hand, the resolution of Sleep() depends on that of hardware
timer and the time slice of operating system. For example,
Windows XP has the default value of 15 ms if it runs on
multiple CPUs. To obtain an accurate sleep, our method
increases the rate of context switches by altering the timequantum from 15 ms to 1 ms. This alternation can be doneusing timeBeginPeriod() and will be done if and only ifguest applications are allocated to the GPU. Due to the same
reason, some PC games might change the time quantum
when they are executed as host applications.

A. Overhead of Task Division
We ﬁrst conﬁrm that our task division strategy controls theexecution time k of a subtask. Table II shows the measured
time k of matrix multiplication with varying the number d oftask divisions. We can see that the time k varies according
to the number d of task divisions, i.e., the number n of
TBs. For example, the original kernel takes K = 299.9 ms
to complete matrix multiplication but the execution time is
reduced to 2.4 ms if the task is divided into 128 subtasks(d = 128). Furthermore, the time k is proportional to thenumber n, as we modeled in Eq. (2). Therefore, we caneasily control the time k if the original time K = 299.9 isgiven to the grid system.
With respect to the overhead of task division, the overhead
reveals when d = 128. In this case, the effective performance
reduces from 57.3 GFLOPS to 55.9 GFLOPS, which is
777

70

Table IIE XECUTION TIME OF MATRIX MULTIPLICATION WITH DIFFERENTNUMBERS OF TASK DIVISIONS . P ERFORMANCE IS MEASURED USINGMACHINE #1.

1
w/ sleep

w/o sleep

Throughput
0.9

60
0.8

k: kernel time (ms)Measured Estimated
299.9
299.9
75.0
75.0
37.5
37.5
18.8
18.7
9.4
9.4
4.8
4.7
2.4
2.3

Performance (GFLOPS)
Kernel Kernel w/ wait
57.3
57.3
57.3
44.5
57.3
37.9
57.1
28.9
57.1
31.6
57.1
15.8
55.9
7.9

0.7
0.6

40

0.5
30

0.4

Kernel thgrouhput

d: # of task
divisions
1 (original)
4
8
16
32
64
128

Frame rate (fps)

50

0.3

20

0.2
10
0.1
0

0
2.4

4.8

9.4

18.8

37.5

75

Execution time k per kernel launch (ms)equivalent to 2.4% performance drop. The effective performance here is given by 2N 3 /dk, where 2N 3 represents
the number of ﬂoating-point operations needed for matrix
multiplication of size N . Thus, the overhead is small for
matrix multiplication. Note here that the entire guest performance can be further reduced from this value due to thesleeping time 1/F . Table II also shows another effective
performance, 2N 3 /d(k + w), which explains the impact of
this waiting overhead. For example, the entire performance
results in 7.9 GFLOPS when d = 128 though the kernel
performance itself reaches 55.9 GFLOPS.
In summary, our task division strategy is useful to control
the execution time k of a subtask. It has a lower overhead but
the waiting overhead dw will reduce the entire performanceof guest applications.

Figure 5. Frame rate of phong shader and throughput of matrix multiplication with different task granularities. Results are obtained on machine #1
and are shown in average.

sleep function plays an important role in achieving smoothrendering for host applications.
Finally, Fig. 6 shows the frame rate of the phong shader,the relative throughput of matrix multiplication and that of
biological sequence alignment. These results are measured
on machine #2. There are two differences in Fig. 6(a), as
compared with results on machine #1 (Fig. 5). Firstly, weobserve lower frame rates on machine #2 though it has
higher performance than machine #1. For example, the frame
rate at k = 24 is approximately 30 fps in Fig. 6(a), which
is 44% lower than that at k = 18.8 in Fig. 5. Secondly, theeffects of the sleep function is reduced when k ≥ 24. We
think that these differences are due to the difference between
Windows XP and Windows 7. The latter has a hardwarebased graphical user interface (GUI) called Windows Aero.
This GUI is implemented using the DirectX graphics library[9]. Therefore, there are two host applications on machine#2: the phong shader and the Windows Aero.
We also ﬁnd similar results for biological sequence alignment, as shown in Fig. 6(b). As compared with matrix
multiplication results, we obtain slightly higher frame rates
when running this biological application. It differs frommatrix multiplication in that the performance is limited by
the instruction issue rate rather than the memory bandwidth.
Therefore, the frame rate can be increased if host and guest
application have different performance bottlenecks.

B. Performance of Multitasking
Figure 5 shows the frame rate of the phong shader and therelative throughput of matrix multiplication, explaining how
the host and guest application performance vary according
to the execution time k per kernel invocation, i.e., thenumber d of task divisions. The relative throughput of 1.0
here corresponds to the maximum performance measured
on dedicated machine #1. The frame rate is shown inaverage. Obviously, there is a tradeoff relation between thehost performance and the guest throughput. For example,
the host performance will be maximized when the task isdecomposed into many subtasks. However, we can see thata throughput of 0.35 can be achieved without degrading therendering performance. Furthermore, the throughput reaches
0.5 if resource owners accept 10% performance loss (54fps, in this case). Thus, we can obtain 50% of dedicatedperformance in a non-dedicated environment.
We also investigated the effect of the alternative execution
strategy, as shown in Fig. 5. We obtain higher, stable frame
rates by calling the sleep function. For example, the frame
rate reaches 54 fps when k = 18.8 but it reduces to 26 fps
if we do not call the sleep function. Accordingly, the kernel
throughput increases from 0.5 but the frame rate becomes
unstable. For example, the rate ranges from 51 fps to 57 fps
if we call the sleep function. In contrast, it ranges from 23
fps to 39 fps if we do not call the function. In this sense, theV. C ONCLUSION
We have presented a cooperative multitasking method
capable of simultaneously running a graphics kernel and a
CUDA kernel on a single GPU. In order to control the frame
rate of the graphics application, our method divides CUDA
tasks into smaller subtasks such that each subtask can becompleted within an idle period. Furthermore, the method
calls a sleep function for every kernel invocation to avoid

778

70

70

1
w/ sleep

w/o sleep

1
w/ sleepThroughput

w/o sleep

Throughput
0.9

0.9

60

60

0.8

0.8

50

40

0.5
30

0.4

Frame rate (fps)

0.6

0.3

20

0.7
0.6

40

0.5
30

0.4

Kernel thgrouhput

0.7
Kernel thgrouhput

Frame rate (fps)

50

0.3

20

0.2

0.2

10

10

0.1

0.1
0

0

0
3

6

12

24

48

0
3

96

6

12

24

48

Execution time k per kernel launch (ms)

Execution time k per kernel launch (ms)

(a)

(b)

96

Figure 6. Frame rate of phong shader and throughput of guest applications with different task granularities. (a) matrix multiplication and (b) biologicalsequence alignment. Results are obtained on machine #2 and are shown in average.

resource starvation due to continuous execution of CUDA
kernels.
In experiments, we have shown that the method successfully controls the frame rate of host applications and thethroughput of guest applications. Our multitasking execution
achieves 35% of guest throughput as compared with exclusive execution. This throughput is achieved without dropping
the original frame rate of 60 fps. The throughput increases
to 50% if host users accept 10% frame loss.
One future work is to extend the method for nonperiodical applications, which dynamically vary the framerate.

[6] F. Ino, Y. Kotani, Y. Munekawa, and K. Hagihara, “Harnessing the power of idle GPUs for acceleration of biologicalsequence alignment,” Parallel Processing Letters, vol. 19,
no. 4, pp. 513–533, Dec. 2009.
[7] Y. Kotani, F. Ino, and K. Hagihara, “A resource selection
method for cycle stealing in the GPU grid,” in Proc. 4th Int’l
Symp. Parallel and Distributed Processing and Applications
Workshops (ISPA’06 Workshops), Dec. 2006, pp. 939–950.
[8] S. Yamagiwa and K. Wada, “Performance study of interference on sharing GPU and CPU resources with multiple
applications,” in Proc. 11th Workshop Advances in Paralleland Distributed Computational Models (APDCM’09), May
2009, 8 pages (CD-ROM).

ACKNOWLEDGMENT[9] Microsoft
Corporation,“DirectX,”//www.microsoft.com/directx/.

This work was partly supported by JSPS Grant-in-Aid
for Scientiﬁc Research (A)(20240002), Young Researchers(B)(19700061), and the Global COE Program “in silico
medicine” at Osaka University.

2007,

http:

[10] D. Shreiner, M. Woo, J. Neider, and T. Davis, OpenGL
Programming Guide, 5th ed. Reading, MA: Addison-Wesley,Aug. 2005.
[11] Microsoft Corporation, “Windows API,” 2009. [Online]. Available: http://msdn.microsoft.com/en-us/library/
cc433218(VS.85).aspx

R EFERENCES[1] J. D. Owens, M. Houston, D. Luebke, S. Green, J. E. Stone,
and J. C. Phillips, “GPU computing,” Proceedings of theIEEE, vol. 96, no. 5, pp. 879–899, May 2008.

[12] Y. Munekawa, F. Ino, and K. Hagihara, “Design and implementation of the Smith-Waterman algorithm on the CUDAcompatible GPU,” in Proc. 8th IEEE Int’l Conf. Bioinformatics and Bioengineering (BIBE’08), Oct. 2008, 6 pages(CD-ROM).

[2] E. Lindholm, J. Nickolls, S. Oberman, and J. Montrym,“NVIDIA tesla: A uniﬁed graphics and computing architecture,” IEEE Micro, vol. 28, no. 2, pp. 39–55, Mar. 2008.

[13] T. F. Smith and M. S. Waterman, “Identiﬁcation of common
molecular subsequences,” J. Molecular Biology, vol. 147, pp.
195–197, 1981.

[3] nVIDIA Corporation, “CUDA Programming Guide Version
2.3,” Jul. 2009, http://developer.nvidia.com/cuda/.
[4] The Folding@Home Project, “Folding@home distributed
computing,” 2008, http://folding.stanford.edu/.

[14] T. Teranishi, “Phong shading sample program,” 2003, http://www.asahi-net.or.jp/∼yw3t-trns/opengl/samples/fshphong/.

[5] Y. Kotani, F. Ino, and K. Hagihara, “A resource selection
system for cycle stealing in GPU grids,” J. Grid Computing,vol. 6, no. 4, pp. 399–416, Dec. 2008.

779