PIM-based computing device for accelerating combined operations

Abstract

To provide a PIM-based computing device for join operational acceleration in which an IDP maximizes join throughput and a PIM-enabled DIMM, instead of a CPU, performs global partitioning of two tables, thereby utilizing a high internal memory bandwidth and computation throughput of the PIM-enabled DIMM and optimizing communication between all-to-all IDPs.SOLUTION: A PIM-based computing device includes: a main memory configured to store first and second data sets; a PIM module configured as an in-line memory module including a memory processor, a working memory, and a memory bank, the memory processor being connected to the working memory and performing a join operation on the first and second data sets loaded into the working memory; and a processor configured to form first and second channels with the main memory and the PIM module, import the first and second data sets through the first channel, and store the data sets in the memory bank through the second channel.SELECTED DRAWING: Figure 1

Description

Technical Field

[0001] The present invention relates to computing technology based on PIM (Processing-In-Memory), and more particularly to an acceleration technique that can improve the performance of associative operations in PIM hardware.

Background Art

[0002] PIM is a next-generation semiconductor that integrates a memory and a processor arithmetic unit inside one chip. This can solve the problems of data bottleneck phenomena and excessive power consumption that occur in existing computing structures where the memory and the processor are separated.

[0003] In a Dual in-line memory module (DIMM) supported by PIM, an IDP (In-DIMM processor) is mounted near the memory bank. When using a DIMM supported by PIM, an application can offload calculations from the host CPU to the IDP. The IDP can directly access the data stored in the memory bank and perform calculations without using host-side resources. In this way, the DIMM supported by PIM reduces the data transfer volume between the CPU and the main memory, utilizes the high internal memory bandwidth and computing throughput of the DIMM supported by PIM, and greatly helps the application to overcome the data bottleneck phenomenon.

[0004] On the other hand, in order to improve the performance of a relational database management system (RDBMS), a technique for accelerating relational operations using a DIMM supported by PIM has been proposed. The prior art relied on a PIM hardware simulator and architectural assumptions for PIM hardware. However, the architectural assumptions of the prior art may have a problem of being incompatible with the DIMM as they are different from the architecture of the DIMM supported by PIM.

[0005] Therefore, to realize the high potential of PIM in actual computing systems, a new coupling algorithm optimized for PIM-supported DIMMs is needed. [Prior art documents] [Patent Documents]

[0006] [Patent Document 1] Republic of Korea Publication Patent No. 10-2022-0094180 [Overview of the project] [Problems that the invention aims to solve]

[0007] The present invention provides a PIM-based computing device for accelerating coupled operations, which is a PIM coupled computation acceleration technique compatible with DIMM architectures, maximizes coupled throughput in IDPs, performs global partitioning of both tables in PIM-supported DIMMs instead of CPUs, leverages the high internal memory bandwidth and computation throughput of PIM-supported DIMMs, and optimizes many-to-many IDP communication. [Means for solving the problem]

[0008] In one embodiment, a PIM-based computing device for accelerating join operations comprises a main memory for storing first and second datasets, an inline memory module including a memory processor, working memory, and a memory bank, wherein the memory processor is connected to the working memory and includes a PIM (Processing-In-Memory) module that performs join operations on the first and second datasets loaded into the working memory, and a processor that forms a first channel and a second channel with the main memory and the PIM module, and takes in the first and second datasets in the first channel and stores them in the memory bank in the second channel.

[0009] The PIM module can implement the inline memory module using multiple dual inline memory modules and support scale-out join operations.

[0010] The PIM module can subgroup the multiple dual inline memory modules and connect them to one of the second channels.

[0011] The PIM module can perform global partitioning via inter-data relocation between the plurality of dual inline memory modules to arrange the first and second datasets.

[0012] The PIM module can perform RnS (Rotate-and-Stream) operations on the data in the memory bank of each of the multiple dual inline memory modules during the mutual data transfer process.

[0013] The PIM module can send and receive specific datasets in units of datasets between the multiple dual inline memory modules and perform the RnS (Rotate-and-Stream) operation.

[0014] The processor may perform an Unordered Scatter-Gather (USG) operation on the first and second data sets in the main memory before performing the global partitioning, and then store them in the memory bank.

[0015] The processor can perform a Stream Transpose Send (STS) procedure during the USG calculation process and distribute the first and second datasets to the memory bank.

[0016] The memory processor can select a hash join, a sort-merge join, or a nested-loop join as the join operation.

[0017] The memory processor can generate a hash table in the working memory, move the first and second datasets in the memory bank to the working memory, perform local partitioning, and then perform hash joins by associating them with the hash table using a hash function.

[0018] The memory processor can perform sort-merge joins on the first and second datasets in the memory bank by performing range partitioning, parallel in-place sort, and parallel merge.

[0019] The memory processor can move the first and second data sets in the memory bank to the working memory, perform internal loop operations on each independently, and then perform nested loop connections.

[0020] In one embodiment, the PIM-based computing device includes a main memory for storing first and second source datasets, a plurality of inline memory modules each comprising a memory processor, working memory, and a memory bank, the memory processor being connected to the working memory and performing combination operations on the first and second processed datasets loaded into the working memory, and a processor for preprocessing the first and second source datasets and storing them in the memory bank.

[0021] The processor can perform an Unordered Scatter-Gather (USG) operation on the first and second source datasets in the main memory and store them in the memory bank.

[0022] The plurality of PIM modules can arrange the first and second processed data sets through inter data relocation for the preprocessed first and second source data sets.

[0023] The plurality of PIM modules can perform RnS (Rotate-and-Stream) operations on the preprocessed first and second source data sets during the inter data relocation process.

Advantages of the Invention

[0024] The disclosed technology can have the following advantages. However, it should not be understood that the scope of the rights of the disclosed technology is limited thereby, as it does not mean that a specific embodiment must include all of the following advantages or only the following advantages.

[0025] The PIM-based computing device for the acceleration of the join operation according to the present invention is a PIM join operation acceleration technique compatible with the DIMM architecture, which maximizes the join throughput with IDP, performs global partitioning of both tables with DIMM instead of the CPU, and can utilize the high internal memory bandwidth and computing throughput of DIMM, and optimize the many-to-many IDP communication.

Brief Description of the Drawings

[0026] [Figure 1] It is a diagram showing a PIM-based computing device according to an embodiment of the present invention. [Figure 2a] It is a diagram for explaining the process of performing a join operation with the memory processor of FIG. 1. [Figure 2b] It is a diagram for explaining the process of performing a join operation with the memory processor of FIG. 1. [Figure 2c] It is a diagram for explaining the process of performing a join operation with the memory processor of FIG. 1. [Figure 3] This is a flowchart illustrating one embodiment of the coupling calculation method for a PIM-based computing device according to the present invention. [Figure 4] This diagram illustrates a dual inline memory module (DIMM) architecture with PIM support. [Figure 5] This figure shows the single-PID performance of various hash functions. [Figure 6a] This figure shows the join execution delay times for different algorithms: hash joins, sort-merge joins, and nested-loop joins. [Figure 6b] This figure shows the join execution delay times for different algorithms: hash joins, sort-merge joins, and nested-loop joins. [Figure 7] This diagram illustrates the execution process of the RnS (Rotate-and-Stream) operation. [Figure 8] This figure shows the bandwidth of RnS mutual IDP communication compared to the existing method. [Figure 9] This diagram illustrates the execution process of the USG (Unordered Scatter-Gather) operation. [Figure 10] This figure shows the working model of PID-JOIN according to the present invention. [Figure 11] This figure illustrates the experimental results according to the present invention. [Figure 12] This figure illustrates the experimental results according to the present invention. [Figure 13] This figure illustrates the experimental results according to the present invention. [Figure 14] This figure illustrates the results of sensitivity research according to the present invention. [Figure 15] This figure illustrates the results of sensitivity research according to the present invention. [Figure 16] This figure illustrates the results of sensitivity research according to the present invention. [Figure 17a] This figure illustrates the results of sensitivity research according to the present invention. [Figure 17b] This figure illustrates the results of sensitivity research according to the present invention. [Modes for carrying out the invention]

[0027] This invention was filed with support based on the following research topic. [This invention was supported by a National Research and Development Project of the Republic of Korea] [Project-Specific Number] 1711193986 [Project Number] 2020-0-01361-004 [Ministry / Agency Name] Ministry of Science and ICT [Name of the (specialized) organization for issue management] Information and Communications Planning and Evaluation Institute [Research Project Name] Training of Innovative Human Resources in Information, Communication, and Broadcasting [Research Project Title] Support for Graduate School of Artificial Intelligence (Yonsei University) [Contribution rate] 1 / 2 [Project Implementation Organization Name] Yonsei University Industry-Academia Cooperation Group [Research Period] 2023.01.01~2023.12.31 [This invention was supported by a National Research and Development Project of the Republic of Korea] [Project-Specific Number] 1711193187 [Project Number] 2021-0-00853-003 [Ministry / Agency Name] Ministry of Science and ICT [Name of the (specialized) organization for issue management] Information and Communications Planning and Evaluation Institute [Research Project Name] Development of Leading Technology for Novel Concept PIM Semiconductors [Research Project Title] Development of a Software Platform for PIM Utilization [Contribution rate] 1 / 2 [Project Implementation Organization Name] Yonsei University Industry-Academia Cooperation Group [Research Period] 2023.01.01~2023.12.31

[0028] The descriptions relating to the present invention are merely embodiments for structural or functional explanation; therefore, the scope of the rights of the present invention should not be construed as being limited by the embodiments described herein. That is, since embodiments can be modified in various ways and can take on various forms, the scope of the rights of the present invention should be understood to include equivalents that can realize the technical idea. Furthermore, the purposes or effects presented in the present invention do not mean that any particular embodiment should include all of them or only such effects; therefore, the scope of the rights of the present invention should not be understood as being limited by them.

[0029] On the other hand, the meanings of the terms used in this application should be understood as follows:

[0030] Terms such as "first," "second," etc., are used to distinguish one component from another, and these terms should not limit the scope of rights. For example, the first component may be named the second component, and similarly, the second component may be named the first component.

[0031] When it is mentioned that one component is "connected" to another, it must be understood that it may be directly connected to the other component, or that other components may exist in between. Conversely, when it is mentioned that one component is "directly connected" to another, it must be understood that there are no other components in between. On the other hand, other expressions describing the relationship between components, namely "between" and "immediately between," or "adjacent to" and "directly adjacent to," must be interpreted in the same way.

[0032] Unless the context clearly indicates otherwise, singular expressions should be understood to include plural expressions, and terms such as “includes” or “possesses” should be understood to indicate the existence of a particular feature, number, step, action, component, part, or combination thereof, without prejudice to the existence or possibility of adding one or more other features, numbers, steps, actions, components, parts, or combination thereof.

[0033] In each step, the identification codes (e.g., a, b, c, etc.) are used for explanatory purposes only and do not indicate the order of the steps. The steps may occur in a different order than specified unless the context clearly indicates a specific order. That is, the steps may occur in the same order as specified, substantially simultaneously, or in the reverse order.

[0034] The present invention is implemented as computer-readable code on a computer-readable recording medium, and the computer-readable recording medium includes all types of recording devices that store data to be read by a computer system. Examples of computer-readable recording media include ROM, RAM, CD-ROM, magnetic tape, floppy disk, and optical data storage devices. Furthermore, the computer-readable recording medium can be distributed across a network of connected computer systems, and computer-readable code can be stored and executed in a distributed manner.

[0035] All terms used herein have the same meaning as that generally understood by a person of ordinary skill in the art to which this invention pertains, unless otherwise specifically defined. Terms defined in commonly used dictionaries should be interpreted in accordance with their meaning in the context of the relevant art and should not be interpreted as having an ideal or overly formal meaning unless explicitly defined herein.

[0036] Figure 1 shows a PIM-based computing device according to one embodiment of the present invention.

[0037] Referring to Figure 1, the PIM-based computing device 100 can be implemented including a main memory 110, a PIM (Processing-In-Memory) module 130, and a processor 150.

[0038] The main memory 110 can store the first and second datasets.

[0039] The PIM module 130 consists of an inline memory module that includes a memory processor 131, working memory 133, and memory bank 135.

[0040] The memory processor 131 is connected to the working memory 133 and can perform join operations on the first and second data sets loaded into the working memory 133. The PIM module 130 can implement the inline memory module as multiple dual inline memory modules (DIMMs) and support scale-out join operations. In one embodiment, the DIMM can consist of a set in which multiple memory processors 131, working memory 133, and memory banks 135 are interconnected.

[0041] The processor 150 forms first and second channels with the main memory 110 and the PIM module 130, and can acquire first and second data sets in the first channel and store them in the memory bank 135 in the second channel.

[0042] The PIM module 130 can subgroup multiple dual inline memory modules (DIMMs) and connect them to one of the second channels. The PIM module 130 can perform global partitioning via inter data relocation between multiple dual inline memory modules (DIMMs) to arrange the first and second data sets.

[0043] Before performing global partitioning, the processor 150 can perform an Unordered Scatter-Gather (USG) operation on the first and second data sets in the main memory 110 and store them in the memory bank 135. The processor 150 can distribute the first and second data sets evenly to the memory bank 135 via the USG operation. The processor 150 can perform a Stream Transpose Send (STS) procedure during the USG operation to distribute the first and second data sets to the memory bank 135.

[0044] Since the PIM module 130 correctly positions the first and second datasets via global partitioning, and because the input and output datasets in memory bank 135 do not require a strict order, the first and second datasets must be distributed evenly across all memory processors 131. USG (Unordered Scatter-Gather) can be implemented as a mechanism to evenly distribute the first and second datasets to memory bank 135. Processor 150 can stream the first and second datasets in main memory 110 during the USG calculation process, then transpose them and transfer them to the PIM module 130. Processor 150 can sequentially access both main memory 110 and the PIM module 130 to evenly distribute the first and second datasets to memory bank 135.

[0045] The PIM module 130 can perform RnS (Rotate-and-Stream) operations on data in the memory bank 135 of each of the multiple dual inline memory modules (DIMMs) during the process of mutual data transfer. The PIM module 130 can send and receive specific data sets in units of data sets between multiple dual inline memory modules (DIMMs) and perform RnS (Rotate-and-Stream) operations.

[0046] The memory processor 131 can select hash join, sort-merge join, or nested-loop join as the join operation. In one embodiment, the memory processor 131 can generate a hash table in the working memory 133, move the first and second data sets in the memory bank 135 to the working memory 133, perform local partitioning, and then perform a hash join by associating them with the hash table using a hash function. In one embodiment, the memory processor 131 can perform sort-merge join by performing range partitioning, parallel in-place sort, and parallel merge on the first and second data sets in the memory bank 135. In one embodiment, the memory processor 131 can perform nested-loop join by moving the first and second data sets in the memory bank 135 to the working memory 133 and independently performing internal loop operations on each of them.

[0047] The PIM-based computing device 100 can be implemented by including a plurality of PIM modules 130. In this case, each of the plurality of PIM modules 130 can consist of an inline memory module including a memory processor 131, working memory 133, and a memory bank 135. The main memory 110 can store the first and second source data sets. The processor 150 can preprocess the first and second source data sets and store them in the memory bank 135. Here, the processor 150 can store the first and second source data sets in the main memory 110 in the memory bank 135 after performing a USG (Unordered Scatter-Gather) operation as preprocessing. The memory processor 131 can perform a join operation on the first and second processed data sets that have been moved to the working memory 133.

[0048] Multiple PIM modules 130 can place the first and second processed datasets onto the pre-processed first and second source datasets via inter data relocation. The multiple PIM modules 130 can perform RnS (Rotate-and-Stream) operations on the pre-processed first and second source datasets during the inter data relocation process.

[0049] Figures 2a to 2c illustrate the process of performing concatenation operations in the memory processor shown in Figure 1.

[0050] Referring to Figures 2a to 2c, the memory processor 131 can perform join operations on the first and second datasets that have been brought up to the working memory 133. Here, the memory processor 131 can select and perform hash joins, sort-merge joins, or nested-loop joins as the join operations.

[0051] Hash join performs a join on the first and second datasets using a hash table. Specifically, referring to Figure 2a, the memory processor 131 generates a hash table in the working memory 133. The hash table consists of key-value pairs. Here, the key and value of the pair can correspond to the join key and the tuple ID of the first dataset, the R tuple, respectively. The memory processor 131 moves the first and second datasets (R, S) from memory bank 135 to the working memory 133 and performs local partitioning. Next, the memory processor 131 performs a hash join by associating them with the hash table 210 using a hash function. The memory processor 131 probes the hash table with the second dataset, the S tuple, identifies the R tuple in the entries of the hash table, and checks whether both tuples satisfy the given join condition.

[0052] Since the memory processor 131 frequently accesses entries in the hash table during the hash join process, generating the hash table in working memory prevents access to the slower memory bank 135.

[0053] The sort-merge join performs range partitioning, parallel in-place sort, and parallel merge on the first and second datasets (R,S) in memory bank 135. Specifically, referring to Figure 2b, the memory processor 131 sorts and merges the R and S tuples, which are the first and second datasets, and then performs the join operation. The memory processor 131 performs range partitioning to group tuples with the same join key value range in order to combine the work sets in the work memory 133. Next, the memory processor 131 performs parallel in-place sort by assigning different partitions to different threads and sorting the assigned partitions simultaneously. At this time, the memory processor 131 uses an internal sorting algorithm to ensure that the size of the work set does not exceed the capacity of the work memory 133. After that, the memory processor 131 merges the tuples of the S partition with the R partition.

[0054] Nested-loop coupling involves moving the first and second data sets in memory bank 135 into working memory 133 and performing independent internal loop operations on each. Specifically, referring to Figure 2c, the outer loop scans an S tuple, and the inner loop scans an R tuple. The memory processor 131 checks the buffer size in working memory 133 required to store the R tuple and allocates a circular buffer in working memory 133 for each thread to store the tuple. Next, the memory processor 131 moves the S tuple into working memory 133, filling the circular buffer with the R tuple, and then performs coupling on the R and S tuples.

[0055] Figure 3 is a flowchart illustrating one embodiment of the PIM-based computing device coupling calculation method according to the present invention.

[0056] Referring to Figure 3, the PIM-based computing device 100 can store the first and second datasets in the main memory 110 (step S310).

[0057] The PIM-based computing device 100 can form channels with the main memory 110 and the PIM module 130 via the processor 150 (step S330). The processor 150 forms first and second channels with the main memory 110 and the PIM module 130, and can take in first and second data sets in the first channel and store them in the memory bank 135 in the PIM module 130 in the second channel. The processor 150 can perform USG (Unordered Scatter-Gather) operations on the first and second data sets in the main memory 110 and store them in the memory bank 135.

[0058] The PIM-based computing device 100 can perform join operations on the first and second data sets that have been moved to the working memory 135 by the memory processor 131 of the PIM module 130 (step S350). The PIM module 130 can consist of an inline memory module including a memory processor 131, working memory 133, and memory bank 135. The PIM module 130 can implement the inline memory module as multiple dual inline memory modules and support scale-out join operations. The PIM module 130 can arrange the first and second data sets by performing global partitioning via inter data relocation between multiple dual inline memory modules. The PIM module can perform RnS (Rotate-and-Stream) operations on the data in the memory bank 135 of each of the multiple dual inline memory modules during the inter data relocation process. The memory processor 131 of the PIM module 130 can select and perform hash join, sort-merge join, or nest-loop join as the join operation.

[0059] The PIM-based computing device according to the present invention will be described in more detail below with reference to Figures 4 to 17b.

[0060] Figure 4 illustrates a dual inline memory module architecture with PIM support.

[0061] Referring to Figure 4, a PIM-enabled dual in-line memory module (DIMM) 400 can connect an In-DIMM Processor (IDP) 430 to a memory bank 410 and support the PIM module 130 of the present invention.

[0062] The DIMM400 consists of multiple chips, each chip comprising multiple banks. The DIMM400 groups the chips into ranks and stores sequential data fragments in one of the ranks. Here, a rank refers to a set of chips that share the same control signals. The DIMM400 can realize a PIM module 130 by connecting an IDP430 to each 64MB memory bank 410. The IDP430 cannot directly access the data stored in the memory bank 410 and must bring the data up as working memory (WRAM) 450.

[0063] Each IDP 430 cannot access data stored in memory bank 410 without first transferring it to WRAM 450; therefore, it cannot directly access data stored in other memory banks. The host CPU, processor 150, must transfer data between the memory banks 410 of different IDP 430s. Such CPU-based inter-IDP 430 communication uses limited host CPU bandwidth and generates significant performance overhead. Due to the rank-level parallel processing of DIMM 400, the performance overhead can become even greater without careful optimization of inter-IDP 430 communication. When transferring data between IDP 430s, all other IDPs belonging to the same rank as the source and destination IDPs must also transfer data, and IDPs of that rank operate in a synchronized manner. Therefore, one of the optimization guidelines for efficiently utilizing DIMM 400 is to minimize inter-IDP communication.

[0064] This invention presents a novel PIM-supported joining algorithm (hereinafter referred to as "PID-JOIN") that leverages the characteristics of the DIMM400 architecture to maximize the performance of join operations using PIM. The PID-JOIN method proposed in this invention can be applied to and evaluated in actual computing systems.

[0065] The PID-JOIN scheme proposed in this invention can perform join operations on first and second datasets. Here, the first and second datasets may correspond to R and S tuples, respectively. The R and S tuples are stored in memory bank 410. Three representative types of equi-join algorithms—hash join, sort-merge join, and nested-loop join—can be prototyped and evaluated as join operations. After designing the algorithms, their performance can be compared with the synthetic workload, and the best-performing algorithm can be selected. For performance comparison, it is assumed that IDP 430 processes the equi-join queries. Furthermore, it is assumed that all keys are initially stored in memory bank 410, and that tuples are stored in memory bank 410 as a result of the equi-join queries.

[0066] To achieve high coupling performance, a single IDP coupling algorithm must be optimized to the key architectural characteristics of the IDP, such as: Firstly, since the IDP can only access data stored in WRAM, the algorithm must fit its working set within the 64KB capacity of WRAM. If the size of the working set exceeds the WRAM capacity, the IDP may experience performance degradation due to frequent data transfers between WRAM and memory banks. Secondly, to maximize the IDP's internal memory bandwidth, the algorithm must transfer large sequential data chunks of at least 2KB in size. Since DRAM-based memory banks provide higher bandwidth for larger sequential data accesses, WRAM must be used to filter smaller or random data accesses. Thirdly, due to the IDP's limited computing capabilities, the algorithm must use operations that the IDP's basic hardware supports (e.g., integer addition). The IDP supports non-native operations such as integer multiplication / division and floating-point arithmetic via emulation, but achieves much lower computation throughput for non-native tasks.

[0067] The hash join algorithm uses a hash table to join both tables, namely R and S. The hash join algorithm first generates a hash table using tuples in R. The hash table consists of (key, value) pairs. The key and value of a pair correspond to the original join key and the tuple ID of the R tuple, respectively. Next, the hash join algorithm probes the hash table using tuples in S, identifies the R tuple in the entries of the hash table, and checks whether both tuples satisfy the given join condition.

[0068] A single IDP hash join algorithm leverages architectural properties to its fullest potential by generating a hash table in working memory, performing local partitioning in R due to the limited capacity of the working memory, then constructing and searching the hash table for each partition in R, and processing each partition. Constructing a single hash table containing all tuples in R results in the hash table occupying a memory bank, leading to memory bank access. The hash join algorithm can achieve high computational throughput by using IDP-friendly hash functions for work that IDP inherently supports. The hash join algorithm uses linear probing as a collision handling method.

[0069] Figure 5 shows the single-PID performance of various hash functions, displaying the hashing throughput and collision rate of the evaluated hash functions.

[0070] Referring to Figure 5, hash functions that utilize integer multiplication, namely Multiplicative Move (Mult), Multiplicative Addition (MultAdd), and Murmur, achieve low hash throughput due to the lack of support from the underlying IDP hardware. The other three hash functions, namely Tabulation, OAT, and CRC32, achieve much higher hash throughput because they utilize XOR-based operations. Furthermore, by measuring the number of hash collisions that occurred while constructing and searching a hash table with 128-K randomly generated tuples, we compared the collision rates of the hash functions and showed that OAT and CRC32 achieve lower collision rates than Tabulation. However, through additional measurements, we found that Tabulation's higher hash throughput corresponds to a higher collision rate and the fastest speed.

[0071] The sort-merge join algorithm first sorts and then merges the R and S tuples. The sort-merge join algorithm used in this invention performs range partitioning, parallel sorting, and parallel merging with respect to the structural characteristics of the IDP. Range partitioning is used to group tuples with the same range of join key values ​​to fit the work set in WRAM. In range partitioning, sorting a tuple in one partition is independent of tuples in other partitions, so that different partitions have their own unique range of join key values. Therefore, after different partitions are assigned to different threads via parallel sorting, the threads sort their assigned partitions simultaneously. An internal sorting algorithm is also used to ensure that the size of the work set does not exceed the WRAM capacity. The sort-merge join algorithm also utilizes multithreading support to load various R and S partitions into WRAM, then assigns R partitions to other threads in the IDP to merge the tuples of one S partition with all R partitions. This algorithm uses fast sorting.

[0072] The nested-loop coupling algorithm iteratively evaluates the coupling condition for all pairs of R and S tuples using nested loops. To maximize memory bank bandwidth, our nested-loop coupling algorithm minimizes data transfer between WRAM and the memory bank by allocating different amounts of WRAM to a buffer that temporarily stores the R tuples, depending on whether all R tuples can fit into WRAM. If all R tuples fit into WRAM, the nested-loop coupling algorithm allocates a WRAM buffer whose capacity matches the tuple capacity to load all R tuples into the buffer and utilizes the remaining WRAM buffer to stream the S tuples. Otherwise, it allocates most of the WRAM to store the S tuples and uses the remaining WRAM as a circular buffer for the R tuples.

[0073] The PID-JOIN method proposed in this invention can compare the latency of executing the join of three single IDP join algorithms on a single IDP in order to identify the single IDP join algorithm with the best performance.

[0074] Figures 6a and 6b show the delay times for different join execution algorithms: hash join, sort-merge join, and nested-loop join, with Figure 6a representing the case of a uniform S and Figure 6b representing the case of a distorted S.

[0075] Referring to Figures 6a and 6b, the hash join algorithm was the fastest in all experiments, efficiently utilizing WRAM and minimizing data transfer between WRAM and memory banks, achieving the best performance. In contrast, other join algorithms experienced difficulties due to frequent memory bank access. Using TAB in terms of the hash function can result in faster join processing speeds than using OAT or CRC32. Although Tabulation achieves a higher collision rate than OAT and CRC32, the higher hash throughput can mitigate the penalty caused by the higher collision rate.

[0076] To fully utilize PIM-supported DIMMs, a single IDP hash join algorithm must be extended to the entire rank, which consists of a minimum of 64 IDPs. Extending the rank level inevitably leads to partitioning issues. A single IDP join algorithm assumes that the entire R and S partitions are stored in memory banks. However, to extend to the entire rank, the tables must first be partitioned so that each IDP processes only one pair of R and S partitions, and there are no dependencies between partitions. This partitioning step, also known as global partitioning, distributes tuples from both tables to the IDPs so that all tuples belonging to the same hash key range belong to the same IDP.

[0077] Due to the non-shared architecture of IDPs, host CPU intervention is unavoidable in inter-IDP communication. This characteristic generates a huge amount of overhead, becoming a performance bottleneck. Therefore, the main challenge for fast inter-IDP communication is the data layout of the DIMM. Due to byte-based interleaving between multiple chips in the rank, the host CPU cannot interpret the source data obtained from a PIM-supported DIMM. To solve this, existing kernel modules for DIMMs can perform byte-level transposition for all cache lines received by the CPU, but this places a heavy computational burden on the CPU, resulting in a performance bottleneck.

[0078] The PID coupling method proposed in this invention can perform RnS (Rotate-and-Stream) operations to solve the high latency and low bandwidth of existing inter-IDP communication. By slightly adjusting the partitioning operation, the CPU load can be significantly reduced to a single rotation. Furthermore, the rotation occurs within a single cache line. As a result, data for inter-IDP communication can be stored only in registers without transferring it to main memory, eliminating the need for main memory access and enabling vector streaming.

[0079] Figure 7 illustrates the execution process of the RnS (Rotate-and-Stream) operation.

[0080] Referring to Figure 7, RnS is shown with an example of eight IDPs, each sending and receiving eight 8-byte tuples. A single tuple is sent and received between the eight IDPs. In the partitioning step, each IDP constitutes partitioned storage in a twisted order. Unlike existing IDP-to-IDP communication, where partitioning bins are always aligned by target chip ID starting from 0, RnS first places its own target bin and then places them in ascending order. For example, the partition for Chip 0 is stored in the order 0, 1, 2, ..., 7, and the partition for Chip 7 is stored in the order 7, 0, 1, ..., 6. Next, a single 64-byte cache line-sized data block is read from the main processor without byte-level transposition. The cache line contains the first 8 bytes of each of the eight chips, but is drawn in a byte-interleaving scheme with tuples of the same hue that are vertically aligned. After performing a bit rotation on each of the 64-byte blocks (the Nth chip has N bytes), inter-IDP communication becomes possible simply by transferring the data to a PIM-supported DIMM.

[0081] Because the rotation occurs within a 64-byte cache line, it can fit into a 512-bit SIMD register. This allows for efficient rotation using vector commands, and more importantly, it does not require main memory. Furthermore, cache access overhead can be minimized by streaming to the cache with non-temporal streaming operations. Here, the RnS implementation loads 64 bytes of data into a PIM-supported DIMM using the non-temporal stream load instruction _mm512_stream_load_si512(), which does not pollute the host CPU cache. The rotation is performed in place using the instruction _mm512_rol_epi64(). Then, once the data is written back to the DIMM, a stream storage operation _mm512_stream_si512() is used to access the DIMM directly without writing the data to the cache.

[0082] Figure 8 shows the bandwidth of RnS mutual IDP communication compared to the existing method.

[0083] Figure 8(a) shows the bandwidth when the total amount of data transferred between 64 IDPs in a single rank is changed from 32 MB to 2 GB. In terms of bandwidth, RnS is geometrically 6.23 times better than existing methods because it reduces the computational load on the host CPU and prevents data from flowing into main memory. Figure 8(b) shows how performance degrades when the number of available CPU cores decreases. Because RnS is less dependent on the CPU, it is less sensitive to available computational resources than existing inter-IDP communication. To perform RnS, each IDP must prepare the data to be transferred in packets. To utilize PIM-supported DIMMs, single-rank inter-IDP communication must be extended to multiple ranks and DIMMs. The PID-JOIN scheme proposed in this invention can adopt a parallel RnS access scheme.

[0084] The PID-JOIN method proposed in this invention does not require a strict order for the input and output tuples in the DIMM's memory banks, as global partitioning ensures that the correct IDPs are used. In the PID-JOIN method, R and S tuples are distributed evenly across all IDPs. The PID-JOIN method implements the USG (Unordered Scatter-Gather) operation, which is a mechanism for distributing R and S tuples evenly to IDPs without requiring arbitrary access to main memory.

[0085] Figure 9 illustrates the execution process of the USG (Unordered Scatter-Gather) operation, particularly showing the scattering process.

[0086] Referring to Figure 9, the USG's working model loads a 64-byte tuple directly into a 512-bit vector register in a streaming manner, removing unnecessary small loads across various areas of main memory. Then, after a simple byte-level transposition, it is transferred to a PIM-supported DIMM, completing the transfer. While existing IDP-based data transfer methods are limited by the requirement to transfer the tuple to the exact target IDP, the USG does not involve region partitioning. As a result, the USG accesses all main memory and PIM-supported DIMMs sequentially, while the tuple can be distributed evenly across the IDPs.

[0087] Figure 10 shows a working model of PID-JOIN according to the present invention.

[0088] Referring to Figure 10, the PID-JOIN method according to the present invention can perform USG and RnS operations between a CPU (Central processing unit) and a PIM (Processing-in-memory) supported DIMM (Dual in-line memory module), and perform join operations on the R and S values ​​of both tables.

[0089] First, the CPU performs a USG operation on the R and S tuples in main memory and distributes them evenly to the PIM-supported DIMMs ((1)). Next, the IDP (In-DIMM Processor) of each PIM-supported DIMM performs global partitioning on its own R and S tuples, grouping the tuples in each partition and forming RnS packets ((2)). All-to-all mutual IDP communication is performed between PIM-supported DIMMs using RnS, so that R and S tuples with the same hash key value are transferred to the same IDP ((3)). Subsequently, the IDP executes a single IDP hash join algorithm on the R and S tuples, and the result of the join is a tuple ((4)). As a result, the tuple is transferred to main memory via the execution of a USG operation ((5)).

[0090] The PID-JOIN method proposed in this invention is evaluated by comparing the execution delay time of existing CPU-based joining algorithms in a CPU-based basic system with that of a PIM-supported system. For workloads, a set of workloads derived from various prior studies on accelerating in-memory joining is used. Join execution delay times are measured for R and S of both tables with varying tuple counts and Zipf coefficients. Each tuple consists of a 32-bit integer join key and a 32-bit tuple ID. In addition, the Zipf factor of S is set to representative Zipf factors used in joining performance evaluation, namely 0.00, 0.25, 0.50, and 0.75, to evaluate the effect of distorted tables.

[0091] First, the execution delay time of the PID-JOIN scheme according to the present invention is evaluated by comparing it with CPU-based join algorithms using uniform R and S values. For the CPU-based join algorithms, SQLite is selected for nested-loop joins, M-PASS for sort-merge joins, and PRO and PRHO for hash joins. The PID-JOIN scheme includes a single IDP hash join algorithm, as well as single IDP sort-merge and nested-loop join algorithms. RnS and USG are applied to all three single IDP join algorithms.

[0092] Figure 11 shows experimental results of the execution delay time of CPU-based joining algorithms and PID-JOIN methods using a uniform R, S, and a variable |R|:|S| ratio.

[0093] Referring to Figure 11, the experimental results show that the PID-JOIN method according to the present invention efficiently utilizes the high internal memory bandwidth and computation throughput of the IDP, outperforming CPU-based joining algorithms for all workloads. In contrast, CPU-based joining algorithms experience limited memory bandwidth issues during the data transfer process between the host CPU and main memory. For the PID-JOIN method applying the fastest single-IDP hash join algorithm among the evaluated joining algorithms, a geometric mean speed improvement of 1.92 times is achieved compared to PRHO. As seen in Figure 11(a), the PID-JOIN method according to the present invention achieves a geometric mean speed improvement of 2.27 times compared to PRHO, with a speed improvement of 1.6 times for 64M tuples. The large speed improvement of the PID-JOIN method is mainly due to the fast hash table build and probing using PIM-supported DIMMs.

[0094] Figures 12 and 13 show the results of the analysis of the execution delay time of the PID-JOIN method.

[0095] Referring to Figure 12, when only a single IDP hash join algorithm is applied, a 1.62x decrease in geometric mean speed is observed due to the high overhead of inter-IDP and host-DIMM data transfer, which accounts for 92.4% of the average total latency. Inter-IDP communication and host-DIMM communication are only 1.79 GB / s and 1.77 GB / s, respectively, resulting in considerable latency. Applying RnS optimizes overall inter-IDP communication, improving the geometric mean speed by 1.15x compared to PRHO. Applying rank-based USG optimizes host-DIMM data transfer, improving the geometric mean speed by 2.27x compared to PRHO. As the |R|:|S| ratio increases, the latency of PRHO and PID-JOIN methods decreases, thus reducing R. Nevertheless, the speed improvement remains relatively constant. When the ratio of |R|:|S| is 1:1, the increase in geometric mean velocity is 2.43 times, and when it is 1:32, the increase in geometric mean velocity is 2.21 times.

[0096] By adjusting the Zipf coefficient of S from 0 to 0.75, the robustness of the PID-JOIN can be evaluated using a distorted table. Since the PID-JOIN method performs global partitioning by hashing, high skewness can lead to load imbalances between IDPs.

[0097] Referring to Figure 13, we see that PID-JOIN is robust to biased tables except for a Zipf coefficient of 0.75. However, when the Zipf coefficient is 0.75, the execution latency of PID-JOIN remains lower than that of PRHO. In contrast, the execution latency of PRHO decreases slightly as the Zipf coefficient increases. This is because the number of duplicate tuples in S increases. Duplicate tuples generate the same hash value and refer to the same item in the hash table. In PID-JOIN, duplicate tuples can induce imbalances between IDPs because global partitioning assigns duplicate tuples to the same IDP. It is expected that such load imbalances between IDPs can be mitigated by using a skew processing method that modifies duplicate tuples to generate other hash values.

[0098] Next, we compare the system cost of the PID-JOIN method according to the present invention with the system cost of a CPU-based PRHO.

[0099] For comparison, we can first collect and sum the manufacturer's suggested retail price (MSRP) for PIM-supported and CPU-based base systems. Then, by dividing the system cost by the maximum coupling throughput achieved in the experiment (i.e., tuples / s), we can obtain the system cost per unit of performance for the two coupling algorithms.

[0100] The comparison showed that PID-JOIN achieved a 26.9% lower system cost per performance than PRHO. PID-JOIN achieved a maximum of 158.4M tuples / second on a $4,020 PIM-supported system, while PRHO on a base system achieved a maximum of 65.2M tuples / second. The results suggest that PID-JOIN may be a more attractive choice for RDBMS than CPU-based join algorithms.

[0101] Sensitivity Studies

[0102] The effect of bucket size on RnS (Rotate-and-Stream)

[0103] To study the impact of RnS packet size on all-to-all inter-IDP communication bandwidth, we measure total bandwidth, effective bandwidth, and fill rate for five packet sizes (8, 16, 32, 64, and 128 bytes) with various S sizes (64M, 128M, 256M, and 512M tuples). Total bandwidth is measured including null packets, while effective bandwidth excludes null packets. The fill rate is the ratio between the two bandwidth values.

[0104] Figure 14 shows how the RnS bandwidth changes depending on the packet size. As mentioned above, RnS uses 8-byte packets (64 bytes in the case of a rank containing 8 IDPs) to achieve fast inter-IDP communication. Ideally, the total bandwidth should be kept constant with packet sizes larger than 8 bytes. Larger packet sizes benefit from sequential memory bank access. However, larger packet sizes result in a slight decrease in effective bandwidth due to the decrease in fill rate.

[0105] PIM kernel call overhead

[0106] This study analyzes the impact of PIM kernel call overhead on PID-JOIN throughput. Every PIM kernel call incurs approximately 12ms of latency overhead in a PIM-supported system. Because PID-JOIN calls different PIM kernels for partitioning, hash table building, and hash table lookup, the overall PIM kernel call overhead remains constant. To ensure that the overhead does not offset the performance benefits of PID-JOIN, the PIM kernels must have a sufficient amount of work that increases proportionally with the table size.

[0107] Figure 15 shows the impact of PIM kernel call overhead on R and S of various sizes. The contribution of call overhead to join execution latency decreases as table size increases. Call overhead can be ignored for 256M tuples. However, call overhead accounts for 14.9% for 64M tuples. Despite the overhead, PID-JOIN still consistently outperforms PRHO for all workloads evaluated.

[0108] Collision handling method

[0109] The single IDP hash join algorithm for PID-JOIN uses linear probes (LPs) for collision handling. However, LPs can introduce significant performance overhead as the number of failed probes increases.

[0110] Figure 16 shows the execution delay analysis and collision ratios for PID-JOIN, LP, QP (quadratic probing), and DH (double hashing).

[0111] Overall, LP has a higher collision rate than QP and DH, but the lowest execution latency. LP's low latency is due to its low computational overhead. When a collision occurs, LP simply probes the next item in the hash table by adding one to the previous key. However, QP requires integer multiplication, which IDP doesn't fundamentally support, and DH involves an additional hash function. The overhead of these additional calculations increases the combined execution latency for QP and DH, making LP the fastest collision handling method. A similar trend can be observed even with distorted S.

[0112] The results also show that local partitioning accounts for up to 48.0% of the PID-JOIN join execution delay. Local partitioning limits the maximum number of tuples per partition, calculated assuming that partition tuples fill a 32KB hash table with a 50% packing ratio, and ensures a working set of partitions that fits in WRAM. In cases of extreme skewness, excessive duplicate tuples may cause local partitioning to not fit the partitioning working set in WRAM.

[0113] Potential speed improvements via TPC-H queries

[0114] We will now use four single-join TPC-H queries (Q4, Q12, Q14, Q19) with a scale factor of 400 to probe the potential speed improvement of PID-JOINs when integrated into a full RDBMS. For the experiment, we will first measure the inter-terminal query execution latency in MonetDB, a well-known in-memory RDBMS, on a CPU-based baseline system. Next, we will identify the execution latency of the join operator and calculate the expected speed improvement for the query by replacing the join execution latency with a PID-JOIN. For each query, we will first extract the tables used as input for the join operator according to the query plan in MonetDB. Then, we will perform a PID-JOIN on the extracted tables from a PIM-supported system. When measuring the join execution latency of the PID-JOIN for the query, we will consider the overhead of host-DIMM data transfer.

[0115] Figures 17a and 17b show an analysis of query execution latency in MonetDB and improvements in expected query speed when integrating PID-JOIN into MonetDB.

[0116] PID-JOIN achieves a geometric mean speed improvement of 2.55 times in query joins. The expected speed improvement of PID-JOIN augmented MonetDB is 1.14 times geometric mean. The results indicate that using PID-JOIN to accelerate joins is a promising option for in-memory RDBMS. The PID-JOIN method proposed in this invention can take various forms of PIM hardware in addition to PIM-supported DIMMs.

[0117] The PID-JOIN method proposed in this invention can maximize join throughput in the IDP by using distributed and partitioned hash join algorithms, and can utilize the high internal memory bandwidth and computation throughput of the DIMM by performing global partitioning of both tables in the DIMM instead of the CPU. At the same time, the overhead of inter-IDP communication can be minimized by using CPU cache streaming and vector instructions during the many-to-many inter-IDP communication process with respect to the CPU. Furthermore, the PID-JOIN method proposed in this invention can minimize data transfer latency between the CPU and the DIMM, thereby accelerating data transfer.

[0118] While the above description has been made with reference to preferred embodiments of the present invention, those skilled in the art will understand that the present invention can be modified and altered in various ways, without departing from the spirit and scope of the invention as described in the following claims. [Explanation of symbols]

[0119] 100 PIM-based computing devices 110 Main Memory 130 PIM modules 131 memory processors 133 Working memory 135 Memory Banks 150 processors

Claims

1. Main memory for storing the first and second datasets, A PIM (Processing-In-Memory) module comprising multiple memory processors, a working memory connected to the memory processors, and an inline memory module including a memory bank, which performs join operations on the first and second data sets transferred to the working memory, The system comprises a processor that forms a first channel and a second channel between the main memory and the PIM module, and takes in the first and second data sets in the first channel and stores them in the memory bank in the second channel, The aforementioned PIM module is The aforementioned inline memory module is implemented using multiple dual inline memory modules to support scale-out join operations. The system is configured to perform global partitioning through interdata relocation between the multiple dual inline memory modules, thereby rearranging the first and second datasets. The PIM-based computing device is characterized in that, in the process of rearranging the first and second datasets, the memory processor generates a hash table in the working memory, moves the first and second datasets in the memory bank to the working memory and performs local partitioning, associates them with the hash table using a hash function, and performs a hash join to rearrange the first and second datasets in the hash table to the memory processor so that all first and second datasets belonging to the same hash key range belong to the same memory processor.

2. The PIM-based computing device according to claim 1, characterized in that the PIM module subdivides the data of each partition and connects it to one of the second channels by global partitioning for the plurality of dual inline memory modules.

3. The PIM module, in the process of rearranging the data, performs RnS (Rotate-and-Stream) operations by bit rotation and streaming to transfer sequential data to the data in each memory bank of the plurality of dual inline memory modules, performs many-to-many communication between the memory processors so that the first and second datasets having the same hash key value are transferred to the same memory processor, then the memory processor executes a hash join algorithm to join the first and second datasets, and transfers the joined first and second datasets to the main memory, characterized in that the PIM-based computing device according to claim 1.

4. The PIM-based computing device according to claim 3, characterized in that the PIM module transmits and receives a specific dataset in units of datasets between the plurality of dual inline memory modules and performs the RnS (Rotate-and-Stream) operation.

5. The PIM-based computing device according to claim 1, characterized in that, before performing the global partitioning, the processor performs a stream that sequentially transfers the first and second datasets in the main memory from the main memory, then performs an Unordered Scatter-Gatter (USG) operation to evenly distribute the first and second datasets to the memory banks, and stores them in the memory banks.

6. The PIM-based computing device according to claim 5, characterized in that the processor, in the process of the USG calculation, directly loads the first and second datasets into vector registers in a streaming manner, performs byte-level transposition, and then transfers them to the memory bank in an STS (Stream Transfer Send) procedure, thereby distributing the first and second datasets to the memory bank.

7. Main memory for storing the first and second source datasets, Multiple PIM (Processing-In-Memory) modules, each comprising an inline memory module including a memory processor, a working memory connected to the memory processor, and a memory bank, and each PIM module including multiple memory processors that perform join operations on the first and second processed data sets transferred to the working memory, A processor that preprocesses the first and second source datasets and stores them in the memory bank, The aforementioned multiple PIM modules are, The system is configured to perform global partitioning via inter-data relocation on the pre-processed first and second source datasets, thereby rearranging the first and second processed datasets. PIM-based computing device characterized in that, in the process of rearranging the first and second processed datasets, the memory processor generates a hash table in the working memory, moves the first and second source datasets in the memory bank to the working memory and performs local partitioning, associates them with the hash table using a hash function, and performs a hash join to rearrange the first and second processed datasets of the hash table to the memory processor so that all first and second source datasets belonging to the same hash key range belong to the same memory processor.

8. The PIM-based computing device according to claim 7, characterized in that the processor performs a stream of sequential transfers from the main memory to the first and second source datasets in the main memory, then performs an Unordered Scatter-Gatter (USG) operation to evenly distribute the first and second source datasets to the memory banks, and stores them in the memory banks.

9. The PIM-based computing device according to claim 7, characterized in that the plurality of PIM modules perform a bit rotation and streaming RnS (Rotate-and-Stream) operation on the pre-processed first and second source datasets in the process of rearranging the data, perform many-to-many communication between the memory processors so that the first and second source datasets having the same hash key value are transferred to the same memory processor, then the memory processor executes a hash join algorithm to join the first and second source datasets, and transfer the joined first and second processed datasets to the main memory.

Images