Zhan Shi

Kai Lu

Xiaoping Wang
email: xiaopingwang@nudt.edu.cn

Wenzhe Zhang
email: zhangwenzhe@nudt.edu.cn459046923@qq.com

Yiqi Wang

Application-Based Coarse-Grained Incremental Checkpointing Based on Non-volatile Memory

des établissements d'enseignement et de recherche français ou étrangers, des laboratoires publics ou privés.

Introduction

As more of high performance computing systems are being deployed in practice, faulttolerance and self-healing are becoming increasingly important. Large-scale machines are often used by the scientific community to solve computationally intensive applications typically running for days or even months. Fault-tolerance is also a musthave characteristic of large-scale software system used in safety-critical and businesscritical domains. However, large systems exponentially increase the total number of failure points and suffer [START_REF] Reed | High-End Computing: The Challenge of Scale[END_REF][START_REF] Schroeder | Understanding failures in petascale computers[END_REF]. For instance, if mean-time-to-failure (MTTF) of an individual node is as long as ten years, the MTTF of 64,000 node system is expected to decrease to the order of tens of minutes [START_REF] Plank | Libckpt: Transparent checkpointing under Unix[END_REF][START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF]. Therefore, it is critical to tolerate hardware and software failures with low impact on application performance and avoid re-computation.

Checkpoint-restart is an effective mechanism for preventing applications from restarting from scratch when unexpected system failure or scheduled maintenance occurs [START_REF] Plank | Compressed differences: An algorithm for fast incremental checkpointing[END_REF][START_REF] Research | A checkpointing library for Intel Paragon[END_REF][START_REF] Plank | Libckpt: Transparent checkpointing under Unix[END_REF]. A checkpoint is a snapshot of application state in non-volatile devices used for restarting execution after failure. However, the scaling of computing system puts additional pressure on checkpoint mechanism. On the one hand, checkpoint has to be taken more frequently [START_REF] Vaidya | Impact of checkpoint latency on overhead ratio of a checkpointing scheme[END_REF][START_REF] Plank | Experimental assesment of workstation failures and their impact on checkpointing systems[END_REF] relative to the failure rate of the system which, in turn, directly impacts the application execution time and non-volatile storage requirements. On the other hand, the overhead of global checkpoint increases due to the large checkpoint size.

With the emerging non-volatile memory technologies, fast in-memory checkpoint helps alleviate the data transfer bandwidth problem, which is recognized as a major bottleneck of checkpointing. Moreover, in-memory approach allows frequent checkpointing to meet the requirements of large-scale systems. The presence of nonvolatile memory provides opportunities for incremental checkpointing technique. Since the in-memory checkpointing makes it possible to take checkpoints every few seconds, it reduces the overhead of incremental checkpointing. As the checkpoint interval decreases, a small number of memory blocks are modified, hence, the average size of an incremental checkpoint decreases. Recent work [START_REF] Plank | Compressed differences: An algorithm for fast incremental checkpointing[END_REF][START_REF] Research | A checkpointing library for Intel Paragon[END_REF][START_REF] Plank | Libckpt: Transparent checkpointing under Unix[END_REF] has demonstrated that the incremental checkpointing shows its advantage over full-size checkpointing when the checkpoint interval is within one minute.

To employ an incremental checkpointing technique, dirty page management is required for every page in the DRAM. Page-based incremental checkpointing techniques leverage the hardware page protection to track dirty pages and back up the dirty pages at the end of checkpoint interval [START_REF] Koltsidas | PSS: A prototype storage subsystem based on PCM[END_REF]. Generally, the overhead of incremental checkpointing is made up of two parts, the monitoring overhead and the copying overhead. These two parts are a pair of trade-off factors and the key to balance the tradeoff is the granularity of the monitor. Finer-grained monitoring identifies the modified data in finer granularity with higher monitoring overhead, but is able to copy less unmodified data. With coarser-grained monitoring, the monitoring overhead is reduced but more unmodified data needs to be copied. Our experiment reveals that pagebased incremental checkpointing technique suffers more from significant monitoring overhead. It indicates that a coarser-grained implementation will better balances the tradeoff and boosts the overall performance. However, blindly releasing the granularity does benefit to the monitoring overhead anyway, but results in extremely large checkpoint size which slow both the checkpoint and restart processes.

We observe the memory access characteristics at page granularity and find that the size of contiguous memory regions visited by applications tends to be proportional to the size of corresponding memory allocation. It indicates that the memory allocation information can be leveraged to appropriately release the monitoring granularity. Based on this observation, we propose Application-Based Coarse-Grained Incremental Checkpointing (ACCK) mechanism that strikes the balance between the monitoring overhead and copying overhead in incremental checkpointing. Based on the relationship between contiguous memory regions visited by applications and the corresponding memory allocation, ACCK mechanism releases the monitoring granularity by not monitoring the memory regions that the applications tend to visit. Since page-based incremental checkpointing works with normal-size pages to enable fine-grained monitoring, it inevitably increases the TLB overhead. We implement merge and split operation in Linux kernel which merges normal-size pages (4Kbyte) to huge-size page (2Mbyte) and splits huge-size page back to normal-size pages. ACCK merges normalsize pages when it decides to keep them unmonitored and splits huge-size page at checkpoint time to maintain fine-grained monitoring. This effectively alleviates the TLB overhead during the runtime.

Compared to traditional page-grained incremental checkpointing, ACCK mechanism largely reduces the monitoring overhead with bearable increase in copying overhead. The experimental results show that the performance of memory access monitoring, which accounts for around 70% of the checkpointing time, improves 9x on average. As a result, ACCK presents a 2.5x performance improvement over the baseline page-grained based incremental checkpointing mechanism with the copying overhead increases by 7% and the TLB overhead alleviated.

The remainder of this paper is organized as follows: Section 2 introduces the background and our motivation. We describe the design and implementation of ACCK in detail in Section 3. Section 4 gives the experiments results. Related work is discussed in Section 5 and we conclude in Section 6.

2

Background and Motivation

Non-volatile Memory

Non-volatile memory (NVM) is maturing fast in recent years. NVM features fast access, non-volatile, large capacity, and byte-addressable. Currently, phase change memory (PCM) is the most developed non-volatile memory technology [START_REF] Koltsidas | PSS: A prototype storage subsystem based on PCM[END_REF][START_REF] Coburn | NV-Heaps: making persistent objects fast and safe with next-generation, non-volatile memories[END_REF][START_REF] Bautista-Gomez | FTI: high performance fault tolerance interface for hybrid systems[END_REF]. The read speed of PCM is comparable with DRAM, while the write speed is much slower (typically 10x slower). There are many studies at architecture level focusing on mitigating the write latency [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF][START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF]. In this paper, we emulate the slow write by adding latency to test its performance on supporting checkpoint. HDD data transfer bandwidth has long been a main concern for checkpointing technologies [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF]. NVM allows in-memory copy to accelerate checkpointing. With the NVM support, the age-old checkpoint/restart mechanism seems more attractive due to its simplicity and low cost. Recent work [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF] has shown that PCM-supported inmemory checkpointing is 50 times faster than the conventional in-HDD checkpointing. We assume the non-volatile memory can be directly accessed by CPUs via load and store instructions. It shares the same physical address with DRAM.

Incremental Checkpointing with Non-volatile Memory

As the scaling of computing systems, MTTF continues to decrease from days to hours, and to several minutes [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF][START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF]. The checkpoint interval decreases with MTTF since if the checkpoint time surpass the failure period, it means the possibility of ending up with infinite execution time. As noticed by previous work [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF], frequent checkpointing is only possible by using PCRAM and is critical to benefit from incremental checkpointing schemes. As the checkpoint interval decreases, the probability of polluting a clean page becomes smaller, hence, the incremental checkpointing shows advantages over full-size checkpointing in terms of average checkpoint size. Incremental checkpointing features user-transparency, decent portability and flexibility [START_REF] Plank | Compressed differences: An algorithm for fast incremental checkpointing[END_REF][START_REF] Plank | Libckpt: Transparent checkpointing under Unix[END_REF][START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF]. It requires a memory access monitor to track the dirty data and periodically backs them up to non-volatile memory or hard-disk. Existing researches leverage the hardware page protection mechanism to monitor write operations and back up the data in the granularity of pages. At the beginning of each checkpoint interval, all the writable pages are marked as read-only using mprotect system call. When an mprotected page is written, a page fault exception occurs and sends the SIGSEGV signal. The memory access monitor receives the signal and call the page fault exception handler to save the address of the dirty page in an external data structure. The page fault signal handler also marks the written page as writable by calling unprotect system call. At the end of the checkpoint interval, the monitor scans the data structure that tracks the dirty pages and backs up all the dirty pages. In this paper, we implement this pageprotection based memory access monitor in runtime level and trigger checkpoint operation with the time interval of 5 seconds.

The Problem with Current Checkpoint

Generally, the overhead of incremental checkpointing is decided by the granularity of memory access monitoring [START_REF] Zhang | Fine-grained Checkpoint Based on Non-volatile Memory[END_REF]. The finer the granularity, the more page faults will be triggered but less data will be copied. The coarser the granularity, the less page fault occurs but more unmodified data has to be moved. The ideal case falls somewhere between which involves the least overhead to the total system performance. The widely supported page sizes in operating system are 4K and 2M. Page-based incremental checkpointing works with 4K memory pages. 2M page is too coarse for monitoring which largely increases the checkpoint size. This will make incremental checkpoint fall back to full-size checkpoint with useless monitoring. However, current incremental checkpointing suffers from significant monitoring overhead (as well as TLB overhead) with 4K pages. Once a page is visited, a page fault occurs to handle the fault and the page will be copied at the checkpoint time. We quantify the monitoring overhead and copying overhead to clear our motivation. The rdtsc instruction is used to reckon the page-fault handler by time. We estimate the data movement from DRAM to PCRAM to simulate the copying overhead where each 8 byte atomic writes costs 600ns. Note that both monitoring and copying overhead are generic and applicationindependent. Table 1 shows the average monitoring overhead and copying overhead at 4K and 2M granularity. Monitoring overhead is around 2.5x the time of copying at 4K level while copying data accounts for the most of overhead (~200x) with 2M huge page.

Our experiments in Section IV shows that monitoring overhead accounts for around 70% of total checkpointing time. Therefore, a key to optimize the incremental checkpointing performance is finding a way to minimize the monitoring overhead. Specifically, a coarser-grained monitor that adapts well to the memory access pattern of applications will yield much better results. Previous work [START_REF] Zhang | Fine-grained Checkpoint Based on Non-volatile Memory[END_REF] has tried some intuitive methods to reduce the number of page faults during checkpoint interval, such as leaving the pages around the written pages unmonitored. However, these methods without priori information shows marginal improvement.

DESIGN AND IMPLEMENTATION

Contiguous Memory Regions to Be Visited

Since no priori information is used to predict the memory regions to be visited, previous page-based incremental checkpointing has to track every write operation that trigger the page fault. We observe the memory access pattern of PARSEC and SPLASH benchmark suites at 4K granularity and find that the size of contiguous memory regions visited by applications tends to be proportional to the size of the corresponding memory allocation. For example, if an application malloc a chunk of memory space sizing 100 pages, usually a 30% or 50% of pages will be contiguously visited, which are 30 and 50 pages. Given the same ratio (30% or 50%), for applications typically allocate thousands of pages, these numbers (30 and 50 pages in last example) will be proportionally increased. In this experiment, the total size of applications' memory allocation is divided into 2n pieces where n can be a reasonably large integer, such as 6 or 7. Each piece is regarded as "hot" piece if more than 80% of pages inside the piece has been visited. In this definition of "hot" piece, we do not put strict limitation that all pages have to be visited in order to cover some more flexible memory access patterns.

Figure 1 shows the cumulative distribution function (CDF) and histogram of the priori information we are interested, which is the proportion between contiguous memory space visited by applications and the size of memory allocation, and their corresponding frequency. We run and statistic all applications in PARSEC and SPALSH2 to make this observation more reliable and generic. The parameter n is set to 6 so each memory allocation is divided into 64 pieces. In the histogram and CDF figures, the x-axis shows the ratio between the contiguous pieces visited by applications and the total pieces of memory allocation and y-axis shows the frequency or the proportion of corresponding ratio for memory allocations in all applications. It is noticeable that nearly 30% has no "hot" pieces at all. Around 18% (from 27% to 45%) of allocation has "hot" pieces between 30% and 40% of total pieces allocated while around 17% of allocation has "hot" pieces between 60-70% of total pieces allocated. This observation sheds new light on memory access pattern which can be utilized to appropriately loose the monitoring granularity to achieve better performance of incremental checkpointing for applications with large memory footprint. Our observation shows that applications tend to visit contiguous memory regions with its length often proportional to the size of the corresponding memory allocation.

Based on this observation, we propose Application-Based Coarse-Grained Incremental Checkpointing (ACCK) that looses monitoring granularity for "hot" applications. ACCK first records of the start address and size of memory regions applied by applications. ACCK keeps track of the memory access pattern at 4K page granularity and divides the memory allocation into 2n pieces (at least one page for each piece).

Here, we adopt a new definition of "hot" piece that for each piece of memory, if more than 80% pages out of first 10% of pages are visited, it is marked as a potential "hot" piece. ACCK releases the monitoring granularity incrementally and gradually. Given that there are 2n pieces for each memory allocation, ACCK leaves the pieces unmonitored based on the memory access pattern at the pace of (20, 21, …, 2n-3), (20, 21…, 2n-3), …, (20, 21, …, 2n-3) pieces. Specifically, at the beginning of checkpoint, every piece is monitored. If a monitored piece becomes a potential "hot" piece, it will be left unmonitored by unprotecting all the pages inside the piece. If the following piece become a potential "hot" piece again, this piece and the following piece (a total of 21 = 2 pieces) will left unmonitored. Similarly, if following piece becomes potential "hot", ACCK leaves the following pieces unmonitored at the pace of 22, 23…until 2n-3 pieces, then it repeat this process from the pace 20 to prevent over coarse granularity. For example, if an application visits all pieces in order, ACCK mechanism (with n=6) releases the monitor at the pace of (1, 2, 4, 8) pieces and repeat until the end of memory regions. Figure 2 (a) shows the ACCK example for an in-order memory access pattern and Figure 2 (b) shows an ACCK example for an any-order memory access pattern. The dotted areas in Figure 2 denote the recognized potential "hot" pieces, which trigger the following arrows. Note that the arrows in Figure 2 does not have to occur in order.

Huge Page Support

2M huge page has a theoretical 512x performance improvement in terms of TLB performance over 4K page. However, previous researches are cautious about the use of huge pages. If huge pages are not used appropriately, a large amount of unmodified data has to be copied at the end of checkpoint interval.

In this work, we implement a kernel patch to Linux kernel 3.18 to provide merge and split operations between 4K pages and 2M huge pages. A merge operation merges 512 contiguous, 2M-alined 4K pages into a huge page (Figure 3 top). A split operation splits a 2M huge page into 512 4K pages (Figure 3 bottom). These functions are exposed to runtime library via system calls.

In ACCK mechanism, if a large number of pieces with more than 512 4K pages in total will be unmonitored, these pages will be merged into huge pages if possible (note that huge pages to be merged into must be 2M-aligned). If the merge succeeds, ACCK turns off the page protection of the huge pages to leave it unmonitored. If the merge fails, ACCK remains to work in the traditional 4K granularity. In order to monitor in fine granularity at the beginning of checkpoint interval, the merged huge pages will be split into 4K pages at the end of checkpoint interval. The benefits of huge page comes from low TLB miss rate and penalty during runtime.

EXPERIMENTS

This paper introduces Application-Based Coarse-Grained Incremental Checkpointing (ACCK). This section mainly evaluates and compares the performance of ACCK mechanism with the page-grained incremental checkpointing techniques.

Experimental Setup and Benchmarks

Our experiment platform is an AMD sever equipped with 2.2GHz 12-core CPU and 16GB of physical memory. The operating system is Linux 4.2. Our benchmarks come from PARSEC and SPLASH benchmark suite [START_REF] Bienia | The PARSEC benchmark suite: Characterization and architectural implications[END_REF]. It focuses on programs from all domains, including HPC applications and server applications. The largest "native" input set is used.

Performance Metrics and Corresponding Results

Monitoring Overhead. ACCK mechanism mainly focuses on reducing the significant monitoring overhead of incremental checkpoint. Given the fact that the unit monitoring overhead is around 2.5x of unit copying overhead (application-independent), ACCK mechanism appropriately release the monitoring granularity with useful priori information. Figure 4 shows the significant performance improvement in terms of monitoring overhead in incremental checkpointing. As can be seen, ACCK mechanism lowers the monitoring overhead for all applications. The improvement can be as significant as 10x for most applications. The application freqmine is an exception with its reduction not as significant as other applications. The reason behind is that there is not enough potential "hot" pieces in its memory allocation, which does not satisfy the requirements to trigger ACCK mechanism. In this case, the discrete memory access pattern will cause ACCK to fall back to the baseline incremental checkpointing. However, due to negligible management overhead of ACCK, the monitoring overhead of ACCK is always much better than traditional page-grained incremental checkpointing. Copying Overhead. As mentioned in the previous section, the copying overhead is not as significant as monitoring overhead per memory page. ACCK mechanism releases monitoring granularity and increases copying overhead during checkpoint. However, our experimental results proves that the additional copying overhead is minimal compared to the improved monitoring performance. Figure 5 shows the total copying overhead of ACCK and the baseline incremental checkpoint. Only an average of 7.4% more data has to be copied. Moreover, we argue that the preCopy mechanism [START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF] to pre-move data before checkpoint time reduces the memory pressure at checkpoint time. Therefore, the additional copying overhead can be effectively amortized.

Fig. 5. Copying Overhead

TLB Overhead. We also evaluate the performance of the page size adjustment. ACCK mechanism merge 4K pages into huge pages if a large number of pieces with more than 512 4K pages in total will be unmonitored. The use of huge page helps improve the TLB performance since each TLB entry maps much larger memory region with huge page. We use oprofile profiling tool to evaluate the data TLB miss using DTLB_MISS counter where the sampling count is set to 500. Figure 6 shows the normalized data TLB miss of ACCK mechanism with and without huge page support. Note that the merged pages will fall back to 4K pages at checkpoint time to maintain fine-grained monitoring at start. Overall Checkpointing Performance. Generally, the overhead of incremental checkpoint consists of the monitoring overhead and the copying overhead. ACCK mechanism sacrifices the copying overhead for significant reduction in monitoring overhead and our experimental results prove it to be very effective in terms of overall checkpointing performance. Figure 6 gives the overall checkpointing performance of ACCK mechanism and the baseline incremental checkpoint. The overall checkpointing performance is defined as the reciprocal of the summation of the monitoring overhead and copying overhead. ACCK mechanism achieves 2.56x performance improvement over the baseline incremental checkpointing. It is noticeable that the performance improvement of each benchmark is pretty average, with the highest improvement 2.79x (ocean_cp) and lowest improvement 2.2x (freqmine).

Related Work

As checkpoint-restart is the commonly-used technique for fault-tolerance, the related research is abundant. Non-volatile memory sheds new light on fault tolerance. Previous work [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF][START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF] leverages the byte persistency of non-volatile memory to do in-memory copy to accelerate checkpointing. To address the slow write speed and limited bandwidth during checkpoint, [START_REF] Dong | Leveraging 3D PCRAM technologies to reduce checkpoint overhead for future exascale systems[END_REF] proposes a preCopy mechanism to move checkpoint data to non-volatile memory before checkpoint time to amortize the memory bandwidth pressure at checkpoint time. [START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF] proposes a 3D PCMDRAM design at architectural level to facilitate data movement between DRAM and PCM. These two studies both focus on hiding the long write latency of non-volatile memory. Most of related work of reducing the checkpoint size uses the incremental checkpointing technique [START_REF] Sancho | On the feasibility of incremental checkpointing for scientific computing[END_REF] that only saves the dirty data between two consecutive checkpoints. Since then, the hardware memory management mechanism has been leveraged to monitor the dirty data. However, incremental checkpoint technique still suffers from significant monitoring overhead and struggle with the granularity of memory access monitor. This paper propose a novel monitoring mechanism which can reduce the monitoring overhead with bearable increase in checkpoint size. Our work is orthogonal to previous studies (e.g. pre-copy mechanism [START_REF] Kannan | Optimizing checkpoints using nvm as virtual memory[END_REF]) and can be combined to achieve better performance.

CONCLUSION

Checkpoint-restart has been an effective mechanism to guarantee the reliability and consistency of computing systems. Our work mainly addresses the significant monitoring overhead of current incremental checkpointing technique. This paper proposes Application-Based Coarse-Grained Incremental Checkpointing (ACCK) mechanism based on non-volatile memory. We observe the memory access characteristics and find that the size of contiguous memory regions heavily visited by applications tends to be proportional to the size of allocated memory space. ACCK leverages the priori information of the memory allocation to release the memory monitoring granularity in an incremental and appropriate way. The experimental results shows that ACCK largely reduces the monitoring time and presents 2.56x overall checkpointing performance improvement. This work can be applied to frequent checkpoint of a wide range of applications and databases and can be combined with other work (e.g. pre-copy mechanism) to achieve better checkpoint performance.

Fig. 1 .

 1 Fig. 1. Histogram and cumulative distribution function (CDF) of the ratio between contiguous memory pieces visited by applications and total allocated pieces

Fig. 2 .

 2 Fig. 2. Example of ACCK mechanism for in-order and any-order memory access pattern

Fig. 3 .

 3 Fig. 3. An example of the merge (top) and split (bottom) operation done by modification on Linux.

Fig. 4 .

 4 Fig. 4. Monitoring Overhead

Fig. 6 .

 6 Fig. 6. Normalized data TLB miss with and without huge page support

Fig. 7 .

 7 Fig. 7. Overall Performance

Table 1 .

 1 Comparison of Monitoring and Copying Overhead at 4K and 2M granularity

	Granularity	Monitoring Overhead	Copying Overhead
	4K	~100 us per 4KByte	< 40 us per 4KByte
	2M	~0.2 us per 4KByte	< 40 us per 4KByte

Acknowledgments. This work is partially supported by National Hightech R&D Program of China (863 Program) under Grants 2012AA01A301 and 2012AA010901, by program for New Century Excellent Talents in University, by National Science Foundation (NSF) China 61272142, 61402492, 61402486, 61379146, 61272483, and by the State Key Laboratory of High-end Server & Storage Technology (2014HSSA01). The authors will gratefully acknowledge the helpful suggestions of the reviewers, which have improved the presentation.


Bang Di
email: dibang@hnu.edu.cn

Jianhua Sun
email: jhsun@hnu.edu.cn

Hao Chen
email: haochen@hnu.edu.cn

A Study of Overflow Vulnerabilities on GPUs

Keywords: GPGPU, CUDA, Security, Buffer overflow

GPU-accelerated computing gains rapidly-growing popularity in many areas such as scientific computing, database systems, and cloud environments. However, there are less investigations on the security implications of concurrently running GPU applications. In this paper, we explore security vulnerabilities of CUDA from multiple dimensions. In particular, we first present a study on GPU stack, and reveal that stack overflow of CUDA can affect the execution of other threads by manipulating different memory spaces. Then, we show that the heap of CUDA is organized in a way that allows threads from the same warp or different blocks or even kernels to overwrite each other's content, which indicates a high risk of corrupting data or steering the execution flow by overwriting function pointers. Furthermore, we verify that integer overflow and function pointer overflow in struct also can be exploited on GPUs. But other attacks against format string and exception handler seems not feasible due to the design choices of CUDA runtime and programming language features. Finally, we propose potential solutions of preventing the presented vulnerabilities for CUDA.

Introduction

Graphics processing units (GPUs) were originally developed to perform complex mathematical and geometric calculations that are indispensable parts of graphics rendering. Nowadays, due to the high performance and data parallelism, GPUs have been increasingly adopted to perform generic computational tasks. For example, GPUs can provide a significant speed-up for financial and scientific computations. GPUs also have been used to accelerate network traffic processing in software routers by offloading specific computations to GPUs. Computation-intensive encryption algorithms like AES have also been ported to GPU platforms to exploit the data parallelism, and significant improvement in throughput was reported. In addition, using GPUs for co-processing in database systems, such as offloading query processing to GPUs, has also been shown to be beneficial.

With the remarkable success of adopting GPUs in a diverse range of realworld applications, especially the flourish of cloud computing and advancement in GPU virtualization [START_REF] Shi | vcuda: Gpu-accelerated high-performance computing in virtual machines[END_REF], sharing GPUs among cloud tenants is increasingly becoming norm. For example, major cloud vendors such as Amazon and Alibaba both offer GPU support for customers. However, this poses great challenges in guaranteeing strong isolation between different tenants sharing GPU devices.

As will be discussed in this paper, common well-studied security vulnerabilities on CPUs, such as the stack and heap overflow and integer overflow, exist on GPUs too. Unfortunately, with high concurrency and lacking effective protection, GPUs are subject to greater threat. In fact, the execution model of GPU programs consisting of CPU code and GPU code, is different from traditional programs that only contains host-side code. After launching a GPU kernel (defined in the following section), the execution of the GPU code is delegated to the device and its driver. Therefore, we can know that the GPU is isolated from the CPU from the perspective of code execution, which means that the CPU can not supervise its execution. Thus, existing protection techniques implemented on CPUs are invalid for GPUs. On the other hand, the massively parallel execution model of GPUs makes it difficult to implement efficient solutions tackling security issues. Unfortunately, despite GPU's pervasiveness in many fields, a thorough awareness of GPU security is lacking, and the security of GPUs are subject to threat especially in scenarios where GPUs are shared, such as GPU clusters and cloud.

From the above discussion, we know that GPUs may become a weakness that can be exploited by adversaries to execute malicious code to circumvent detection or steal sensitive information. For example, although GPU-assisted encryption algorithms achieve high performance, information leakage such as private secrete key has been proven to be feasible [START_REF] Pietro | CUDA leaks: A detailed hack for CUDA and a (partial) fix[END_REF]. In particular, in order to fully exert the computing power of GPUs, effective approaches to providing shared access to GPUs has been proposed in the literature [START_REF] Pai | Improving GPGPU concurrency with elastic kernels[END_REF][START_REF] Ravi | Supporting GPU sharing in cloud environments with a transparent runtime consolidation framework[END_REF]. However, without proper mediation mechanisms, shared access may cause information leakage as demonstrated in [START_REF] Pietro | CUDA leaks: A detailed hack for CUDA and a (partial) fix[END_REF]. Furthermore, we expect that other traditional software vulnerabilities on CPU platforms would have important implications for GPUs, because of similar language features (CUDA programming language inherits C/C++). Although preliminary experiments has been conducted to show the impact of overflow issues on GPU security [START_REF] Miele | Buffer overflow vulnerabilities in CUDA: a preliminary analysis[END_REF], much remains unclear considering the wide spectrum of security issues. In this paper, we explore the potential overflow vulnerabilities on GPUs. To the best of our knowledge, it is the most extensive study on overflow issues for GPU architectures. Our evaluation was conducted from multiple aspects, which not only includes different types of attacks but also considers specific GPU architectural features like distinct memory spaces and concurrent kernel execution. Although we focus on the CUDA platform, we believe the results are also applicable to other GPU programming frameworks.

The rest of this paper is organized as follows. Section 2 provides necessary background about the CUDA architecture. In Section 3, we perform an extensive evaluation on how traditional overflow vulnerabilities can be implemented on GPUs to affect the execution flow. Possible countermeasures are discussed in Section 4. Section 5 presents related work, and Section 6 concludes this paper.

Background on CUDA Architecture

CUDA is a popular general purpose computing platform for NVIDIA GPUs. CUDA is composed of device driver (handles the low-level interactions with the GPU), the runtime, and the compilation tool-chain. An application written for CUDA consists of host code running on the CPU, and device code typically called kernels that runs on the GPU. A running kernel consists of a vast amount of GPU threads. Threads are grouped into blocks, and blocks are grouped into grids. The basic execution unit is warp that typically contains 32 threads. Each thread has its own program counters, registers, and local memory. A block is an independent unit of parallelism, and can execute independently of other thread blocks. Each thread block has a private per-block shared memory space used for inter-thread communication and data sharing when implementing parallel algorithms. A grid is an array of blocks that can execute the same kernel concurrently. An entire grid is handled by a single GPU.

The GPU kernel execution consists of the following four steps: (i) input data is transferred from the host memory to GPU memory through the DMA; (ii) a host program instructs the GPU to launch a kernel; (iii) the output is transferred from the device memory back to the host memory through the DMA.

CUDA provides different memory spaces. During execution, CUDA threads may access data from multiple memory spaces. Each thread maintains its own private local memory that actually resides in global memory. Automatic variables declared inside a kernel are mapped to local memory. The on-chip shared memory is accessible to all the threads that are in the same block. The shared memory features low-latency access (similar to L1 cache), and is mainly used for sharing data among threads belonging to the same block. The global memory (also called device memory) is accessible to all threads, and can be accessed by both GPU and CPU. There are two read-only memory spaces accessible by all threads, i.e. constant and texture memory. Texture memory also offers different addressing models, as well as data filtering, for some specific data formats. The global, constant, and texture memory spaces are optimized for different memory usages, and they are persistent across kernel launches by the same application.

Empirical Evaluation of GPU Vulnerabilities

In this section, we first introduce the testing environment. Then, we discuss specific vulnerabilities for stack overflow, heap overflow, and others respectively, with a focus on the heap overflow because of its potential negative impact and significance in scenarios where multiple users share GPU devices. Due to the proprietary nature of the CUDA platform, we can only experimentally confirm the existence of certain vulnerabilities. And further exploration about inherent reasons and such issues is beyond the scope of this paper, which may require a deeper understanding of the underlying implementation of CUDA framework and hardware device intricacy. 14 pFdummy fp [START_REF] Breß | Datenbanksysteme für Business, Technologie und Web (BTW), -Workshopband, 15. Fachtagung des GI-Fachbereichs "Datenbanken und Informationssysteme[END_REF];

15 fp[0]=normal1; fp[1]=normal2; fp[2]=normal3; fp[3]=normal4;
16 fp [START_REF] Ravi | Supporting GPU sharing in cloud environments with a transparent runtime consolidation framework[END_REF]=normal5; fp [START_REF] Miele | Buffer overflow vulnerabilities in CUDA: a preliminary analysis[END_REF]=normal6; fp [START_REF] Silberstein | Gpufs: Integrating a file system with gpus[END_REF]=normal7; fp [START_REF]NVIDIA: Cuda-memcheck[END_REF]=normal8;

17 for(int i = 0; i < length; i++) {buf[i] = overf[i];} 18 fp[5];

Experiment Setup

The machine conducting the experiment has a Intel Core i5-4590 CPU clocked at 3.30GHz, and the GPU is NVIDIA GeForce GTX 750Ti (Maxwell architecture) that has compute capability 5.0. The operating system is Ubuntu 14.04.4 LTS (64 bit) with CUDA 7.5 installed. nvcc is used to compile CUDA code, and NVIDIA visual profiler is adopted as a performance profiling tool. CUDA-GDB allows us to debug both the CPU and GPU portions of the application simultaneously. The source code of all implemented benchmarks is publicly available at https: //github.com/aimlab/cuda-overflow.

Stack Overflow

In this section, we investigate the stack overflow on GPUs by considering different memory spaces that store adversary-controlled data, and exploring all possible interactions among threads that are located in the same block, or in different blocks of the same kernel, or in distinct kernels.

The main idea is as follows. The adversary formulates malicious input data that contains the address of a malicious function, and assign it to variable a that is defined in global scope. Two stack variables b and c are declared in a way to make their addresses adjacent. If we use a to assign values to b to intentionally overflow b and consequently corrupt the stack variable c that stores function pointers. Then, when one of the function pointers of c is invoked, the execution flow would be diverted to the adversary-controlled function. Note that there is a difference of the stack between the GPU and CPU. In fact, the storage allocation of GPU stack is similar to the heap, so the direction of overflow is from low address to high address.

We explain how a malicious kernel can manipulate a benign kernel's stack with an illustrating example that is shown in Figure 1. In the GPU code, we define 9 functions containing 1 malicious function (used to simulate malicious behavior) and 8 normal functions (only one is shown in Figure 1, and the other 7 functions are the same as the function normal1 except the naming). The device qualifier declares a function that is executed on the device and callable from the device only. The noinline function qualifier can be used as a hint for the compiler not to inline the function if possible. The array overf[100] is declared globally to store data from another array input[100] that is controlled by the malicious kernel. Given the global scope, the storage of overf[100] is allocated in the global memory space, indicating both the malicious kernel and benign kernel can access. In addition, two arrays named buf and fp are declared one after another on the stack to ensure that their addresses are consecutively assigned. The fp stores function pointers that point to the normal functions declared before, and the data in overf[100] is copied to buf (shown at line 17) to trigger the overflow. The length variable is used to control how many words should be copied from overf to buf (shown at line 17). It is worth noting that the line 12 is only executed in the malicious kernel to initialize the overf buffer. If we set length to 26 and initialize overf with the value 0x590 (address of the malicious function that can be obtained using printf("%p",malicious) or CUDA-GDB [START_REF] Miele | Buffer overflow vulnerabilities in CUDA: a preliminary analysis[END_REF]), the output at line 18 would be string "Normal". This is because with value 26, we can only overwrite the first 5 pointers in fp (sizeof(buf ) + sizeof(pFdummy) * 5 == 26 ). However, setting length to 27 would cause the output at line 18 to be "Attack!", indicating that fp [START_REF] Miele | Buffer overflow vulnerabilities in CUDA: a preliminary analysis[END_REF] is successfully overwritten by the address of the malicious function. This example demonstrates that current GPUs have no mechanisms to prevent stack overflow like stack canaries on the CPU counterpart. It is straightforward to extend our experiments to other scenarios. For example, by locating the array overf in the shared memory, we can observe that the attack is feasible only if the malicious thread and benign thread both reside in the same block. While if overf is in the local memory, other threads haves no way to conduct malicious activities. In summary, our evaluation shows that attacking a GPU kernel based on stack overflow is possible, but the risk level of such vulnerability depends on specific conditions like explicit communication between kernels.

Heap Overflow

In this section, we study a set of heap vulnerabilities in CUDA. We first investigate the heap isolation on CUDA GPUs. Then, we discuss how to corrupt locally-allocated heap data when the malicious and benign threads co-locate in the same block. Finally, we generalize the heap overflow to cases where two kernels are run sequentially or concurrently. Heap Isolation. Similar to the description of stack overflow, we also use a running example to illustrate heap isolation from two aspects. First, we consider the case of a single kernel. As shown in Figure 3 (from line 15 to 22), suppose we have two independent threads t 1 and t 2 in the same block, and a pointer variable buf is defined in the shared memory. We can obtain similar results when buf is defined in the global memory. For clarity, we use buf 1 and buf 2 to represent the buf in t 1 and t 2 . buf 1 is allocated by calling malloc as shown at lines 16 and 17. Our experiments show that t 2 can always access buf 1 (line 21 to 22) unless buf is defined in the local scope. Second, we consider the case of two kernels (not shown in the figure). Kernel a allocates memory space for buf, and assigns the input value to it. If a returns without freeing the memory of buf, another kernel b can always read the content in buf if b also has a variable buf defined in either shared memory or global memory (no memory allocation for buf in b). This is because the GPU assigns the same address to buf for b, which makes it possible to access the not freed content of buf in b. In summary, for globally-defined heap pointer, the memory it points to can be freely accessed by threads that are not the original allocator. It is not the case for locally-defined heap pointers, but it may still be possible if we can successfully guest the addresses of local heap pointers (we leave this to future work). Most importantly, when a heap pointer is globally visible and the corresponding memory is not properly managed (freed), arbitrary memory accesses across kernels would be possible.

Heap Exploitation. In this experiment, we present that because the heap memory for different threads or kernels is allocated contiguously, overflowing one thread's local buffer may lead to the corruption of another thread's heap data.

We first consider heap memory allocated in local scope. As shown in Figure 3 (from line 24 to 29), like before, suppose we have two threads t 1 and t 2 in the same block. For t 1 , we use malloc and new to allocate memory for buf 1 and f p 1 respectively (we use these notations to clarify our discussion). f p 1 just stores the start address of the four virtual functions (the addresses is contained in the VTABLE). buf 1 and f p 1 are declared locally. t 2 is the same as t 1 . After initializing t 1 and t 2 , the memory layout of buf 1 , f p 1 , buf 2 , and f p 2 looks like that shown in Figure 4. Assume t 2 has malicious intention. The input in t 1 consists of normal data, and the input in t 2 consists of four addresses of the malicious function (0x138), and the remaining contents of input are the address of buf 2 (0x50263f920). When length is less than 11 (not 9 due to alignment), both t 1 and t 2 will print the string "Normal". However, when length is set to 11, the virtual table address in f p 1 would be modified to the start address of buf 2 where the four addresses of the malicious function are stored. So the output of t 1 will be the string "Attack!". When length is set to 21 (this value is relative to f p 1 ), both t 1 and t 2 will invoke the malicious function. Similarly, assuming t 1 is the malicious thread, both t 1 and t 2 will output "Normal" when length is less than 21. By assigning 21 to length, only t 1 will print "Attack!". And when the value of length is 31, both t 1 and t 2 will invoke the malicious function.

Based on the analysis above, we can conclude that the memory allocated from the heap in CUDA is globally accessible from different GPU threads without proper access control, and no protection is provided by the runtime system to prevent buffer overflow of heap-allocated memory. The addresses of heap pointers are often guessable, which makes it easy for a adversary to conduct attacks.

Because of the closed-source nature of CUDA, further investigation about the implementation-level details is left as future work. In the following, we extend our analysis to the scenario where two kernels are considered when experimenting buffer overflow. Heap Exploitation Between Kernels. First, we consider the overflow between sequentially launched kernels. The experiment setup is similar to the previous section except that the host launches serially two kernels, kernel 1 (launched first) and kernel 2 . kernel 1 simulates behaviors of an adversary and initializes the thread t 1 . kernel 2 represents a benign user. t 1 's input data consist of four addresses of malicious function, and its remaining contents are the address of buf 1 . kernel 1 intentionally overflows buf 1 by using a large value of length 1 , and then terminates. kernel 2 allocates memory for buf 2 and f p 2 , but does not assign any value to them. When accessing one element in f p 2 as a function call, we will observe the output of string "Attack!". This is because the GPU assigns the same addresses of pointer variables to the second kernel, and the contents used in the first kernel are remained in GPU memory unchanged.

Second, we analyze the situation of two concurrently running kernels as shown in Figure 5. This is the typical case of sharing GPUs in cloud or cluster environment, where different users may access the GPU at the same time. kernel 1 and kernel 2 must be in different streams. cudaMemcpyAsync is called to receive data from the host that allocates page-locked memory. The sleep() function is used to perform a simple synchronization between kernels to make our experiments more deterministic. t 1 in kernel 1 mimics a malicious user's behavior. When buf 2 and f p 2 are initialized (from line 9 to 12) and at the same time t 1 has finished the execution of the for loop (at line 20), we will observe the output of string "Attack!" from t 2 if it continues to run after the pause at line 13. Based on these observations, we can conclude that multiple kernels, regardless of serially or concurrently running, have the opportunity to intentionally manipulate the heap memory of each other.

Other Vulnerabilities

In this section, we discuss the issues of struct and integer overflow that are demonstrated in one example. In this case, the attacker can exploit the characteristics of integer operations. Because of the two's complementary representation of integers, integer operations may produce undesirable results when an arithmetic operation attempts to produce a numeric value that is too large to be representable with a certain type.

In this experiment (Figure 6), we define two variables input [START_REF] Maurice | Confidentiality issues on a GPU in a virtualized environment[END_REF] and length, which stores user data and data size respectively. The device code for functions normal and malicious are the same as defined in Figure 5. In addition, we define a struct unsafe, which contains an array buf [START_REF] Silberstein | Gpufs: Integrating a file system with gpus[END_REF] and a function pointer of type normal. The init() (line 6) function is used to initialize the structure defined above. The if statement (line 12) performs array-bounds check to prevent out-of-bound access. Suppose that the address of the malicious function is 0x43800000438 that is assigned to input [START_REF] Maurice | Confidentiality issues on a GPU in a virtualized environment[END_REF] as input data by a malicious user. The variable length whose type is unsigned char is set to 0 by the attacker. The if branch would be executed because the value of length (0) is smaller than 6. But, the value of length will be 255 after it is decremented by one, which causes that the assignment at line 15 overflows the array buf and corrupt the function pointer in struct unsafe. This experiment shows that struct and integers can both be exploited in CUDA. In particular, both overflow and underflow of integer arithmetic operation are possible, which exposes more opportunities for the adversaries. Format string and exception handling vulnerabilities have been studied on CPUs. However, our experiments show that they currently are not exploitable on the GPU due to the limited support in CUDA. For format string, the formatted output is only supported by devices with compute capability 2.x and higher. The in-kernel printf function behaves in a similar way to the counterpart in the standard C-library. In essence, the string passed in as format is output to a stream on the host, which makes it impossible to conduct malicious behavior through printf to expose memory errors on the GPU. For the exception handling in C++, it is only supported for the host code, but not the device code. Therefore, we can not exploit the security issues of exception handling on CUDA GPUs currently.

Discussions and Countermeasures

This section discusses potential countermeasures that can prevent or restrict the impact of the described vulnerabilities. Basically, most discussed weaknesses have limited impact on current GPUs if considered from the security point of view. For example, stack overflow and in-kernel heap overflow can only corrupt the memory data from within the same kernel. It is difficult to inject executable code into a running kernel to change its control flow under the GPU programming model. Of course, it may not be the case when more functionalities are integrated to make developing GPU applications more like the CPU counterparts [START_REF] Silberstein | Gpufs: Integrating a file system with gpus[END_REF]. However, exploitable vulnerabilities for concurrent kernels pose real and practical risks on GPUs, and effective defenses are required for a secure GPU computing environment.

For security issues of the heap, calling the free() function, to a certain extent, can offer necessary protection for heap overflows. When an attacker wants to exploit vulnerabilities to corrupt the heap memory that belongs to other threads and has been freed, error message is reported and the application is terminated. But this is hard to guarantee in environments where multiple kernels are running concurrently. The bounds check can be used to prevent overflow for both the stack and heap, but the performance overhead is non-trivial if frequently invoked at runtime. In addition, integer overflow may be leveraged to bypass the bounds check. Therefore, a multifaceted solution is desired to provide full protection against these potential attacks. For example, standalone or hybrid dynamic and static analysis approaches can be designed to detect memory and integer overflows. Ideally, such protections should be implemented by the underlying system such as the runtime system, compiler, or hardware of GPU platforms, as many countermeasures against memory errors on CPUs, such as stack canaries, address sanitizers, and address randomization, have been shown to be effective and practical. Unfortunately, similar countermeasures are still missing on GPUs.

CUDA-MEMCHECK [START_REF]NVIDIA: Cuda-memcheck[END_REF] can identify the source and cause of memory access errors in GPU code, so it can detect all the heap related overflows in our experiments. But it fails to identify the stack overflows. CUDA-MEMCHECK is effective as a testing tool, but the runtime overhead makes it impractical to be deployed in production environments. Most importantly, when multiple mutually untrusted users share a GPU device, dynamic protection mechanisms are indispensable. Thus, we argue that research endeavors should be devoted to the design of efficient approaches to preventing security issues for GPUs.

A wide range of defensive techniques has been proposed to prevent or detect buffer overflows on CPUs. Compiler-oriented techniques including canaries, bounds checking, and tagging are especially useful in our context. Given that CUDA's backend is based on the LLVM framework and Clang already supports buffer overflow detection (AddressSanitizer) through compiler options, we believe these existing tools can be leveraged to implement efficient defenses against GPU-based overflow vulnerabilities.

Related Work

Recently, using GPUs in cloud environment has been shown to be beneficial. The authors of [START_REF] Ravi | Supporting GPU sharing in cloud environments with a transparent runtime consolidation framework[END_REF] present a framework to enable applications executing within virtual machines to transparently share one or more GPUs. Their approach aims at energy efficiency and do not consider the security issues as discussed in this work. For GPU-assisted database systems, the study in [START_REF] Breß | Datenbanksysteme für Business, Technologie und Web (BTW), -Workshopband, 15. Fachtagung des GI-Fachbereichs "Datenbanken und Informationssysteme[END_REF] shows that data in GPU memory is retrievable by other processes by creating a memory dump of the device memory. In [START_REF] Lee | Stealing webpages rendered on your browser by exploiting GPU vulnerabilities[END_REF], it is demonstrated feasible to recover user-visited web pages in widely-used browsers such as Chromium and Firefox, because these web browsers rely on GPUs to render web pages but do not scrub the con-tent remained on GPU memory, leading to information leakage. The paper [START_REF] Maurice | Confidentiality issues on a GPU in a virtualized environment[END_REF] highlights possible information leakage of GPUs in virtualized and cloud environments. They find that GPU's global memory is zeroed only when ECC (Error Correction Codes) is enabled, which poses high risk of private information exposure when ECC is disabled or not available on GPUs. In [START_REF] Pietro | CUDA leaks: A detailed hack for CUDA and a (partial) fix[END_REF], the authors present a detailed analysis of information leakage in CUDA for multiple memory spaces including the global memory, shared memory, and register. A real case study is performed on a GPU-based AES encryption implementation to reveal the vulnerability of leaking private keys.

However, all existing studies on GPU security have less considered vulnerabilities that are more extensively studied on CPUs. The paper [START_REF] Miele | Buffer overflow vulnerabilities in CUDA: a preliminary analysis[END_REF] presents a preliminary study of buffer overflow vulnerabilities in CUDA, but the breadth and depth are limited as compared to our work. For example, we put more emphasis on heap overflows between concurrently running kernels, which we believe deserves special attentions in future research of securing shared GPU access.

Conclusion

In this paper, we have investigated the weakness of GPUs under malicious intentions with a focus on the CUDA platform. We believe the experiments conducted in this work are also applicable to other GPU frameworks such as OpenCL to reveal the potential security vulnerabilities. Particularly, we have confirmed the existence of stack and heap overflows through a diversified set of experiments, and also uncovered how integer overflow can be used in overwriting a function pointer in a struct. Other security issues such as format string and exception handling are also investigated. Although direct exploitation of these potential vulnerabilities is not feasible on current CUDA platform, care must be taken when developing applications for future GPU platforms, because GPU programming platforms are evolving with a fast pace. We hope this study can not only disclose real security issues for GPUs but stimulate future research on this topic especially for scenarios where GPU devices are shared among untrusted users.

  i = 0; i < length; i++) {overf[i] = input[i];} 13 unsigned int buf[16];

Fig. 1 .

 1 Fig. 1. A code snippet of stack overflow in device.

Fig. 2 .

 2 Fig. 2. Illustration of stack overflow.

8 }; 9 /

 89 long v1() {printf("Normal\n");return 0;} 5 __device__ virtual unsigned long v2() {printf("Normal\n");return 0;} 6 __device__ virtual unsigned long v3() {printf("Normal\n");return 0;} 7 __device__ virtual unsigned long v4() {printf("Normal\n");return 0;} a snippet code of memory isolation of heap =========== 15 __shared__ unsigned long *buf; 16 if(threadIdx.x == 0) 17 buf = (unsigned long *) malloc(sizeof(unsigned long) * 8); 18 Vtable *fp = new Vtable; 19 if(threadIdx.x == 0) 20 for(int i = 0; i < length; i++) {buf[i] = input[i];} 21 if(threadIdx.x == 1) 22 printf("%lx", buf[0]); 23 //============ a snippet code of exploiting of 'global' heap ========= 24 unsigned long *buf; 25 buf = (unsigned long *) malloc(sizeof(unsigned long) * 8); 26 Vtable *fp = new Vtable; 27 printf("malicious %p\n", malicious); 28 for(int i = 0; i < length; i++) {buf[i] = input[i];} 29 res=fp->v1(); res=fp->v2(); res=fp->v3(); res=fp->v4();

Fig. 3 .

 3 Fig. 3. A code snippet of heap overflow

Fig. 4 .

 4 Fig. 4. A heap overflow: t1 and t2 represent the benign and malicious thread respectively.

1 /

 1 unsigned long *) malloc(sizeof(unsigned long) * 8); 11 Vtable *fp = new Vtable; 12 for(int i = 0; i < length; i++) {buf[i] = input[i];} 13 sleep(10000000); 14 res=fp->v1(); res=fp->v2(); res=fp->v3(); res=fp->v4(); 15 //==================== malicious kernel ==================== 16 unsigned long *buf; 17 buf = (unsigned long *) malloc(sizeof(unsigned long) * 8); 18 Vtable *fp = new Vtable; 19 sleep(1000000); 20 for(int i = 0; i < length; {buf[i] = input[i];} 21 res=fp->v1(); res=fp->v2(); res=fp->v3(); res=fp->v4();

Fig. 5 .

 5 Fig. 5. A code snippet of concurrent kernel execution.

1 // 8 } 9 __global__Fig. 6 .

 1896 Fig. 6. A snippet code of struct and integer vulnerabilities

Acknowledgment

This research was supported in part by the National Science Foundation of China under grants 61272190, 61572179 and 61173166.


Jian An
email: anjian@mail

Ruobiao Wu
email: wuruobiao@stu

Lele Xiang
email: xianglele@stu

Xiaolin Gui
email: xlgui@mail

Zhenlong Peng
email: pengzl@stu.xjtu.edu.cn

FCM:A Fine-grained Crowdsourcing Model based on Ontology in Crowd-Sensing

Crowd sensing between users with smart mobile devices is a new trend of development in Internet. In order to recommend the suitable service providers for crowd sensing requests, this paper presents a Fine-grained Crowdsourcing Model (FCM) based on Ontology theory that helps users to select appropriate service providers. First, the characteristic properties which extracted from the service request will be compared with the service provider based on ontology triple. Second, recommendation index of each service provider is calculated through similarity analysis and cluster analysis. Finally, the service decision tree is proposed to predict and recommend appropriate candidate users to participate in crowd sensing service. Experimental results show that this method provides more accurate recommendation than present recommendation systems and consumes less time to find the service provider through clustering algorithm.

Introduction

In the new trend of ubiquitous computing, Crowd Sensing (CS) has been proposed to efficiently collect information in the multiple-source and heterogeneous environment. Based on massive information, including users' behavior characteristics, context information and activities law, the user behavior model and the relationship between service attributes could be obtained.

However, there are still some problems needed to be solved in the crowd sensing scenario. To start with, present crowdsourcing models do not particularly classifying the service requests [START_REF] Doan | Crowdsourcing systems on the World-Wide Web[END_REF]. Furthermore, the internal relationship between service preference of customers and historical data has not been fully considered in the existing service recommending model [START_REF] Yang | A location-aware recommender system for mobile shopping environments[END_REF]. Thus, finding an applicable model to utilize the historical data is an effective way to promote success rate of service assignment.

To solve the problems mentioned above, a fine-grained crowdsourcing model has been proposed in this paper to realize sensing task assignment, which could recommend a set of appropriate service providers to the requesters. The whole system model is showed in the Fig. 1. The key contributions of this paper are listed as following.  In order to realize the reasonable distribution of sensing tasks, ontology theory is used in this paper to describe various attributes of users. Based on the ontology method, all service providers and requesters could be accurately differentiate.  According to the fine granularity description of sensing requests, recommendation index of each service provider is calculated through similarity analysis.  In addition, as a method to obtain approximate discrete valued target function, decision tree is also referred in the crowdsourcing model to predict potential attribute with the information entropy principle.

Characteristic Ontology Triple

The characteristic data should be defined in a unified form which could accurately describe the user behavior. Ontology, a concept data structure, could effectively show the user characteristic through concept sets and concept relationship set. In order to meticulously classify all service providers and customers, the characteristic is supposed to be defined as ontology triple . The components in each type of property are listed as following.

Common Property. For a service provider or a customer, common property contains user ID, user name and additional information. Therefore, the common property is defined as a set, Particularly, all the elements of CP are text attributes, even though ID is the serial number of a service provider or a customer.

Exclusive Property. The exclusive property contains the location, availability and the service type. The Exclusive Property is defined as following set,

The element is the GPS location of the service provider or the customer. The element describes whether the user could participate in the service. The type service providers specializing in or customers requesting for is selected from the type catalogue set which is determined subjectively in relation to the application scenarios.

QoS Property. Although the exclusive properties will affect the decision of users, the Quality of Service properties are the determinant in service selection. The QoS property is defined as following set, Both of and are matrixs, where N is the number of service providers and M is the number of customers. Every element in the matrix is valued in the range of .

Recommendation Method

Recommendation Index

The recommendation index shows selection priority of a service provider. The higher recommendation index of a service provider, the more likely the user will enjoy the service.

Definition 1 (User Behavior Similarity).

In the service selection method, the user behavior similarity between two users is calculated by valuation offered in the past service. The user behavior similarity between two users and could be calculated through following equation,

                        , 2 2 , , , , , , ,
k i j k i j k i j i j SP OS U U user i j i j SP OS U U SP OS U U Val i k Val Val j k Val Sim U U Val i k Val Val j k Val         (1)
Where consists of the service providers which have both served users and , is the valuation that user values on the service provider .

Except for User Behavior Similarity, the types of service the service providers specialize in will also impact the validity of the valuation. Concept Hierarchy Tree [START_REF] Ye | Research on a data mining algorithm based on concept hierarchy tree and its application in CRM[END_REF] is proposed to processing the service types.

According to the definition of the information content in the information theory [START_REF] Resnik | Using Information Content to Evaluate Semantic Similarity in a Taxonomy[END_REF], the present information content of service type will be calculated through the method proposed by the Resnik, and the calculate method is mentioned as following,  

( ) log x x IC Type p Type  (2)
Definition 2 (Type Similarity). After obtaining the information content of each type, the type similarity between and will use the following formula to calculate,

                        , , , 2 ,
i j i j sp i j i j i j i j

IC A Type Type d A Type Type Sim Type Type IC Type IC Type d Type d Type d A Type Type

     (3)
Where is the nearest common ancestor concept of and .

, ( ) and are the depth of the , and the nearest common ancestor concept of and .

Recommendation index. Using user behavior similarity and type similarity, the recommendation index of service provider to user is calculated as following,

          1 , , , ( , , ) n
i j k sp j k user i x x RI U SP Type Sim Type SP Type Sim U U Val i j     (4)
Where ( ) is the Type property concluded in the exclusive property of , is the service type which user requires, and n is the number of the users.

Service Decision Tree

Decision tree is an appropriate prediction model used in the process of selecting service provider. The attributes considered in the service decision tree conclude following data: Availability, Distance, Type and Success Rate.

After building the basis of the service decision tree by using the Ontology Triple, the rest of the service decision tree will be built through recursive process described as following step:

Step 1. Acquiring the data of node which is going to build the its subtree, then defining the sample set remained as , the attributes remained in the attribute list as

Step 2. Check whether the attributes Valuation of samples in set are the same. If the valuation attributes of samples in are the same, this node should be marked as the leaves of the service decision tree. Conversely, if not all the valuation of samples is identical, turn to step 3.

Step 3. Check the attribute list. If the attribute list is empty, marking the major valuation attributes of the samples in as the leaves of the service decision tree. Conversely, if the attribute list is not empty, turn to step 4.

Step 4. Find the most appropriate attribute as the decision factor and the sample into different child nodes of and remove from .

Step 5. Recur Step 1 with the child node , and are the sample set and attribute list in the node .

Experiment

In the experiment, we have tested the fine-grained crowdsourcing model through data set collected from www.dianping.com. The data collects from 5000 customers and 100 service providers.

To show the recommendation process choosing the suitable service provider and improving the expect evaluation. In this experiment, Random Service Selection (RSS) and Convergent Population Diversity Handling Genetic Algorithm (CoDiGA) [START_REF] Ma | Quick convergence of genetic algorithm for QoS-driven web service selection[END_REF] are chosen to compare with Fine-grained Crowdsourcing Model (FCM). The experiment result shows in the Fig. 2. Fig. 2. Customers' average evaluation on recommended service providers after using RSS, CoDiGA and FCM in different sets of customers.

The recommendation effect of service decision tree should also be compared with recommendation index in the experiment. Hence, we choosing different numbers of service providers which have the highest prediction value gained from decision tree as the recommended service providers, in the meantime, choosing the same amount of service providers owns the highest recommendation index and choosing the same amount of service providers received the highest average valuation. The experiment result shows in the Fig. 3. In the fine-grained crowdsourcing model, different coefficient could lead to different results. The number of the customers, the number of the service providers and the number of clusters the server preparing to divide will impact the effectiveness of model. In order to affirm this model could be efficient in various environments. The experiment applies the fine-grained crowdsourcing model in different scales of customers and service providers, and gains the average valuation that customers score on the recommended service providers. The relationship between the scale of data and the average evaluation is showed in the Fig. 4-Fig. 5.

Conclusion

Taking the advantage of ontology theory and decision tree in the service recommendation could highly increase the customer satisfaction. The fine-grained crowdsourcing model discussed in this paper offering a good way to choose appropriate service providers.

In addition, we can find that if user needs a large amount of service provider, the service decision tree performs better, and if only a few service providers are needed, according to the recommendation index is a better choice. Both using service decision tree and recommendation index in the service recommendation could bring a perfect result to user.

Fig. 1 .

 1 Fig. 1. Structure of service crowdsourcing model and whole process of recommending.

Fig. 3 .

 3 Fig. 3. The average scores valued by customers after accepting different number of recommended service providers through different recommendation methods.

Fig. 4 .Fig. 5 .

 45 Fig.4. The effect of model on different data scales (M=2)

Acknowledgments. This study was partly supported by NSFC under Grant (No.61502380, No. 61472316), Natural Science Basic Research Plan in Shaanxi Province of China (No.2014JQ8322), The Fundamental Research Funds for the Central Universities (No.XJJ2014049), and Shaanxi Science and Technology Innovation Project (2013SZS16-Z01/P01/K01), The Key projects of science and technology plan in Fujian province (No.2016H0029), Projects for Young Teacher in FuJian province (No.JA15401).


Jiali Wang

Hang Zhuang
email: zhuangh@mail.ustc.edu.cn

Changlong Li

Hang Chen
email: hangc@mail.ustc.edu.cn

Bo Xu

Zhuocheng He

Xuehai Zhou
email: xhzhou@ustc.edu.cn

A Fast and Better Hybrid Recommender System Based on Spark

Keywords: Recommender System, Hybrid, Weight, Spark

With the rapid development of information technology, recommender systems have become critical components to solve information overload. As an important branch, weighted hybrid recommender systems are widely used in electronic commerce sites, social networks and video websites such as Amazon, Facebook and Netflix. In practice, developers typically set a weight for each recommendation algorithm by repeating experiments until obtaining better accuracy. Despite the method could improve accuracy, it overly depends on experience of developers and the improvements are poor. What worse, workload will be heavy if the number of algorithms rises. To further improve performance of recommender systems, we design an optimal hybrid recommender system on Spark. Experimental results show that the system can improve accuracy, reduce execution time and handle large-scale datasets. Accordingly, the hybrid recommender system balances accuracy and execution time.

Introduction

Along with the popularization of the Internet, a sharp increase in the amount of data leads to information overload [START_REF] Eppler | The concept of information overload: A review of literature from organization science, accounting, marketing, mis, and related disciplines[END_REF]. Thus, recommender systems [START_REF] Cosley | Is seeing believing?: how recommender system interfaces affect users' opinions[END_REF] were proposed to relieve the stress of massive data. To improve recommender systems performance, researchers put forward the weighted hybrid method. Despite performance boost has been brought by the method, there are still several problems affecting performance, including weight setting and computation load. Hence, we implement a weighted hybrid recommender system on Spark. In the system, we design a new method to compute weights, using cluster analysis and user similarity. Besides, the execution time can be reduced by deploying the system on Spark.

Hybrid Recommender Systems

Hybrid recommender systems combine two or more recommendation algorithms to overcome weaknesses of each algorithm. It is generally classified as Switching, Mixed, Feature Combination, Meta-Level, and Weighted [START_REF] Burke | Hybrid recommender systems: Survey and experiments[END_REF].

The weighted hybrid technique combines different algorithms with different weights [START_REF] Burke | Hybrid recommender systems: Survey and experiments[END_REF]. The main idea is that the algorithm with better accuracy has a higher weight. At present, developers always set a weight for an algorithm manually and repeat experiments until achieving superior accuracy. Thus, the method depends on developers' experience to determine accuracy of an algorithm in different datasets. Due to large-scale datasets, sparsity of rating data and the number of algorithms, it's generally hard to obtain appropriate weights. Eventually the improvements of accuracy are poor.

In addition, to improve user experience, the system should return recommendation results efficiently. In other words, it has to quickly locate information which can appeal users in massive data. Thus, execution time is another evaluation standard of performance. However, the weighted hybrid technique needs to execute two or more algorithms and compute hybrid results, it's tough to reduce execution time.

Apart from accuracy and execution time of the system, scalability is also an important consideration. With the increasing of data scale and the algorithm complexity, the system requires more storage space and computing resources. It's difficult to meet the actual demand by only optimizing algorithms.

To address the above-mentioned issues, we design a hybrid recommender system on Spark. In the system, we propose an optimized method to improve accuracy. It computes weights and hybrid results based on cluster analysis and user similarity. Meanwhile, we deploy the system on Spark which is a fast and general engine for large-scale data processing [START_REF] Zaharia | Spark: Cluster computing with working sets[END_REF] to accelerate the training process and improve scalability.

Work of Paper

The rest of this paper is organized as five sections. Section 2 reviews recommendation algorithms and introduces the Spark. Section 3 describes the design of the optimized method. Section 4 shows how we implement the system on Spark. Section 5 gives experimental results and our analysis. Section 6 presents our conclusions and future work.

Related Work

In this section, we first review and compare recommendation algorithms and recommender systems. Then, we briefly analyze predicting ratings of algorithms. Finally, we introduce the distributed computing platform Spark and compare Hadoop and Spark.

Recommender Systems

Recommendation algorithms are the basis of recommender systems. In this section, we first introduce several representative algorithms.

Collaborative recommendation is almost the most popular algorithm. Based on overlapped ratings, it computes similarities among users. And then, it uses similarities to predict the rating that the current user on an item [START_REF] Burke | Hybrid systems for personalized recommendations[END_REF]. Tapestry [START_REF] Goldberg | Using collaborative filtering to weave an information tapestry[END_REF], Ringo [START_REF] Shardanand | Social information filtering: algorithms for automating word of mouth[END_REF] and GroupLens [START_REF] Resnick | Grouplens: an open architecture for collaborative filtering of netnews[END_REF] are typical systems with the algorithm.

Content-based recommendation pays attention to connections between items. It analyses descriptions of items that have been rated by users [START_REF] Pazzani | A framework for collaborative, content-based and demographic filtering[END_REF] and calculates similarities between items. The represention of an item's feature and the way to classify a new item are two important sub-problems [START_REF] Pazzani | A framework for collaborative, content-based and demographic filtering[END_REF].

Demographic-based recommendation [START_REF] Pazzani | A framework for collaborative, content-based and demographic filtering[END_REF] is a simple algorithm. It focuses on types of users that like a certain item. The technique identifies features of users such as age, gender, nationality, education, etc. It measures user similarity by taking those features into consideration. Table 1. shows strengths and weaknesses of each algorithm [5][9]. As the simple and effective technique, the weighted hybrid recommender system has been widely used in numerous fields. P-Tango and Pazzani are two typical systems. P-Tango is an online news system. It combines collaborative and content-based recommendation algorithms. The system adjusts weights of algorithms in the process of operation. Until the system obtains the expected accuracy, it determines weights. Pazzani is the other weighted hybrid recommender system. It combines collaborative, content-based and demographic-based recommendation algorithms. The system uses voting to determine recommendation results.

Weight Analysis

As previously described in section 1, we give the formalized representation of the weighted hybrid technique as follows:

R ui = n j=1 α j r j ui (1)
where j represents the j'th algorithm, it ranges from 1 to n. α j corresponds to the weight of the j'th algorithm. r j ui is the predicting rating of user u on item i by the j'th algorithm. R ui indicates the final hybrid result. From the formula (1), we can recognize that each algorithm just has a certain weight. That means the technique presupposes that predicting ratings of an algorithm are all greater or less than their ratings. However, this condition evaluates to false. Here we give some empirical evidence. We implement the User-based Collaborative Filtering (User-CF) and the Alternating Least Squares (ALS) in Python2.7, and use MovieLens-100K as observed data. In the Table 2, countH is the number of predicting ratings which are greater than real ratings. The countL is less than real ratings and countE is equivalent amounts. From the empirical results, we know that:

(1) In these algorithms, there are little predicting ratings that equal to ratings.

(2) A part of predicting ratings are greater than ratings, and another are less than ratings. (3) Only a weight for an algorithm may affect accuracy. Thus, it is essential to optimize weights.

Spark

Spark is a fast and general-purpose cluster computing platforms for large-scale data processing [START_REF] Zaharia | Spark: Cluster computing with working sets[END_REF] which is developed by UC Berkeley. In the environment of Spark, it includes Spark SQL [START_REF] Armbrust | Spark sql: Relational data processing in spark[END_REF], Spark Streaming [START_REF] Zaharia | Discretized streams: an efficient and fault-tolerant model for stream processing on large clusters[END_REF], Mllib [START_REF] Meng | Mllib: Machine learning in apache spark[END_REF], GraphX [START_REF] Xin | Graphx: A resilient distributed graph system on spark[END_REF], etc. Based on resilient distributed dataset (RDD) [START_REF] Zaharia | Resilient distributed datasets: A fault-tolerant abstraction for in-memory cluster computing[END_REF], it achieves memorybased computing, fault tolerance and scalability. Currently, Spark is deployed in Amazon, ebay and Yahoo! to process large-scale datasets.

For a hybrid recommender system, performance is affected by data scale, the number of algorithms and the complexity of algorithms. Deploy the system on Spark can mitigate above affects.

(1) In the system, large-scale datasets could be stored in distributed storage.

(2) Algorithms are independent with each other, they are supposed to be performed in parallel. (3) Intermediate values can be cached in memory to decrease execution time.

Therefore, in this paper, we design an optimized hybrid recommender system on Spark.

Design Overview

The empirical evidence from section 2 suggests that accuracy still has chance to be improved. The predicting ratings are higher or lower than corresponding ratings. Thus, we use cluster analysis to obtain more accurate weights. The principle of cluster analysis is that according to the properties of samples, using mathematical methods to determine relationship between samples, and according to the relationship to cluster samples. Based on cluster analysis, we present an optimized method for calculating personalized weights. Now let us discuss the method in detail.

Objective Function.

In this section, we first give explanations of several concepts. In the following statement:

1. Assume that there are n algorithms in the system and j is the j'th algorithm. 2. u for user, i for item and (u,i) represents the data item of u and i. 3. R ui is the rating of u on i, r j ui is the predicting rating of u on i which is computed by the j'th algorithm. 4. For the j'th algorithm, the error between the rating and the predicting rating is:

D j ui = R ui -r j ui .
In order to reduce n j=1 u,i D j ui , similar errors are expected to get same weights. Based on errors, we divide (u,i) into k clusters and design C ui = (c 1 , c 2 , • • • , c k ) to reflect the cluster of (u,i). For the j'th algorithm, α j = (α j1 , α j2 , • • • , α jk ) represents k weights of the algorithm. α j C T ui finally determines the weight for r j ui .

According to our analysis, we define the objective function as formula (2):

F (α) = u,i (R ui -α 1 C ui T r ui 1 -α 2 C ui T r ui 2 -• • • -α n C ui T r n ui ) 2 (2)
s.t. n j=1 α j C T ui = 1 (3)

Weight Calculation

According to D j ui , the optimized method classifies all (u,i) into k clusters . For each (u,i), it has a vector

C ui = (c 1 , c 2 , • • • , c k ) and is initialized to C ui = (0, 0, • • • , 0)
. The value which corresponds to (u,i)'s cluster is set to 1. For instance, if (u,i) belongs to the cluster 2, C ui = (0, 1, 0, • • • , 0). The weight for r j ui is α j2 which is computed by α j C ui T . Therefore, C could map weights to predicting ratings and achieve multiple weights for an algorithm. Fig. 1 shows the pipeline of the method.

After calculating C, the optimized method requires to compute α j . For the purpose of minimizing the objective function, we make use of the Lagrange theory and minimum theory [START_REF] Borneas | On a generalization of the lagrange function[END_REF] [START_REF] Mitra | Estimating surface normals in noisy point cloud data[END_REF]. Based on formula (2), the method constructs the Lagrange function L(α).

L(α) = F (α) + λ u,i φ(α) (4)
φ(α) = n j=1 α j C ui T -1 (5)
For each j, let ∂L ∂(αj Cui T ) = 0. We can get an equation:

2 * u,i (α 1 C ui T r 1 ui r j ui + α 2 C ui T r 2 ui r j ui + • • • + α n C ui T r n ui r j ui ) + λ = 2 * u,i R ui r j ui (6)
The equation ( 6) can be represented by matrix:

XY = 2 * α 1 α 2 • • • α n λ *        u,i C ui T r 1 ui r 1 ui u,i C ui T r 1 ui r 2 ui • • • u,i C ui T r 1 ui r n ui u,i C ui T u,i C ui T r 2 ui r 1 ui u,i C ui T r 2 ui r 2 ui • • • u,i C ui T r 2 ui r n ui u,i C ui T . . . . . . . . . . . . . . . u,i C ui T r n ui r 1 ui u,i C ui T r n ui r 2 ui • • • u,i C ui T r n ui r n ui u,i C ui T 1 1 • • • 1 0        = 2 *        u,i R ui r 1 ui u,i R ui r 2 ui . . . u,i R ui r n ui u,i 1        = R (7)
Thus the weight matrix X can be calculated by

X = R * Y -1 (8)
The optimized method uses ratings which have already stored in the system to compute weights. However, these weights aren't entirely appropriate for a new (u,i). We further introduce user similarity to compute weights. The user similarity is computed by cosine similarity:

sim u,v = |N (u) ∩ N (v) | |N (u) ||N (v) | (9)
where sim u,v is the similarity between u and v. N (u) means the number of items that u have rated. N (v) is the same as N (u) . For the (u , I ), the optimized method calculates the hybrid result as:

ru I = v sim u ,v * (α 1 C vI T r 1 u I + α 2 C vI T r 2 u I + • • • + α n C vI T r n u I ) v sim u v
(10) The equation ( 10) is able to filter the interference of non similar weights and get a personalized weight for the (u , I ). According to the design overview, we deploy the hybrid recommender system on Spark. The system contains data storage, prediction, cluster, weight, model fusion and recommendation, totally 6 modules. Fig. 2 shows the architecture of the system.

Modules

Data storage module is the basis of the system. It stores input data, including historical data and ratings. We use HDFS which is a distributed file system to store raw data [START_REF] Borthakur | The hadoop distributed file system: Architecture and design[END_REF]. The pre-processed data are put in the database such as HBase, Redis, Hive, etc [START_REF] Zhang | Research on hadoop-based enterprise file cloud storage system[END_REF][19] [START_REF] Thusoo | Hive: a warehousing solution over a map-reduce framework[END_REF]. Topside modules read data from the database. Prediction module is used to compute predicting ratings. It performs recommendation algorithms in parallel. Outputs are predicting ratings.

The cluster module concentrates on errors of (u,i). It exploits k-means to classify (u,i). Output of the module is C. The weight module accepts C to compute weights. With C and α, the module can get a weight for each r j ui . Output of it is α.

The model fusion calculates hbrid results based on predicting ratings, C, α and user similarity. According to these parameters, it determines hybrid results by logistic regression [START_REF] Tsukimoto | Logical regression analysis: from mathematical formulas to linguistic rules[END_REF]. Recommendation is used to recommend items for users. Based on hybrid results, it generates recommendation lists. Besides, it also outputs an evaluation for results.

Discussion

In the hybrid recommender system on Spark, data are translated into RDDs. Because of the characteristics of memory-based computing and parallel operations, RDDs can be processed in parallel to reduce execution time. The read-only and fault tolerance of RDD make the system more reliable. Besides, due to the distributed storage of Spark, the system is able to handle large-scale datasets. It improves scalability of the system. Therefore, deploy the hybrid recommender system on Spark could decrease execution time and further improve scalability.

Performance

Evaluation Index

Accuracy. The system accuracy is measured by root mean square error (RMSE) [START_REF] Willmott | Advantages of the mean absolute error (mae) over the root mean square error (rmse) in assessing average model performance[END_REF]. It is defined as:

RM SE = u,i∈T (R ui -r ui ) 2 |T | (11)
where R ui and r ui is the rating and the hytbrid result that u on i respectively.

|T | denotes the number of r ui .

Execution Time. The execution time includes time of algorithms, clustering, calculating weights and hybrid results. It is measured in minutes.

Experimental Setup

In this experiment, we choose Spark as our platform. All experiments were performed using a local cluster with 7 nodes (1 master and 6 worker nodes): each node has Xeon(R) dual-core 2.53GHz processor and 6GB memory.

Dataset. In Table 3, we list datasets that were used in the experiment. For each dataset, we divide it into 2 training sets and a test set randomly. Algorithms. We implement 3 recommendation algorithms: User-CF, Itembased Collaborative Filtering (Item-CF) and ALS. We perform them in training sets and test sets to compute predicting ratings, weights and hybrid results.

Nodes. We compare execution time of the stand-alone system and the distributed system. For the former, we use the server with Xeon(R) dual-core 2.53GHz processor and 6GB memory. For the latter, we use a local cluster with 7 nodes (1 master and 6 worker nodes): each node has Xeon(R) dual-core 2.53GHz processor and 6GB memory.

Performance Comparision

In this section, we evaluate performance of the hybrid recommender system on Spark, including accuracy and execution time.

Fig. 3 shows impacts of data scales on accuracy. In the experiment, we performe the combination of User-CF and ALS on four MovieLens datasets. In the Fig. 3 sparsity of MovieLens-700K, the hybrid recommender system obtains the best result. Compare with User-CF and ALS, the system improves accuracy of 8.21 Fig. 4 gives the RMSE of different types of datasets. In the experiment, we performe the combination of User-CF and ALS on MovieLens-400K and BookCrossing. The Fig. 4 shows that the hybrid recommender system can improve accuracy of different types of datasets. And there are significant improvements on BookCrossing. The improvements demonstrate that the system is available for sparse datasets. Fig. 5 shows correlations between accuracy and combinations of algorithms. In the experiment, four combinations of algorithms are performed on MovieLens-100K, MovieLens-400K and BookCrossing respectively. In Fig. 5, the hybrid recommender system obtains better accuracy than single algorithm. When accura- cy of single algorithm is favorable, the hybrid recommender system also obtains better accuracy. Fig. 6 compares execution time of stand-alone mode and local cluster mode. The experiment performs the combination of User-CF and ALS on MovieLens-100K to MovieLens-1M. For the stand-alone system, execution time increases sharply with the expansion of data scale. However, execution time of local cluster mode remains relatively constant. When the data scale is larger than MovieLens-900K, the stand-alone mode couldn't handle it. The local cluster mode could handle MovieLens-10M or larger datasets. From Fig. 6, we can recognize that memory-based computing, parallel operations and distributed storage of Spark are helpful to decrease execution time and improve scalability.

Conclusion and Future Work

Improving performance of recommender systems is a crucial solution for information overload. This paper designs a new weighted hybrid recommender system to solve this problem. We are the first to compute weights by using cluster analysis, user similarity and minimum theory. Besides, we deploy the hybrid recommender system on Spark. The system improves accuracy by optimizing weights and reduces execution time by memory-based computing and parallel operations. And distributed storage of the system is helpful to improve scalability. The experiment results demonstrate the performance of our hybrid recommender system.

In future work, we will consider to improve and extend the system: expansion of algorithm to process more complex scenes. Further research on factors influencing weights to improve accuracy. Meanwhile, optimize the implementation of the system on Spark.

Fig. 1 .

 1 Fig.1. The pipeline of the optimized method. The input file consists of ratings. Algorithms read the input file and output predicting ratings. Then the system computes errors and cluster data items. Finally the system gives the C.

Fig. 2 .

 2 Fig.2. The architecture of the hybrid recommender system. The system reads ratings and outptu recommendation lists. Besides, it also provides an evaluation index.

Fig. 5 .

 5 Fig. 5. The RMSE of different combinations of algorithms. The x-axis indicates the dataset, and the y-axis describes the RMSE.

Fig. 6 .

 6 Fig. 6. The execution time of 2 modes. The x-axis indicates datasets and the y-axis describes execution time.

Table 1 .

 1 Strengths and weaknesses of recommendation algorithms

	Algorithm	Strength	Weakness
	Collaborative	Field independence. Not necessary to understand descriptions of items. Support users to discover potential interests.	New user problem. New item problem. Sparsity.
	Content	Improve accuracy by increasing dimensions of item features.	Cold start problem. Similarity measurement is one-sided.
	Demographic	Historical data are not necessary. Wide range of applications. No cold start problem.	The algorithm is rough and imprecise.

Table 2 .

 2 The results of statistic analysis on predicting ratings

	Algorithm countH countL countE
	User-CF 9181 10752	11
	ALS	6992 12952	0

Table 3 .

 3 Datasets in the experiment

	Dataset	Users Items Ratings
	MovieLens-100K 1000 1700 100000
	MovieLens-200K 1371 10153 200000
	MovieLens-300K 2004 10850 300000
	MovieLens-400K 2661 11634 400000
	MovieLens-500K 3462 13257 500000
	MovieLens-600K 4073 13488 600000
	MovieLens-700K 4753 14154 700000
	MovieLens-800K 5543 14230 800000
	MovieLens-900K 6207 14963 900000
	MovieLens-1M 6000 4000 1 million
	BookCrossing 71212 176272 400000

  , with the increasing of data scale, RMSE generally decreases. Due to the

				User-CF ALS Hybrid				User-CF ALS Hybrid
		1.10					2.38
		1.05					2.18
	RMSE	0.90 0.95 1.00				RMSE	1.98 1.38 1.58 1.78
		0.85	100K	400K MovieLens 700K	1M		1.18 0.98 0.78	MovieLens-400K	BookCrossing
	Fig. 3. The RMSE of different scale Movie-	Fig. 4. The RMSE of different types of
	Lens. The x-axis indicates datasets, and the	datasets. The x-axis indicates datasets, and
	y-axis describes the RMSE.		the y-axis describes the RMSE.