右上角可以切换中文。

This is a Python-free site.

I am currently not recruiting international students.

Domestic students should refer to the Chinese version of this announcement.

(Auto-Translation via LLM)

The Taalas HC1 has recently attracted widespread attention. This article attempts to decode its architecture from a technical perspective to serve as a reference for industry peers. The analysis is based on published patents and interviews from Taalas. Because its architecture has not been made public, many details require reverse engineering and speculation, so accuracy cannot be guaranteed.

The HC1 features two major architectural advancements:

1. It upgrades the naive matrix multiplier to a mechanism specifically tailored for quantized operations, eliminating redundant calculations. The computing mechanism can be summarized as “Multiply-Select-Add” (rather than the standard “Multiply-Add”).
2. It uses ROM as the storage medium for model weights, achieving very high density. Modifying the model weights only requires altering two upper-layer via masks.

The new computing mechanism improves efficiency by leveraging quantization effects, rather than the constant nature brought about by fixed weights. Therefore, a very clear boundary still exists between memory and computation in the HC1 architecture. This means the act of memory access during HC1’s operation is still present; only the accessed target has changed. In terms of density, this solution is already excellent; however, the current computing mechanism has not yet fully utilized the constancy of the weights, leaving further room for exploration in terms of energy efficiency. Its weights need to be read from the ROM row by row, which limits peak performance capabilities.

In the context of linear algebra, every real-valued quantity can take continuous values; but in computers, numerical precision is always limited. Especially after quantizing and compressing a model, quantization effects become far more prominent. For example, if every quantity in a weight matrix is quantized to a 2-bit representation, regardless of how the number system is designed, each number can only take a maximum of 4 different values. Under a binary integer system, these 4 values are 0, 1, 2, and 3. These four specific values can be arbitrarily designed to form number systems other than binary to help quantize more accurately, but there are at most four, which is strictly limited by the information capacity of two bits. Therefore, without loss of generality, this article assumes the values are , , , and . Later, we will assume the weight precision is 2 bits to simplify the illustrations; readers can easily generalize this to higher precisions.

When an input neuron activation value is to be multiplied with such a weight matrix, because the number of operations exceeds the number of value types, the operations inevitably generate a massive amount of repetition. For example, assuming the first column of the weight matrix is , a certain input neuron activation value will be multiplied with these weights respectively, and the products will participate in different accumulation processes. During this process, it can be observed that is calculated 7 times, is calculated 2 times, and is calculated 5 times.

To address how to eliminate this repetition to save hardware overhead, academia has proposed various solutions. The strategy adopted by Taalas is the LUT (Look-Up Table) scheme, which eliminates repeated multiplications through pre-computation and selection. Since the weight value space is extremely small after quantization, the products of the input activation value and all possible weight values can be entirely pre-computed once, listing all potentially used products. Every subsequent individual multiplication operation can then be replaced by selecting a ready-made value from the pre-computed results. Reviewing the 14 patents published by Taalas, its core mechanism is primarily documented in Large parameter set computation accelerator… (US20250123802A1).

1. Multiply: First, pre-computation is performed, multiplying with the constants , , , and respectively. Every input neuron activation value will be pre-multiplied with all possible weight values. To further compress the area, the operations do not need to be completed in parallel. can be input bit-by-bit from LSB to MSB, and the multiplication with fixed values consists of delays (DFF) and carry-save adders (CSA). All signal lines can be simplified to a single bit. The output signal lines are extended to form a parallel bus structure for selection in subsequent steps.

2. Select: Thus, the operation of with the vector becomes a selection. 16 horizontal signal lines are set up to carry the operation results, forming a crossbar switch with the pre-computation output signal lines. Each horizontal line needs to form a connection at the intersection corresponding to the vector value to extract the required pre-computed signal. Since 16 items are too many, the latter part is omitted in the illustration; its structure is repetitive.

3. Add: At this point, the structure has completed these 16 operations without repetition. These operations belong to 16 different output neuron activation values, thus requiring 16 different accumulators to accumulate them.

In this way, the structure completes the multiplication of an input scalar and a weight vector. The next step needs to calculate a similar operation process for and another vector, and accumulate the results together. Since the connectivity of the intersections depends on the value of the weight vector, and this value changes between the two operation steps, the intersections cannot be “hardwired” here. They are changed to reconfigurable crossbar switches. Note that Bajic’s media interview (https://www.nextplatform.com/compute/2026/02/19/taalas-etches-ai-models-onto-transistors-to-rocket-boost-inference/4092140) mentions “we can store four bits away and do the multiply related to it – everything – with a single transistor”; this statement can easily cause ambiguity literally. Combined with the patent scheme, a more reasonable explanation from an engineering implementation perspective is that “store four bits away” has nothing to do with the half-sentence after “and”. He believes that all operations (including multiplication) related to each weight are completed by a single transistor. To achieve the effect he described, we implement the crossbar switch as controlled by a single transistor for switching on and off. Although there are four transistors in each row in this example, indeed only one is conducting at a time and thus participating in the operation.

Usually, using PTL to construct transmission gates requires a pair of transistors, but using a single transistor here is reasonable. A single transistor will only cause a threshold voltage drop; adding a Buf outside the crossbar switch array and before the accumulator to restore the signal is sufficient. Bajic claimed the company hand-drew a portion of the layout, which is very likely for handling this specific part.

The gate of the transistor needs to be provided with a one-hot control vector, which can be read directly from the ROM. Thus, the complete matrix operation architecture is essentially formed. Surrounded by auxiliary functions like vocabulary lookup tables, random sampling, normalization, softmax, positional encoding, and activation functions, paired with a medium-capacity SRAM for storing KV (the model’s own context is very short), the entire model is fully implemented.

Converting each 2-bit weight to a 4-bit one-hot encoding and writing it to ROM requires a relatively large ROM capacity, which contradicts the “store four bits away” claim (HC1 claims to use 3/6-bit mixed quantization; if it were a one-hot representation, it should “store eight bits” at least). Therefore, I believe what is burned into the ROM should be the original weight values, which are then converted to one-hot via a 3-to-8 decoder. The decoder is a one-dimensional structure outside the matrix, and its total overhead is not significant.

Taalas claims that customizing the Mask ROM requires changing two layers of photomasks. This is highly likely for two layers of vias, whose structure should be as illustrated: VDD and GND controlled by wordlines complementarily fill the upper and lower layers of Trenches. The middle Trench serves as the bitline for reading out, and vias are placed above or below the bitline in a complementary manner to represent 0 or 1.

Therefore, the truly hardwired part in Taalas’s scheme is only the vias of the ROM. The entire architecture can still be clearly divided into a storage part and a computing part; memory and computation are structurally separated. It does not eliminate memory access in a broad sense, but merely changes the medium of the accessed target from SRAM (Groq) or HBM (Nvidia) to ROM, and the accessed data is also highly likely to be unprocessed original weight values. If accessing memory on the same chip is not considered memory access, then the goal of eliminating this narrow “off-chip memory access” has actually already been achieved by the Groq architecture.

The computing part is a classic quantized computing architecture, very similar to UNPU, Cambricon-P, etc. Utilizing these architectures to improve the efficiency of quantized matrix multiplication does not require assuming the weights are fixed. I have provided an architecture diagram comparison for reference.

Under this architecture, because structures like pre-multiplication, accumulation, decoders, and Bufs are all located at the periphery of the array and are one-dimensional structures, they can all be regarded as preparatory or finishing work for the actual computation. Only the two-dimensional crossbar switch is regarded as the true computation. From this perspective, the statement “a single transistor completes all operations” also makes some sense, though slightly exaggerated. From a stricter perspective, it only undertakes the selection/gating function in the operation.

The computing part can also be directly used to process variable-value matrices, as long as the source of the crossbar switch array control signal is switched from ROM to SRAM. I suspect this is how HC1 supports LoRA. This architecture can achieve an area balance between computing devices (mainly limited by transistor components) and memory devices (mainly limited by high-level metal vias) by adjusting size parameters. However, the computation needs to be completed by reading the ROM row by row, so processing a complete matrix requires traversing all addresses of the entire ROM, which limits the peak computing power. If multiple matrices are written into the same ROM, it becomes very difficult for this architecture to implement pipeline parallelism.

In terms of density, since computation and ROM storage are separated, the ratio can be adjusted freely, so the area is mainly limited by the ROM. Assuming the Half Pitch of the vias used for ROM encoding is 62 nm, and the ROM bit-cell density is 0.015 μm², an 8B model with an average of 4.0 bpw would require 480 mm², which basically matches the implementation. If the layers are reduced to 38 nm, the density becomes 0.0058 μm², and the model would only need about 180 mm², but it cannot be completed in a single exposure, so the model update cost would double. Since hardwired chips usually do not require ultra-large-scale mass production, if it can be coordinated with the manufacturing process to introduce e-beam lithography in the production line to achieve maskless via processing, it will help realize high-density variable vias at a low cost. Imagine that changing models for chips no longer requires redesigning; merely introducing one or two layers of e-beam drilling processes costing a few thousand yuan per wafer could continuously output chips equipped with various models. The final cost barrier for the commercialization of hardwired language processors would be resolved.

Strictly from an architectural design perspective, the HC1 still has many areas open to debate, but it genuinely realizes a vision that most people cannot yet accept. It is itself just a Proof of Concept, and evaluation of it should not focus too much on its currently demonstrated AI capabilities. The mindset of this team shifted early on, and they have proven to have found an effective technical path. There are not many obstacles left on the subsequent path toward realizing disruptive products. Hardwired chips for the latest cutting-edge models will soon be realized, and by then, it will have a profound impact on the existing AI computing paradigm and chip industry landscape.

Translated by LLMs.

Computer science is one of the fastest-developing disciplines in human history. Yet, trailing behind its leapfrogging technological advancements, its underlying mental models can sometimes appear outdated. In our daily research, we are accustomed to empirical and quantitative expressions; this article, however, attempts to untangle the evolutionary logic of computing architectures through the lens of the philosophy of science. The critical discussion herein does not aim to diminish the achievements of our predecessors under the old paradigm, but rather hopes to provide readers with a rational new dimension for understanding the future of general-purpose computing through the collision of fundamental concepts.

I

Humanity’s exploration of science can never be detached from the objective historical conditions of its time. It is precisely these conditions that determine our initial perspective for understanding the world and profoundly shape the underlying epistemology of different disciplines.

Physics is a discipline concerning existing reality. Before humans had the ability to smash atoms, the sun, moon, and stars had already been in motion for billions of years. Therefore, humanity’s cognitive path in physics inevitably progressed from the macroscopic to the microscopic. Early physical models began with Kepler’s laws of planetary motion, which were later explained by Newtonian mechanics, while the understanding of microscopic elementary particles had to wait hundreds of years until the modern era. Historical conditions dictated that humanity could only gradually build microscopic models after macroscopic models had been perfected.

Consequently, in physics, any epistemologically greedy and radical reductionist thought would directly challenge the prevailing consensus of the time, thereby easily subjecting itself to thorough critique. Such thought suggests: since universal gravitation has been discovered, Kepler’s laws are no longer needed; since the relativistic view of space-time is established, Newton’s absolute space-time is a fallacy; since all interactions can be reduced to a few fundamental forces, and all matter to elementary particles, the exploration that began with a falling apple is merely a superfluous transition towards the Standard Model. It assumes that because more fundamental laws have emerged, the value of macroscopic scientific laws at higher cognitive levels can be negated.

However, in computer science, humanity’s cognitive path progresses from the microscopic to the macroscopic. Computer science is a discipline concerning artifacts. Constrained by the engineering capabilities of their historical periods, computer scientists lacked the historical conditions of physicists to preemptively build defense lines of macroscopic understanding, making them prone to preconceived education in reductionist viewpoints. This has resulted in the epistemological reductionism—which has been thoroughly critiqued in physics—being long adopted as the mainstream in computer science.

Alan Turing was a great mathematician and a pioneer of computing and artificial intelligence. As a branch of mathematical logic, his theory established the ontology of computation, but simultaneously harbored massive epistemological limitations that were never critiqued. His theory holds that basic arithmetic and logic can compute all computable problems—translated into a physics context, this statement is akin to saying “the outcome of World War II can be explained by the motion of quarks and electrons.” On the one hand, this view is logically rigorous, demonstrating the universality and elegance of the theory; on the other hand, it negates the significance of other cognitive levels and subtly redirects their value back to itself.

Turing believed that a general-purpose logical computer could achieve human intelligence through programming; Chomsky believed that the human mind could be reduced to a set of explicit “transformational-generative grammar” rules; Feigenbaum believed that expert systems could be constructed by stacking formal logic and knowledge bases. Driven by this reductionist conviction, generation after generation of scientists flew like moths to a flame into an endlessly unresolvable problem: spending their brief lives writing finite rules in an attempt to fit immensely complex intelligent functions. However, the Kolmogorov complexity of a closed computing system will not spontaneously increase, meaning that finite code and logical rules can never conjure up out of thin air an amount of information and complexity that surpasses their own structure. Infinite theoretical possibilities ultimately still require realization through finite human effort. Therefore, a physicist would never attempt to recreate the Battle of Stalingrad by manipulating quarks; yet, the ontological success of computer scientists long obscured their epistemological primitiveness, rendering their failure inevitable.

Ultimately, these convictions shattered alongside the final efforts of Japan’s Fifth Generation Computer Project, but they also left behind some world-changing byproducts: Turing’s theory is the cornerstone of today’s information society, and the general-purpose digital logic computer is a tool humanity can no longer live without; Chomsky’s theory, though unable to explain the human mind, forged the branch of “Formal Languages and Automata,” becoming the foundation of modern compilers; Feigenbaum won the Turing Award in 1994 for his contributions to early artificial intelligence systems.

II

In the field of computer architecture, computational reductionism long drove a reactionary force that suppressed the birth of new architectures. Under this mindset, the CPU—designed specifically for basic mathematical operations and logical jumps—was the best practice for realizing a general-purpose computer. Since the CPU already possessed complete functionality—guaranteed by Turing—developing other architectures simultaneously was deemed meaningless. Up until the early 21st century, the CPU was almost the sole object of study in this discipline.

The suppression of emerging architectures by this mindset was devastating. Because all computation was ideologically reduced a priori, the value of any architecture distinct from the CPU was preemptively negated, allowed to exist at best only as a “slave” device. You could design a specialized video encoder and attach it to the periphery of the CPU as a second-class citizen to supplement the CPU’s performance. Under quantitative research methods, its value was reflected in the score improvement of a single test: 464.h264ref. Because it only affected one out of an infinite number of programs, it would be deemed inconsequential; because programs evolve, it would be viewed as a temporary workaround.

The CPU’s unique universality was considered the only foolproof strategy for enduring adaptation to program changes. New architectures might conciliatorily pivot to prove that they too possessed Turing completeness, hoping thereby to receive equally serious treatment. However, they would fall into a concept specifically designed to trial such conciliatory thought: the “Turing tar-pit”—a term coined by the first Turing Award winner, Alan Perlis, referring to a universality that is only theoretically possible but practically meaningless.

Under the reductionist viewpoint, deep learning was also long alienated as just a mundane instance among a vast array of programs, viewed as 052.alvinn—essentially no different from 464.h264ref. Therefore, developing a processor chip specifically for program 052 was considered short-sighted, as people believed it was impossible to predict whether tomorrow’s intelligent algorithms would evolve from deep neural networks into support vector machines, rendering the chip obsolete before it was even born. In the face of rapidly rising deep learning, reductionism’s adherence to universality led the entire industry into a paradox: the stronger the momentum of deep learning, the more meaningless it was to research deep learning processors. This stalemate lasted until 2011, culminating in the bizarre spectacle of Google using tens of thousands of CPUs to train a model to recognize cats.

However, the evolution of history has proven that deep learning processors ultimately transcended their subordinate status, developing into a profoundly impactful industry. This is because deep learning is not a simple stacking of existing computational programs, but rather represents another level of human understanding of computation. It possesses an equally profound theoretical foundation and constitutes an entirely different general-purpose computing paradigm:

• Type I Universality (Mathematical Logic): Theoretically based on the Universal Turing Machine (the Hilbert Entscheidungsproblem), it relies on discrete states, discrete computation, and symbolic representation, aiming to decide all logically definable propositions in mathematics recursively. It achieves universality by passively accepting programming and excels at handling various heavy, repetitive logic that can be clearly described by rules.
• Type II Universality (Deep Learning): Theoretically based on the Universal Approximation Theorem (Hilbert’s Thirteenth Problem), it relies on real-number states, real-number computation, and data distribution probabilities, possessing the ability to approximate any continuous function defined on a Euclidean space. It achieves universality through backpropagation training and can solve intelligent tasks whose rules humans cannot explicitly describe but can learn from data.

Recently, a rather authoritative narrative describing deep learning processors as “Domain-Specific Architectures” (DSAs) has become widely circulated. In my view, the concept of “Domain-Specific Architecture” is a product of compromise between the alienation under old concepts and the reality brought about by new ones. On the one hand, it acknowledges the past mistake of alienating deep learning as “a mundane program,” making room for the development of deep learning processors; on the other hand, it is only willing to correct it to “a mundane domain” so as to maintain the supremacy of the CPU (the Type I general-purpose architecture). Deep learning processors have been around for over a decade; has the new golden age of diverse “Domain-Specific Architectures” arrived as promised? Which viewpoint better explains reality is something every reader can objectively judge for themselves.

Google’s task of training a model to recognize cats is almost impossible for a single CPU to complete in any reasonable timeframe, yet it is a breeze for a deep learning processor. However, metaphysical computational reductionists insist that only the former is general-purpose—they believe the CPU can even support artificial intelligence a century from now. The absurdity in this view is akin to believing a physicist can diagnose and treat diseases. From the perspective of intelligent computing, the CPU is nothing more than a “Turing tar-pit.”

III

Some time ago, I was commissioned by Professor Chen Yunji to help summarize and review this history of scientific development. I was amazed to discover that history is a cycle—I once thought I was walking a lonely path, only to find the footprints of predecessors everywhere.

In my view, today’s deep learning Large Language Models (LLMs) have already initially achieved Artificial General Intelligence (AGI), though this has not yet formed a broad consensus. Therefore, I predict that, based on LLMs, a third paradigm of general-purpose computing is about to emerge: cognitive intelligence based on natural language. It achieves universality based on natural language and context comprehension, excelling at macroscopic planning, common sense, logical reasoning, and interaction. This is currently the only known computational pathway capable of achieving high-level cognitive functions, already performing very closely to human intelligence.

By researching highly efficient, low-cost hardwired solutions, a Language Processor that operates directly on natural language tokens can achieve another thousand-fold increase in efficiency compared to deep learning processors (including GPUs). A thousand-fold efficiency increase is no longer just a quantitative change; it unlocks a new level of computation, realizing a silicon-based brain with advanced cognitive capabilities capable of real-time response. This level of computational efficiency was previously widely believed to require a switch to analog circuits, optical computing, new materials, or other disparate technical routes, but our solution is implemented based on standard CMOS processes, possessing the potential for rapid practical application.

Technically, we have basically found a viable path and are continuously advancing it, but the real difficulty lies in reversing the entrenched concepts formed over a long period. Just as deep learning processors had to face a history of alienation, at a time when the industry’s attention is entirely focused on building massive “Stargates” with astronomical investments, this novel Language Processor must also first overcome misunderstandings. Due to a lack of broad awareness regarding the possible formation of a “Third Paradigm of General-Purpose Computing,” most system researchers habitually hold up the yardstick of old concepts.

• Misconception 1: The Language Processor implements only one model and is therefore inconsequential.

  Just as some once thought deep learning was merely program number one in SPEC, today some might think a single LLM is just another mundane model in MLPerf—like Llama 3.1 8B, for instance. However, language capability itself is an important form of general intelligence, and its ability and value in solving general problems need no further elaboration. Exactly which specific model is used to achieve this language capability has already become a secondary technical issue, akin to “choosing which vendor supplies the CPU.”

• Misconception 2: It is impossible to predict whether tomorrow’s model architecture will evolve from Transformer to Mamba, rendering the chip obsolete before it is born.

  In both eras, detractors claimed that current algorithms were in a period of violent turbulence, making it dangerous to bet on a specific algorithm. However, by summarizing the fundamental rules of operations, deep learning processors were able to achieve broad support for deep learning algorithms. They did not perish due to architectural shifts from AlexNet to ResNet to ViT, and running support vector machines on a deep learning processor was never a difficult task to begin with. The Language Processor does not simply haphazardly bake a model into hardware; instead, it extracts the crucial computational patterns for targeted design, providing hardwired templates for matrix-vector operations that are universal to the dominant feed-forward networks and projection components, supplemented by controllable auxiliary units to jointly complete the computation. A “Sea of Neurons” substrate is entirely capable of simultaneously supporting multiple Transformers or other neural network models, and supporting Mamba presents no technical difficulty.

• Misconception 3: Because models update rapidly, the Language Processor will quickly become obsolete.

  This is currently the hardest prejudice to break in the industry. Faced with the ever-changing development of LLMs, people instinctively retreat to an architecture with absolute flexibility to find a sense of security. In the past, this safe harbor was the CPU; today, this safe harbor has become the GPU, implying that “the stronger the momentum of LLM development, the more meaningless it is to research Language Processors.”

  To adapt to rapid development, GPUs have built a massive and heavy software stack (driver - runtime - programming language - operator library - framework - inference system). The astronomical cost of maintaining this entire software ecosystem is currently the GPU’s strongest monopoly barrier. In the era of LLMs, however, this entire expensive software ecosystem ultimately serves only the centralized deployment of a handful of models like GPT, Gemini, and DeepSeek, resulting in very low resource allocation efficiency.

  In a Language Processor, the hardware operates directly at the token level, eliminating the cost of the software ecosystem. Under current technical solutions, hardware tooling costs can also be well controlled. The cost of resetting a set of hardware in tandem with a model update can be made significantly lower than the cost of training the new model itself. “Achieving model updates by replacing hardware” has never historically appeared as a realistic solution, so it will take more time for it to gain widespread acceptance.

  On the other hand, assuming Moore’s Law accelerates and integrated circuit performance doubles every few months, a CPU’s process technology would no longer be cutting-edge before it even hits the market. Under such circumstances, would we stop manufacturing CPUs altogether? Please note that updating a Language Processor will not change its natural language interface; it will only improve its performance on certain tasks. The only thing capable of obsoleting a Language Processor through performance is another Language Processor.

IV

Turing spent his entire life debating his conviction that artificial machines would eventually achieve human intelligence. He tirelessly wrote articles refuting, one by one, arguments from theology, animism, “strawberries and cream,” telepathy, and other views that seem utterly absurd today. He optimistically wrote: “I believe that at the end of the century the use of words and general educated opinion will have altered so much that one will be able to speak of machines thinking without expecting to be contradicted.”

Turing himself once firmly believed that universality already implied that logical computers could attain human intelligence, but in his later years, he also realized the difficulty of simply reducing human intelligence to writing programs on a computer. He wrote something to the effect of: “A general-purpose logical computer is just an extremely disciplined form; to make a machine produce human intelligence, one should build an unorganized machine and then educate it through a system of rewards and punishments.” This new machine form bears a striking resemblance to today’s deep learning. However, because it provided no direct help to the quantitative metric races of CPUs studied in computer architecture, his viewpoint was only selectively absorbed—specifically, the first half—by today’s computer science education system. This shows that in a young discipline like computer science, the knowledge we were previously taught is not necessarily the entire truth. Innovation means constantly re-hammering every established consensus.

Although Turing did not directly realize artificial intelligence, without the glorious development of general-purpose logical computers as tools in the past, there would be no birth and renaissance of deep learning, let alone the prototype of Artificial General Intelligence we are fortunate enough to witness today. After nearly a century of low-level, microscopic stacking, computer science has finally, via LLMs, truly touched the macroscopic “cognitive” level. Following the advent of Language Processors, our creations no longer merely compute numbers, but begin to think deeply through the medium of language. This marks an important turning point in the evolution of the entire discipline from the “engineering” of computers toward “cognitive science.”

History has handed this important turning point to our generation of computer architects. Under existing concepts, the Language Processor architecture seems so incomprehensible; yet, when re-examined through a shifted conceptual lens, it is so intuitive. Hopefully, ten years from now, when people look back at this architecture, they will all feel that it is nothing but a matter of course.

(Automatically Translated via LLMs)

Our paper on the “Hardwired LPU” has been accepted by ASPLOS 2026, a top-tier conference in computer architecture. In this paper, we propose a novel processor morphology: the Hardwired LPU. By directly hardwiring Large Language Models (LLMs) into the chip’s circuit structure, we have achieved an energy efficiency improvement of over 1,000x compared to existing GPUs.

Original Paper: [DOI][PDF] (DOI link may not be active yet)

In the field of integrated circuits, hardwiring neural networks is not a new concept. Since the 1980s, this idea has been repeatedly proposed but has failed to make an industrial impact due to a lack of universality. Applying this historical experience to the future has formed a kind of inertial wisdom: “Look, those who pursued extreme specialized hardware all failed. Nvidia dominates precisely because GPUs sit at the perfect ‘Sweet Spot’ between generality and efficiency.”

However, in the real world, “generality” is an engineering economics proposition: it requires that the value created by a computer must cover its engineering costs. If this condition is not met, that type of computer simply will not exist. Furthermore, there is no immutable physical law stating that computer hardware must support software updates to function. Imagine if you had the sudden idea to build a machine that could only calculate a specific hash function. To a textbook, this would be heresy; but in reality, as long as it is profitable, what does it matter if it cannot be updated? Years later, open-minded and agile actors made fortunes building such machines (ASIC miners), while bystanders could only question them and miss the opportunity: “What if the protocol changes to PoS tomorrow? Are you going to build a new machine?”

Thus, the value of a processor lies not in how many programs it can support or how many model generations it can survive, but in whether it has found a task with sufficiently high value—even if it is just one task. Today, historical conditions have quietly changed. Humanity has obtained the first neural network truly worth hardwiring: ChatGPT. The application value created by LLMs needs no further elaboration, but we can quantify it intuitively: a dedicated cryptocurrency miner earns less than one cent per second; whereas, at current pricing, a single Hardwired LPU node could generate revenue of up to 20 CNY per second.

Even so, there will always be skeptics: “LLMs iterate so fast. What if the algorithm changes to Mamba tomorrow? Are you going to replace the hardware?”

Why not? In the past, computers emphasized software updates because the marginal cost of software iteration was essentially zero. But today, the capital investment required to train a competitive LLM has historically surpassed the cost of mask revisions for advanced process chips. Under these conditions, the model—being more expensive—is the enterprise’s core asset, while the chip, relative to the model, has become a “consumable.” If the application value of an LLM is sufficient to drive companies to release new models quarterly or even monthly, then retiring a batch of custom processors along with the old model results in negligible economic loss. Compared to the fractional extra cost of discarding chips, the benefits brought by specialized hardware are transformative. One Hardwired LPU node can replace a GPU cluster composed of approximately 2,000 H100s. Due to the drastic reduction in the number of chips required, data center construction, operations, electricity, and embodied carbon emissions can be reduced by hundreds or thousands of times. Facing the disorderly growth of inference demand, manufacturing “monthly disposable” hardware and achieving model updates through physical replacement will inevitably become a more economical and environmentally friendly choice.

In our view, although it is technically implemented in the most immutable way, the Hardwired LPU is still a General-purpose processor. It implements an LLM of general utility, so the hardware automatically inherits general utility. It implements token processing directly at the hardware level, where instructions and data are input as prompts to program general problems. Its software-hardware interface (ISA) is natural language, which cannot be exclusively monopolized by any single company, nor does it require the assistance of software like CUDA, naturally eliminating the problem of ecosystem monopoly. Its computational energy efficiency can reach the 1 POPS/W level, forming the fourth paradigm of general-purpose processors (Token, POPS/W) following CPU (Scalar, MOPS/W), GPU (Vector, GOPS/W), and NPU (Matrix, TOPS/W).

While the original concept existed, the technical solution was not easily achieved in a short time. For those without long-term persistence in exploration, the idea sounded overly impractical until the entire path was mapped out. Consequently, since the birth of ChatGPT, individuals and companies have constantly proposed similar ideas, only to abandon them quickly. To our knowledge, UIUC, ARM, UMich, and the University of Tokyo are working on similar technologies but have not reached a level capable of supporting LLMs. Startups like Etched and Taalas have publicly announced plans to launch similar products but have not yet succeeded. Our team began researching reasonable technical solutions for realizing Hardwired LPUs in March 2023. By traversing multiple research fields—algorithms, architecture, and process technology—and organically combining two key innovations, we finally reduced the estimated cost to a reasonable range:

• Hardwired-Neurons (HN) Architecture: Through a mathematical transformation (which the team calls Primitivization), model weight information is extracted from the 2D planar device array and expressed instead by the topological structure of metal wires in 3D space. Each weight only needs to be encoded as a connection relation of a metal wire. This method utilizes a resource in architectural design that is usually only considered during physical design—the vertical space of metal wiring in the chip’s back-end-of-line (BEOL) process—achieving an approximate 15x increase in encoding density.

• Sea-of-Neurons (SoN) Design Methodology: Following HN’s expression of model weights as metal wires, we further constrained the circuit structures related to model weights to only the M8-M11 metal layers. This means that out of the 70 mask layers required for each chip, 60 layers can be reused across different models. SoN reduces the mask costs for the first tape-out and subsequent revisions by 86% and 92% respectively, making chip customization for specific models economically attractive.

The academic prototype system “HNLPU” has been verified through post-layout simulation. Under the conditions of a 5 nm process and an area of 16×827 mm², it fully implements the OpenAI open-source model GPT-OSS 120B, maintaining the model’s native FP4 numerical precision. In a single-card comparison, its performance and energy efficiency reached 5,555 times and 1,047 times that of the Nvidia H100, respectively. In a large-scale deployment scenario with annual hardware updates, HNLPU offers advantages of 1,496x in operating costs (OpEx), 41.7x in total cost of ownership (TCO), and 357.2x in total carbon emissions compared to a GPU cluster.

The Hardwired LPU is a processor paradigm defined by the following: it possesses general, programmable devices (for calculating attention, positional encoding, activation, sampling, etc.), a memory hierarchy for storing key-value sequences, a high-speed interconnect system spanning multiple chips, and a massive constant matrix operation unit (the HN array). The common structure of the HN array is prefabricated by upstream companies. After downstream enterprises determine the LLM model weights and place an order, metallization is completed to write the weights, allowing for the rapid and low-cost production of ultra-efficient chips equipped with various models. We predict this new industrial model will replace the GPU’s monopoly in the long-term, large-scale inference deployment market, though it requires the strongest possible deep integration between model companies, chip design firms, and the semiconductor manufacturing supply chain. Although many challenges remain, the history of Structured ASICs and Gate Arrays (Intel still operates an eASIC business), which once held a place in the chip industry, proves that this path is inherently viable.

This year, we expect to continue releasing fine-tuning architecture schemes and new layout designs to achieve low-cost model fault tolerance and hot-patching, further reducing the revision cost of changing models by more than ten times. Simultaneously, we are actively advancing the trial production of proof-of-concept prototype systems. Supported by the National Key R&D Program, we are fabricating a hardwired human-computer interaction processor chip based on the HN structure using domestic process technology. It is expected to achieve a computational energy efficiency of no less than 1.5 POPS/W @ FP4. This chip is poised to become the world’s most energy-efficient standard CMOS digital logic chip.

We invite colleagues to follow our latest progress.

(This article contains excerpts from the draft of “Hardwired LPU: A New Processor Paradigm in the Era of Large Language Models.” Original authors: ZHAO Yongwei, HAO Yifan, LIU Yang, CHEN Yi, CHEN Yunji)

(ZH-EN Translation by DeepSeek-V3-0324)

Background

Recently, the national key lab planned to initiate work on intelligent circuit design, and while preparing materials, we often needed circuit schematics as examples. However, there were no readily available tools to conveniently draw these schematics. Existing schematic drawing tools require manually placing each component and connecting each wire one by one. Even drawing a small-scale example involves tedious operations. For instance, a 4-bit adder contains hundreds of transistors, and manually placing components for larger examples becomes nearly as labor-intensive as designing a 4004 processor from scratch. Moreover, real synthesis and placement and routing (PNR) tools are not only cumbersome to invoke but also produce visually unappealing diagrams. As a result, I have always resorted to writing my own scripts to generate these diagrams.

Initially, I used Python with PIL (Python Imaging Library) to draw the schematics. I wrote a few basic functions to draw transistors and logic gates at specified positions, but the placement and routing still had to be designed manually, resulting in low efficiency. I then decided to add simple placement and routing logic to the script to automate the entire workflow from input circuit netlists to output schematics. Here is the final result:

For computationally intensive tasks like placement and routing, I immediately abandoned Python in favor of C++ to avoid waiting an entire day to generate a single diagram. The routing algorithm reused a maze routing implementation I had previously written for another project. I then pasted the routing code into DeepSeek and asked it to complete the placement code. DeepSeek-V3-0324 demonstrated impressive coding capabilities, producing functional and complete code that worked right away. With the help of the large language model, along with some research and clever ideas, abandoning PIL did not cause much pain. DeepSeek provided code for outputting PNGs, extracting BDF bitmap fonts, and helped debug a bug in the diagonal line rasterization algorithm.

Thus, I quickly implemented this ~800-line automated circuit placement and routing program that takes BLIF as input and outputs PNGs. It is entirely self-contained, with dependencies limited to the C++ Standard Template Library (STL). Even as someone who has been programming since childhood, this coding experience was entirely new to me, showcasing how large language models are reshaping the programming experience. Without such models, reading RFCs for PNG, BDF, DEFLATE, and writing low-level functions would have been extremely painful, and I would never have attempted it.

Abandoning Python and returning to C++ has never been easier, and presumably, the same applies to abandoning CUDA. In this era, ecosystem barriers are being weakened across the board—an exciting trend!

Code

The code uses C++26. Currently, only GCC 15.0 can compile it, but since 15.0 has not yet been released, I compiled the latest GCC from source.

This is my first extensive use of the ranges library, and the experience has been great—it often saves a lot of unnecessary verbosity. The current ranges implementation still lacks some features, such as a reverse drop (removing elements from the end of a range). It’s also inconvenient to determine whether a loop iteration is the first or last, making the old-fashioned way of writing loops sometimes preferable.

The C++ standard also lacks some key containers, such as a trie—why hasn’t this been added yet? The performance of some older containers has long been criticized, and it’s time for a refresh. The latest standard provides flat_map, but in my testing, it showed no performance advantage in my use case. I didn’t see an unordered_flat_map. Hopefully, these new features will avoid repeating the mistakes of regex.

Here’s an analysis of the code by DeepSeek-R1:

Usage

The input is a netlist in BLIF format. You can use the open-source tool Yosys to synthesize it from Verilog.

For example, I first asked DeepSeek to write a 4-bit adder as an example:

1
2
3
4
5
6
7
8
9
10
11
12

module adder4 (  
    input [3:0] a,      // 4-bit input a  
    input [3:0] b,      // 4-bit input b  
    input cin,         // Carry input  
    output [3:0] sum,  // 4-bit sum output  
    output cout        // Carry output  
);  

// Implement the adder using continuous assignment  
assign {cout, sum} = a + b + cin;  

endmodule  

Next, create a standard cell library. Currently, the code only supports NOT and NAND gates, so we need to instruct Yosys to use only these two cell types.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

library(demo) {  
  cell(NOT) {  
    area: 3;  
    pin(A) { direction: input; }  
    pin(Y) { direction: output;  
              function: "A'"; }  
  }  
  cell(NAND) {  
    area: 4;  
    pin(A) { direction: input; }  
    pin(B) { direction: input; }  
    pin(Y) { direction: output;  
             function: "(A*B)'"; }  
  }  
}  

Launch Yosys and run the following commands:

1
2
3
4
5
6
7

read_verilog adder4.v  
hierarchy -check  
abc -liberty nandcell.lib  
synth  
splitnets -ports  
opt  
write_blif -top adder4  

The last step outputs the netlist in BLIF format. Copy and save it to a file named adder4.blif.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58

# Generated by Yosys 0.51+104 (git sha1 c08f72b80, g++ 13.3.0-6ubuntu2~24.04 -fPIC -O3)  

.model adder4  
.inputs a[0] a[1] a[2] a[3] b[0] b[1] b[2] b[3] cin  
.outputs sum[0] sum[1] sum[2] sum[3] cout  
.names $false  
.names $true  
1  
.names $undef  
.subckt NOT A=a[3] Y=$abc$170$new_n15  
.subckt NOT A=b[3] Y=$abc$170$new_n16  
.subckt NOT A=a[2] Y=$abc$170$new_n17  
.subckt NOT A=b[2] Y=$abc$170$new_n18  
.subckt NOT A=a[1] Y=$abc$170$new_n19  
.subckt NOT A=b[1] Y=$abc$170$new_n20  
.subckt NOT A=a[0] Y=$abc$170$new_n21  
.subckt NOT A=b[0] Y=$abc$170$new_n22  
.subckt NOT A=cin Y=$abc$170$new_n23  
.subckt NAND A=a[3] B=b[3] Y=$abc$170$new_n24  
.subckt NAND A=a[2] B=b[2] Y=$abc$170$new_n25  
.subckt NAND A=$abc$170$new_n17 B=$abc$170$new_n18 Y=$abc$170$new_n26  
.subckt NAND A=$abc$170$new_n25 B=$abc$170$new_n26 Y=$abc$170$new_n27  
.subckt NOT A=$abc$170$new_n27 Y=$abc$170$new_n28  
.subckt NAND A=a[1] B=b[1] Y=$abc$170$new_n29  
.subckt NAND A=a[0] B=b[0] Y=$abc$170$new_n30  
.subckt NAND A=$abc$170$new_n21 B=$abc$170$new_n22 Y=$abc$170$new_n31  
.subckt NAND A=$abc$170$new_n30 B=$abc$170$new_n31 Y=$abc$170$new_n32  
.subckt NOT A=$abc$170$new_n32 Y=$abc$170$new_n33  
.subckt NAND A=cin B=$abc$170$new_n33 Y=$abc$170$new_n34  
.subckt NAND A=$abc$170$new_n30 B=$abc$170$new_n34 Y=$abc$170$new_n35  
.subckt NOT A=$abc$170$new_n35 Y=$abc$170$new_n36  
.subckt NAND A=$abc$170$new_n19 B=$abc$170$new_n20 Y=$abc$170$new_n37  
.subckt NAND A=$abc$170$new_n29 B=$abc$170$new_n37 Y=$abc$170$new_n38  
.subckt NOT A=$abc$170$new_n38 Y=$abc$170$new_n39  
.subckt NAND A=$abc$170$new_n35 B=$abc$170$new_n39 Y=$abc$170$new_n40  
.subckt NAND A=$abc$170$new_n29 B=$abc$170$new_n40 Y=$abc$170$new_n41  
.subckt NOT A=$abc$170$new_n41 Y=$abc$170$new_n42  
.subckt NAND A=$abc$170$new_n28 B=$abc$170$new_n41 Y=$abc$170$new_n43  
.subckt NAND A=$abc$170$new_n25 B=$abc$170$new_n43 Y=$abc$170$new_n44  
.subckt NOT A=$abc$170$new_n44 Y=$abc$170$new_n45  
.subckt NAND A=$abc$170$new_n15 B=$abc$170$new_n16 Y=$abc$170$new_n46  
.subckt NAND A=$abc$170$new_n24 B=$abc$170$new_n46 Y=$abc$170$new_n47  
.subckt NOT A=$abc$170$new_n47 Y=$abc$170$new_n48  
.subckt NAND A=$abc$170$new_n44 B=$abc$170$new_n48 Y=$abc$170$new_n49  
.subckt NAND A=$abc$170$new_n24 B=$abc$170$new_n49 Y=cout  
.subckt NAND A=$abc$170$new_n23 B=$abc$170$new_n32 Y=$abc$170$new_n51  
.subckt NAND A=$abc$170$new_n34 B=$abc$170$new_n51 Y=$abc$170$new_n52  
.subckt NOT A=$abc$170$new_n52 Y=sum[0]  
.subckt NAND A=$abc$170$new_n36 B=$abc$170$new_n38 Y=$abc$170$new_n54  
.subckt NAND A=$abc$170$new_n40 B=$abc$170$new_n54 Y=$abc$170$new_n55  
.subckt NOT A=$abc$170$new_n55 Y=sum[1]  
.subckt NAND A=$abc$170$new_n27 B=$abc$170$new_n42 Y=$abc$170$new_n57  
.subckt NAND A=$abc$170$new_n43 B=$abc$170$new_n57 Y=$abc$170$new_n58  
.subckt NOT A=$abc$170$new_n58 Y=sum[2]  
.subckt NAND A=$abc$170$new_n45 B=$abc$170$new_n47 Y=$abc$170$new_n60  
.subckt NAND A=$abc$170$new_n49 B=$abc$170$new_n60 Y=$abc$170$new_n61  
.subckt NOT A=$abc$170$new_n61 Y=sum[3]  
.end  

Compile cirschem.cpp: g++ cirschem.cpp -std=c++26 -O3

Run: ./a.out adder4.blif

After waiting about half a minute for placement and routing to complete, the output will be saved as schematic.png.

(ZH-EN Translation by DeepSeek-V3-0324)

With the rise of large language models, 4-bit quantized models have become the preferred choice for deployment due to computational and storage budget constraints. However, we observe that traditional computing architectures suffer from significant computational redundancy when handling low-precision matrix multiplication.
We propose the Cambricon-C architecture, which aims to revolutionize the implementation of 4-bit matrix computation through an innovative Primitivized Matrix Multiplication (PMM) algorithm.

An in-depth analysis of 4-bit matrix multiplication reveals astonishing redundancy: in the 4-bit quantized version of Llama2-7B, computing the activation value of a single output neuron requires 11,008 4-bit multiplications followed by accumulation. Since 4-bit data has at most 16 possible values, the pigeonhole principle dictates that over 97% of these multiplications must be repeated. Traditional matrix processors based on MAC units redundantly compute these identical products, resulting in enormous energy waste.

The PMM algorithm proposes: By leveraging the inverse distributive property of multiply-accumulate operations, we first count the occurrences of each multiplication and then perform a unified weighted summation, significantly reducing computational intensity.

1. Precomputation to Eliminate Redundant Multiplications: Precompute all possible 4-bit product combinations (256 in total) as a lookup table. During actual computation, only table lookup is required instead of repeated calculations.
2. Reduced Addition Intensity: Transform accumulation operations into incremental updates of 256 counters, replacing traditional adder trees with unary successor operations (i.e., counting), drastically cutting addition overhead.

Based on the PMM algorithm, we designed the Cambricon-C architecture:

• Quarter-Square Multiplication (QSM): We revisited the 19th-century quarter-square multiplication technique, reducing the number of counters from 256 to 29 (referred to as the “R29 scheme”), achieving an 8.6× improvement in area efficiency.
• Ripple Counters: Implemented counters using self-timed D flip-flop chains, reducing dynamic power by 51% compared to SRAM-based solutions.

Experimental results show:

• A single PE achieves 1.97× higher energy efficiency than traditional MAC units.
• A full 32×32 array accelerator delivers a 1.25× end-to-end energy efficiency improvement for LLaMA2 inference.
• Performance advantages scale with matrix size.

Cambricon-C pioneers a new paradigm of “primitivized computation”. By decomposing complex operations into fundamental counting operations, we redefine the efficiency limits of low-precision computing:
The future efficiency of accelerators will no longer depend on intricate MAC unit designs but rather on holistic optimization of entire matrix operations.
In the near future, we will publish more work on “primitivized computation”, continuing to drive innovation in the microarchitecture of deep learning processors. We welcome attention and collaboration from peers.

Published at MICRO 2024. [DOI] [PDF]

Problem

Given a Gomoku board situation, identify all threat points. Threat points refer to the positions where both sides can form a winning situation after placing a piece, such as the end of a four-in-a-row (winning, immediately forming a five-in-a-row), and both ends of a living three (forming a living four after placing a piece, if the opponent does not win in this step, it will be a winning situation in the next turn). For example, the positions marked with “x” in the following board.

┌┬┬┬┬┬┬┬┬┬┬┬┬┬┐

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼x┼┼┤

├┼┼┼┼┼┼┼┼┼●┼┼┼┤

├┼┼┼x○○○○●┼┼┼┼┤

├┼┼┼┼┼┼┼●┼┼┼┼┼┤

├┼┼┼┼┼┼x┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

├┼┼┼┼┼┼┼┼┼┼┼┼┼┤

└┴┴┴┴┴┴┴┴┴┴┴┴┴┘

During the search of AI programs, threat points must be quickly found in order to prune other moves. Because placing a piece at other positions when there is a threat point will always lead to a disadvantage.

Naive Method

A simple solution is to traverse according to the rules. I believe most people solve it this way when they first start writing programs, and I am no exception. After all, premature optimization is evil. But eventually, you will find that the speed of this process is crucial.

The first solution that comes to mind is vectorization. Use a bitboard to represent the situation, dividing the board into black and white, each board with 225 bits, which can be placed in uint16_t [16], or a YMM register __m256i. Through AVX instructions such as _mm256_slli_epi16/_mm256_srli_epi16 (VPSLLW/VPSRLW) to shift and align the entire board, and then perform logical AND to find living four, blocked four, living three, etc. This saves the overhead of enumerating the board, but due to the complexity of the threatening shapes (blocked four and living three are not necessarily continuous), the program is not simple and still takes a relatively long time.

Ternary Coding

In 2014, I independently discovered a solution and applied it to the evaluation function of kalscope v2 “Dolanaar”. This program was a practice work for me when I was an undergraduate, but it also achieved competitive strength at that time.

I treat each Yang line (horizontal or vertical line) on the board separately. There are 15 places on a line, I list all possible situations, calculate the threat points in advance to form a lookup table. However, if represented by a bitboard, the encoding of a line can reach , and the size of the lookup table reaches 2 GB. Note that there is a lot of redundancy in the bitboard encoding, with 2 bits encoding 3 situations: empty “┼” 0b00, white piece “○” 0b01, black piece “●” 0b10, and 0b11 is illegal. To achieve compact encoding, ternary counting is needed, 0 represents empty, 1 represents white piece, 2 represents black piece, each ternary digit represents a place.

The conversion to ternary involves a lot of division and modulus operations, thus it is unrealistic. Note that the board situation does not appear suddenly, but is always played gradually from an empty board, and ternary representation can be incrementally maintained. The empty line is encoded as 0; placing a white stone at the -th position is equal to adding to the encoding; placing a black piece at the -th position is equal to adding to the encoding. On backtrack, subtract the increased value from encoding. To implement efficiently, we only need to keep another lookup table of the powers of three, power3.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29

constexpr int power3[16] {
 1, 3, 9, 27, 81, 243, 729, 2187, 6561, 19683, 59049, 177147, 531441, 1594323, 4782969, 14348907
};
int triboard_h[15] {};
int triboard_v[15] {};
uint16_t threat_lut[14348907] { ... }; // pre-generated.

void make(int x, int y, int color) {
    triboard_h[y] += power3[x] << color;
    triboard_v[x] += power3[y] << color;
}
void unmake(int x, int y, int color) {
    triboard_h[y] -= power3[x] << color;
    triboard_v[x] -= power3[y] << color;
}
void find_threat() {
    for (int y = 0; y < 15; y++) {
        uint16_t threat = threat_lut[triboard_h[y]];
        while (threat) {
            int x = std::countr_zero(threat);

            // do something to the threat (x,y).

            threat &= threat - 1;
        }
    }
    // another loop for vertical lines.
    // TODO: how about diagonal / anti-diagonal lines?
}

The lookup table threat_lut tightly encodes all possible situations of a Yang line on a 15-line board, occupying only 27.4 MB of memory.

This method is designed specifically for a standard 15-line board, but at that time, the Gomocup rules were using 20-line boards, which made this method invalid. Otherwise, I might have opted in the competition. I am not sure if anyone else has discovered a similar method earlier by me, but I did notice that Rapfi, recent champion of Gomocup, also implemented this technique in it.

This method is not perfect. It is only for a fixed-length Yang line, and using it for the 20-line rules of Gomocup will result in an oversized LUT. In addition, the length of each Yin line (diagonal or anti-diagonal line) is different, and if you forcibly apply the lookup table for the Yang line to Yin lines, it will lead to false positives. For example, on a Yin line with a total length of only 4 places, a living three is not a threat because the board boundary will naturally block it. How should this be handled?

I did not give an answer in 2014, but left an unsolved problem.

Skew Ternary Coding

Last year, I guided doctoral student Guo, Hongrui to complete Cambricon-U, which implemented the in-place counting in SRAM using skew binary numbers. Skew numbers can not only be used for binary, but also for ternary, solving the abovementioned problem.

Skew ternary allows a digit “3” to temporarily exist in the number without carrying over. When counting, if there is no digit “3” existing, increment the least significant digit by 1; If there is a digit “3”, reset the “3” to “0” and carry over to the next digit. Under such counting rules, it is easy to find that skew ternary has the following two properties:

• There is at most one digit “3” in the number;
• If there is a “3”, all less significant digits must be “0”.

We use “0” to represent an empty place, “1” to represent a white stone, “2” to represent a black stone, and “3” to represent the board boundary. Such a coding can naturally and compactly encode various lengths of Yin lines into the lookup table. When the length of the Yin line is less than 15, only the high digits are used, and the boundary is set to “3”, and the less significant digits are “0”. The items in the lookup table are compiled in the order of skew ternary counting:

The lookup table has a total of 21523359 items, tightly encoding all situations of length-15 Yang lines and Yin lines from length 1 to 15, occupying 41.1 MB of memory. The code is slightly modified, changing the ternary weight power3 to the skew ternary weight basesk3, which is the sequence OEIS A261547.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46

constexpr int basesk3[16] {
    1, 4, 13, 40, 121, 364, 1093, 3280, 9841, 29524, 88573, 265720, 797161, 2391484, 7174453, 21523360
};
int sk3board_h[15] {};
int sk3board_v[15] {};
int sk3board_d[29] {};
int sk3board_a[29] {};
uint16_t threat_lut[21523359] { ... }; // pre-generated.
   
// set board bounds on init. 
void init() {
    for (int i = 0; i < 15; i++) {
        sk3board_d[i] = basesk3[13-i] * 3;
        sk3board_a[i] = basesk3[13-i] * 3;
    }
    for (int i = 15; i < 29; i++)
        sk3board_d[i] = basesk3[i-15] * 3;
        sk3board_a[i] = basesk3[i-15] * 3;
    }
}

void make(int x, int y, int color) {
    sk3board_h[y]      += basesk3[x]             << color;
    sk3board_v[x]      += basesk3[y]             << color;
    sk3board_d[14-x+y] += basesk3[x<y?y:x]       << color;
    sk3board_a[x+y]    += basesk3[x+y<15?14-y:x] << color;
}
void unmake(int x, int y, int color) {
    sk3board_h[y]      -= basesk3[x]             << color;
    sk3board_v[x]      -= basesk3[y]             << color;
    sk3board_d[14-x+y] -= basesk3[x<y?y:x]       << color;
    sk3board_a[x+y]    -= basesk3[x+y<15?14-y:x] << color;
}
void find_threat() {
    for (int y = 0; y < 15; y++) {
        uint16_t threat = threat_lut[sk3board_h[y]];
        while (threat) {
            int x = std::countr_zero(threat);

            // do something to the threat (x,y).

            threat &= threat - 1;
        }
    }
    // to process vertical, diagonal and anti-diagonal lines
}

We have completely transformed the very cumbersome threat point enumerating problem into incrementally maintaining the board representation in four directions, and then performing table lookups. A lookup table of about 41 MB can now as a whole fit in the L3 cache, thus the speed is extremely fast. In the above example, find_threat can be further vectorized, and the overhead of the entire operation can be compressed to about a hundred clock cycles.

The same idea can be used for winning checks, evaluation functions, etc. If you are preparing to use an architecture similar to AlphaGo, you might as well try to change the first layer of the neural network to 1D convolution along the four lines, and this method can directly save the first layer convolution operation with lookup.

Quick thoughts after Cambricon-U.

RIM shows our perspective of PIM as researchers who did not studied in the field. When it comes to the concept of PIM, people intuitively first thought of the von Neumann bottleneck, and believe that PIM is of course the only way to solve the bottleneck. In addition to amateur CS enthusiasts, there are also professional scholars who write various materials in this caliber.

This is not true.

In RIM’s journey, we showed why previous PIM devices could not do counting. Counting (1-ary successor function) is the primitive operation that defined Peano arithmetic, and is the very basis of arithmetic. We implemented the counting function in PIM for the first time by using a numeral representation that is unfamiliar to scholars in the field. But this kind of trick can only be done once, and the nature will not allow us for a second time, because it is proven that no representation can efficiently implement both addition and multiplication*.

Computation implies flowing of information, while PIM wants them to stay in place. Contradictory. Combined with Mr. Yao’s conclusion, we may not be able to efficiently complete all the basic operations in-place in storage. If I were to predict, PIM would never really have anything to do with the von Neumann bottleneck.

* Andrew C. Yao. 1981. The entropic limitations on VLSI computations (Extended Abstract). STOC’81. [DOI]

At the MICRO 2023 conference, we demonstrated the invention of a new type of process-in-memory device “Random Increment Memory (RIM)” and applied it to the stochastic computing (unary computing) architecture.

In academia, PIM is currently a hot research field, and researchers often place great hopes on breaking through the von Neumann bottleneck with subversive architecture. However, the concept of PIM has gradually converged to specifically refer to the use of new materials and devices such as memristors and ReRAM for matrix multiplication. In 2023, from the perspective of outsiders like us who have not studied in PIM before, PIM has developed into a strange direction: yet developed PIM devices can process neural networks and simulate brains, but still unable to do the most basic operations such as counting and addition. In the research of Cambricon-Q, we left a regret: we designed NDPO to complete the weight update on the near-memory side, but it was unable to achieve true in-memory accumulation, so the reviewer criticized “it can only be called near-memory computing, not in-memory computing.” From this, we began to think about how to complete in-place accumulation in the memory.

We quickly realized that addition could not be done in-place with binary representation. This is because although addition is simple, there is carry propagation: even adding 1 to the value in memory may cause all bits to flip. Therefore, a counter requires a full-width adder to complete the self-increment operation (1-ary successor function), thus all data bits in the memory must be activated for potential usage.

Fortunately, the self-increment operation results in only two bit flips on average. We need to find a numeral representation that limits the number of bits flipping in the worst case. Therefore we introduced Skew binary number system to replace binary numbers. The skew binary was originally proposed for building new data structures, such as the Brodal heap, to limit the worst-case time complexity when a heap merges. It is very similar to the case here, that is, limiting carry propagation.

We base on SRAM in conventional CMOS technology to design RIM. We store digits (in skew binary) in column direction, and use an additional column of SRAM cells to store the skew bit of each digit (that is, where the digit “2” is in the skew number). The self-increment operations are performed as follows:

• If any digit in the current skew number is “2”, set “2” to “0” and increase the next digit by 1;
• Otherwise, increase the least significant digit by 1.

Although the row-index of “2” in each skew number are different (the skew counting rules promise that at most one digit can be “2”), the cells to be operated on will be randomly distributed in the memory array. It cannot be activated according to the row selection of SRAM. But we can use the latched skew bit to activate the corresponding cells to operate, so that cells located in different rows can be activated in the same cycle!

Finally, we achieved a 24T RIM cell. Not using new materials, but built entirely from CMOS. RIM can process random self-increment on stored data: in the same cycle, each data can self-increment or remain unchanged on demand.

We apply RIM in stochastic computation (unary computation). A major dilemma in stochastic computing is the cost of conversion between unary numbers and binary numbers. Converting binary numbers to unary numbers requires a random number generator, and converting back requires a counter. Because unary numbers are long (up to thousands of bits), after computing in unary, the intermediate results must be converted back to binary for buffering. As a result, counting operations can account for 78% energy consumption. We use RIM to replace the counters in uSystolic and propose the Cambricon-U architecture, which significantly reduced the energy consumption of counting operations. This work solves a key problem of stochastic deep learning processors, boosting the application of these technologies.

Published on MICRO 2023. [DOI] [PDF]