Byzantine-linearizable Single-Writer Multi-Reader Register in the GPU-Inspired Shared Memory Model

Motivation

• GPUs emphasize throughput across many lock-step threads and offer limited fine-grained shared-memory access control, while multi-tenancy and third-party code expose them to Byzantine-like behavior.To a shared object, a buggy, compromised, or misbehaving co-resident thread is indistinguishable from a Byzantine process.Source: §1 · p. 1
• GPU threads may invoke writes despite classical SWMR ownership assumptions, but lock-step execution and bounded scheduling may help mitigate Byzantine interference.The paper studies this tension in a model of a single thread block on one Streaming Multiprocessor [5].Source: §1 · pp. 1–2

Contributions

• The paper formulates Byzantine-linearizable SWMR registers in the GPU-inspired shared-memory model [5] and presents three authenticated constructions: CAReg, SnapObject, and idealized signature-aware memory.Except for the idealized construction, base objects store and return values without inspecting signatures; processes perform authentication.Source: §1 · p. 2
• CAReg uses constant memory with probabilistic resilience, while SnapObject retains a bounded window of authenticated writes to eliminate invalid reads.The CAReg analysis considers scheduler bias and Byzantine interference.Source: §1 · p. 2
• SnapObject has a lock-free linked-list implementation and three wait-free implementations; the wait-free variants prevent adversarial schedulers from starving any correct process, including the writer.The wait-free variants use a circular array, a Herlihy Wing inspired queue with snapshot support, or a fixed-size atomic array; the realizations trade off memory and read complexity.Source: §1 · p. 2

• Atomic registers make each operation appear instantaneous between invocation and response, respect real-time precedence, and support higher-level objects.Source: §2 · p. 2
• Crash-tolerant atomic registers have been extensively studied, while Kowalski et al. established tight resilience bounds for Byzantine SWMR registers in asynchronous message-passing systems.Source: §2 · p. 2
• Cohen and Keidar [3] defined Byzantine linearizability by extending the correct-process projection of an execution with Byzantine operations to obtain a linearizable history.They applied the framework to reliable broadcast, atomic snapshot, and asset transfer from SWMR registers.Source: §2 · p. 2
• Hu and Toueg [9] showed that wait-free linearizable SWMR registers cannot be implemented from crash-tolerant SWMR registers under Byzantine failures, and introduced verifiable and sticky registers [10], later improved by Kshemkalyani et al..Prior results target fully asynchronous systems; the GPU-inspired model of Georgiou et al. [5] captures phase-based execution and lock-step scheduling, whose impact on Byzantine linearizability and progress had not been studied.Source: §2 · p. 2
• The proposed GPU-model register assumes no memory access control, pure asynchrony, or sticky/verifiable memory, instead using authentication, phase-based lock-step execution, and mandatory writer scheduling [10].Source: §2 · p. 2
• Wait-free SnapObject realizations order insertions with unconditional fetch-and-increment, rather than CAS retry loops or helping, so correct-process operations finish within bounded own steps despite Byzantine interference and adversarial scheduling.Source: §2 · p. 2

GPU-Constrained Shared Memory Model

• The model follows Georgiou et al. [5] and focuses on n GPU-thread processes in one block on one SM, grouped into lock-step SIMT warps.Warp size is p (typically 32 on modern NVIDIA GPUs), and the block contains ⌈n/p⌉ warps.Source: §3 · p. 3

Failure Model and Phase-Based Execution

• Crash processes halt, while Byzantine processes may deviate arbitrarily, but cannot bypass warp scheduling, impersonate others, or exceed their scheduled phase [5].Each phase schedules exactly one warp for bounded instructions; divergence is serialized into operation groups that execute sequentially within the phase.Source: §3 · p. 3

Scheduler Fairness and Bias

• The scheduler may be unbiased or favor Byzantine-containing phases, with bias factor γ ≥ 1; γ = 1 denotes an unbiased scheduler.A cycle is a sequence of phases in which every process is scheduled at least once, spanning at least ⌈n/p⌉ phases.Source: §3 · p. 3

Shared Memory

• Shared memory is reliable and accessible only through object interfaces, but has no access control: any process may invoke any exposed operation [3] [5].Processes cannot access memory outside their addressable scope, including another thread’s private registers or local memory; removing phase-based scheduling and lock-step execution yields a classical asynchronous shared-memory system. Independent Thread Scheduling weakens strict intra-warp lock-step, and extension to it is left for future work.Source: §3 · p. 3

On Replay

• A valid signature authenticates a value’s origin but not its freshness, so Byzantine processes can replay previously signed values.Unless stated otherwise, Byzantine linearizability and resilience bounds hold against general adversaries that may replay; CAReg’s freshness and resilience bounds instead assume non-replaying adversaries.Source: §3 · p. 4

4 Byzantine Linearizability and SWMR Register

• Linearizability requires an execution, after resolving some pending calls, to be equivalent to a legal sequential history that preserves real-time precedence.A history is well-formed when each process’s invocations and responses alternate, beginning with an invocation.Source: §4 · p. 4
• Byzantine linearizability [3] requires an extension that preserves correct processes’ projected behavior and is itself linearizable, despite arbitrary Byzantine behavior.The projection removes events generated by Byzantine processes.Source: §4 · p. 4
• An SWMR register has one writer and multiple readers, and each read returns the value of the most recent preceding write.Base objects are non-corruptible and accessed only through operations; except for the idealized Section 7 object, they do not inspect or verify signatures.Source: §4 · p. 4
• The hardware primitives used are atomic compare-and-swap (CAS) and fetch-and-increment (FAI).CAS conditionally updates a memory location when it matches an expected value; FAI increments a location and returns its previous value.Source: §4 · p. 4
• Within a phase, correct readers executing the same read instruction observe the same shared-memory state and return the same value, as illustrated in Figure 2.Source: §4 · pp. 4–5, Figure 2 · p. 5

5.1 Design Features

• The design limits adversarial interference by restricting Byzantine processes to register operations and requiring the writer to rewrite its current value in every phase.The shared memory is non-corruptible, the underlying register is crash-tolerant, linearizable, and wait-free, and signatures let readers distinguish writer-generated values.Source: §5 · p. 5, §5.1 · p. 5

5.2 Algorithmic Solution

• CAReg signs each writer value before storing it, and readers return the value only when its signature verifies; otherwise they return ⊥.Atomicity, authentication, and phase-based execution together support atomic semantics for correct processes despite Byzantine behavior.Source: §5.2 · pp. 5–6

5.3 Correctness Overview

• CAReg preserves atomicity through the underlying register, value validity through authentication, and intra-phase reader consistency through lock-step execution.The construction also satisfies the ordering conditions for Byzantine linearizability; full proofs are in Appendix A.3.Source: §5.3 · p. 6

5.4 Probabilistic Resilience Analysis

• Adversarial scheduling and register contention can probabilistically suppress legitimate values, with spreading Byzantine readers across phases (β = 1) favoring the adversary.The analysis defines P(⊥) as the probability a correct reader returns ⊥ during a cycle and considers biased or unbiased scheduling and register arbitration.Source: §5.4 · p. 6

5.5 Discussion and Limitations

• The single-cell CAReg design can return ⊥ under contention and cannot detect replay of an older validly signed value, which may yield a stale read.P(⊥) remains strictly positive and grows with the number of Byzantine processes, even when the writer, scheduler, and register are unbiased.Source: §5.5 · p. 6
• A window of the most recent p+1 entries guarantees a correct reader encounters an authenticated value, yielding P(⊥) = 0 and resilience up to f ≤ n−1.The writer contributes one authenticated value per phase, while at most p Byzantine insertions can occur within a phase.Source: §5.5 · pp. 6–7

6 Byzantine-Linearizable SWMR Register Using a SnapObject

• SnapObject appends inserts in linearization order and returns the p+1 most recent values, preserving multiple candidates instead of overwriting one shared location.The section assumes a fair, unbiased scheduler; readers distinguish the initial register value NULL from ⊥, which means no valid signed value was found.Source: §6 · p. 7

6.1 Design Features

• The single-insert-per-process-per-phase assumption bounds writes to p+1 entries per phase, while allowing longer phases requires increasing queue capacity.The authors say this bound fits short, synchronous phases and avoids reliance on sticky objects otherwise needed with unbounded concurrent Byzantine writes.Source: §6.1 · p. 7

6.2 Algorithmic Solution

• Writers sign each value with a sequence number, while readers verify returned entries and select the valid value with the largest sequence number, rejecting stale replays.The sequence number prevents replayed writer-signed entries from appearing fresh; a read returns ⊥ if it finds no valid entry.Source: §6.2 · p. 7

6.3.1 Linked-List Based Atomic Object.

• The linked-list realization inserts at the head with CAS and reads at most p+1 recent entries, making insertion lock-free and ReadLatest wait-free with respect to p.Optional reclamation detaches the suffix after Δ = O(n) nodes to bound memory without reclaiming nodes still in use.Source: §6.3.1 · p. 8, Figure 3 · p. 8

6.4 Wait-Free Implementations

• The three wait-free SnapObject realizations replace the linked list’s CAS retry loop with an unconditional fetch-and-increment.Besides the circular array, they include a Herlihy-Wing-inspired snapshot queue with unbounded memory and a constant-memory atomic array that is correct only within one phase.Source: §6.4 · p. 8, §6.4.1 · p. 9

6.4.1 Wait-Free Circular-Array Based Atomic Object.

• With a correct writer, every correct reader returns ⊥ with probability 0, and the construction is Byzantine-linearizable for up to n−1 Byzantine readers if at least one reader is correct.The circular-array implementation obtains these guarantees while using a window sized to cover the writer’s latest value and Byzantine insertions in a phase.Source: §6.4.1 · pp. 8–9

6.5 Discussion

• The four SnapObject realizations all provide deterministic, replay-resistant reads, but trade off progress, memory use, and whether operations may span phases.The linked list is lock-free with reclaimed O(n) memory; the circular array is wait-free with bounded memory and cross-phase operations; the atomic array uses constant memory but requires single-phase operations; the snapshot queue has unbounded memory.Source: §6.5 · p. 9

7 Idealized Signature-Aware Memory

• The object stores the writer’s public key and last accepted sequence number, rejecting replayed signed pairs while retaining constant memory and one read per phase.A write must have a valid signature and a sequence number strictly greater than the stored one, making the construction robust against the general adversary.Source: §7 · p. 9
• Because in-memory signature verification is complex and implementation-specific, this stronger abstraction primarily marks a design-space upper boundary rather than a directly deployable implementation.The paper discusses possible implementation considerations in Appendix A.7.Source: §7 · p. 9
• Internal signature verification eliminates Byzantine overwrites, achieving maximal f-resilience (n-1) and P(⊥) = 0.The memory object accepts writes only when the writer’s signature is valid.Source: §7 · p. 9, Table 2 · p. 10

8 Conclusion and Future Work

• The three constructions trade off memory and synchronization: CAReg uses constant memory and one read per phase with probabilistic resilience, while SnapObject variants provide deterministic resilience with P(⊥) = 0 at higher costs.The signature-aware abstraction also provides constant memory and deterministic resilience, but requires stronger hardware assumptions.Source: §8 · p. 10
• CAReg stores one value and guarantees freshness only against non-replaying adversaries, whereas SnapObject-based and signature-aware constructions bind sequence numbers to signed values and tolerate general adversaries.The sequence number lets readers compare authenticated writes for freshness.Source: §8 · p. 10
• Lock-step execution, bounded instruction phases, and structured scheduling can partially compensate for the absence of access control in Byzantine shared-memory systems.The work formalizes phase-aware Byzantine linearizability for GPU execution and studies atomic semantics for correct processes despite arbitrary writes by Byzantine processes.Source: §8 · p. 10
• Future work includes building Byzantine-tolerant stacks, queues, and snapshots from the SWMR registers, and studying synchronization power and a Byzantine consensus hierarchy in the GPU-inspired model.The proposed hierarchy would be analogous to consensus-number hierarchies in classical shared memory.Source: §8 · p. 10
• Another open direction is extending the model beyond a single SM to account for GPU synchronization across SMs, where Byzantine linearizability and resilience remain challenging.Source: §8 · p. 10

A.1.1 Authentication using Signatures.

• Each authorized writer receives a private signing key to sign messages and a corresponding public verification key is distributed to readers.The writer keeps the private key secret.Source: §A.1.1 · p. 11
• A signature is bound to its specific message, so even a one-bit modification fails verification and Byzantine readers cannot alter or fabricate authenticated messages without the writer’s private key.Forging a valid signature without that key is computationally infeasible under standard security assumptions.Source: §A.1.1 · p. 11
• Well-established public-key schemes that can realize digital signatures include RSA and DSA.Source: §A.1.1 · p. 11

A.1.2 Authentication with the Help of CUDA Kernel.

• A thread’s globally unique identifier can be computed as tid = blockIdx.x * blockDim.x + threadIdx.x.CUDA exposes blockIdx, blockDim, and threadIdx within a kernel.Source: §A.1.2 · p. 11
• CUDA initializes these indexing variables at kernel launch and exposes them as read-only, so user code cannot modify them.The variables are special registers such as blockIdx, blockDim, and threadIdx.Source: §A.1.2 · p. 11
• The write method appends the computed thread ID to the value, and readers discard values whose ID does not match the pre-designated writer’s.Because the method cannot be altered by Byzantine threads and hardware-assigned IDs cannot be forged or altered within a kernel execution, a malicious thread cannot impersonate another thread’s identifier.Source: §A.1.2 · p. 11

A.2 SWMR Byzantine Linearizability in the GPU-inspired Model

• A read completing in phase j must return a value written by a write that completed in phase j or earlier.This is the No Future Read condition.Source: §A.2 · p. 11
• A read cannot return a value already superseded by a more recent completed write before the read’s phase.This is the No Read from Distant-Past condition.Source: §A.2 · p. 11
• Correct readers must not observe an older write after a later read has already observed a newer write.This is the No New–Old Inversion condition.Source: §A.2 · p. 11
• Correct readers completing reads in the same phase must return the same value.The model executes processes performing the same instruction in lock-step within each phase.Source: §A.2 · p. 11
• These properties hold even when operations complete within one phase and when the writer is Byzantine, provided written values are distinguishable.Correct-writer values are assumed unique, an assumption that can be enforced with sequence numbers.Source: §A.2 · p. 11

• The section states that Algorithm 1 implements a Byzantine-linearizable SWMR register according to Definition 3.Source: §A.3 · p. 11
• The proof first establishes that CAReg remains linearizable even when Byzantine processes issue method invocations.Source: §A.3 · p. 11
• It then shows that every non-⊥ value returned by a correct reader was written by the writer.Even when consecutive writes contain the same value, the reader should return the value associated with the second write; a Byzantine writer's value may be useless to the application without violating event ordering.Source: §A.3 · p. 11
• The final proof step verifies No Future Read, No Read from Distant-Past Value, No New–Old Inversion, and Intra-Phase Reader Consistency, which together imply Byzantine linearizability.Source: §A.3 · pp. 11–12

A.3.1 Object-Level Linearizability.

• Any completed method invocation has a unique linearization point within one phase, and the history of completed invocations remains linearizable whether callers are correct or Byzantine.CAReg’s completed reads and writes are therefore linearizable in the Byzantine GPU-inspired shared-memory model.Source: §A.3.1 · p. 12

A.3.2 Value Validity.

• A correct read returning v≠⊥ returns a value written by the authorized writer, although a Byzantine authorized writer may share its private key.Readers accept a value only when its signature verifies under the authorized writer’s public key; replaying a writer-signed value does not violate this validity property.Source: §A.3.2 · pp. 12–13

A.3.3 Byzantine linearizability. • No Future Read.

• Under a non-replaying adversary, correct reads cannot return future or superseded writes, and later reads cannot regress to older writes; replay can break the latter guarantees.These ordering properties supplement Byzantine linearizability, which the section says rests on value validity alone.Source: §A.3.3 · p. 13
• Correct reads that complete in the same phase return the same value, because lock-step execution makes their register reads and deterministic local computations identical.The reads observe the same logical register state under CAReg linearizability.Source: §A.3.3 · p. 14

A.3.4 Byzantine Linearizability.

• Algorithm 1 implements a Byzantine-linearizable SWMR register: completed histories admit a legal sequential order that preserves correct-process real-time order and projection.The construction inserts Byzantine writes immediately before correct reads to account for the values they return. Under a non-replaying adversary, additional ordering conditions also hold.Source: §A.3.4 · p. 14

A.4 Analysis of Crash-Tolerant Atomic Register

• With unbiased CAReg and scheduler, P(⊥)=fp/[n(β+1)]; with CAReg biased toward Byzantine writes and scheduler bias γ, P(⊥) is (fγ/β)/((n/p)+((γ−1)f/β)).In an individual Byzantine phase, unbiased CAReg gives P(⊥)=β/(β+1); a correct phase has P(⊥)=0.Source: §A.4 · pp. 14–15

A.5 Resilience Analysis of CAReg

• CAReg tolerates a number of Byzantine readers that scales linearly with system size n and inversely with phase size p.This resilience claim is stated for the unbiased atomic-register and unbiased-scheduler setting.Source: §A.5 · p. 15
• A correct read’s probability of returning ⊥ depends on Byzantine-process and scheduler bias, and scheduler bias directly degrades CAReg’s resilience.The bound is defined for scheduler bias factor γ ≥ 1 and target probability q ∈ [0, 1].Source: §A.5 · p. 15

A.5.1 Resilience Analysis (𝛽= 1).

• With β = 1 and q = 1/5, the analysis concludes that unbiased CAReg and scheduling yield strong linear resilience guarantees.The q value is an illustrative example; the register’s resilience is parameterized by q.Source: §A.5.1 · p. 16

Case 1: No Lock-Step with the Writer.

• Without lock-step with the writer, a read is valid exactly when no Byzantine write falls between the writer’s write and the read; P(⊥ | no lock-step) = β/(β + 2).Valid reads occur when the read precedes all writes or immediately follows the writer’s write.Source: Case 1: No Lock-Step with the Writer · p. 16

Case 2: Lock-Step Execution with the Writer.

• With lock-step execution, P(⊥ | Byzantine phase, CAReg unbiased) = β/(β + 1), because Byzantine processes must compete with the writer in the same phase.This is a greater probability of ⊥ than when Byzantine readers execute independently of the writer.Source: Case 2: Lock-Step Execution with the… · p. 16

A.6 SnapObject Implementation

• SnapObject’s four implementations share an abstract interface but use different data structures and operational assumptions, trading simplicity against implementation feasibility.The linked-list and wait-free circular-array atomic objects are detailed in the main paper; this appendix details the other two designs.Source: §A.6 · pp. 16–17

A.6.1 Linked-List Based Atomic Object.

• The linked-list SnapObject inserts signed value-sequence entries with CAS and returns a p + 1-entry window; verification selects the latest writer-signed value and rejects replayed older entries.The successful CAS linearizes writes, and reading the head linearizes reads; the object is linearizable, and once the first write completes, correct reads return the writer’s most recent value rather than ⊥.Source: §A.6.1 · pp. 17–18
• Algorithm 3’s linked-list object implements a Byzantine-linearizable SWMR register, with correct reads returning the writer’s most recent preceding write.ReadLatest returns the most recent p+1 successfully inserted values, newest first, or all existing values if fewer are available.Source: §A.6.1 · p. 19

A.6.2 Queue with snap.

• Algorithm 5’s Herlihy–Wing queue with snap [8] preserves the ordering properties under Byzantine processes while allowing readers to retrieve recent values despite empty array positions.Readers may scan up to Δ extra positions to account for inserts that incremented the counter but did not complete; entries are never reclaimed, so storage grows with insertions.Source: §A.6.2 · p. 20

A.6.3 Array-Based Atomic Object.

• Algorithm 6 is not linearizable in the asynchronous model because a read can combine array entries observed before and after concurrent writes, yielding values that never coexisted.Its read scans a circular array sequentially, so no unique linearization point may exist for that result.Source: §A.6.3 · p. 21
• Under phase-based lock-step execution, Algorithm 6 is linearizable, and correct reads return the correct writer’s latest preceding value, never ⊥ after the first write.The bounded array stores p+1 entries; the phase model ensures writes finish before reads in the subsequent phase.Source: §A.6.3 · p. 21
• Algorithm 6 implements a Byzantine-linearizable SWMR register under phase-based lock-step execution.Within a read phase, lock-step threads observe the same array state and produce identical outputs.Source: §A.6.3 · p. 22

Correctness of SnapObject

• SnapObject is Byzantine-linearizable: the linked-list, wait-free circular-array, and Queue-with-snap implementations are linearizable under concurrency, while the array-based implementation requires operations to remain within one phase.Phase-based lock-step execution and ordered semantics supply the properties needed for Byzantine linearizability.Source: Correctness of SnapObject · p. 22

A.6.4 Memory Reclamation for the Linked-List Object.

• Retaining Δ=(p+1)(δ−1)n/p nodes prevents premature reclamation; memory complexity is O(n) when n is much larger than p and δ, and p=32 yields Δ=207n.Garbage collection truncates the list at the bound using CAS and reclaims the detached suffix.Source: §A.6.4 · pp. 22–23

A.6.5 Wait-Free Circular-Array Based Atomic Object.

• WF-Array completes Insert in two shared-memory operations and ReadLatest in W+1; its ticketed ring rejects stale laps, and readers verify entries and select the valid value with the largest sequence number.Insert uses a ticket fetch-and-increment followed by a store; ReadLatest scans the recent ticket window.Source: §A.6.5 · p. 23
• With a correct writer and at least one correct reader, WF-Array is Byzantine-linearizable for f≤n−1, and correct readers return ⊥ with probability 0.The correctness guarantee depends on the writer being correct; the reader population may otherwise be Byzantine.Source: §A.6.5 · p. 24
• Reinserted writer-signed entries retain their original sequence numbers, so replay does not violate the ordering properties under any SnapObject realization.The register’s ordering guarantees hold against the general adversary, including Byzantine processes that replay older entries.Source: §A.6.5 · p. 25

A.7 Idealized Signature-Aware Memory

• The idealized signature-aware memory design assumes shared memory, an atomic register (CAReg), and signatures.Source: §A.7 · p. 25

A.7.2 Algorithmic Solution.

• Its Write atomically updates the stored value and sequence only when the signature verifies and the sequence exceeds the stored one; Read returns the stored value.Rejected writes have no effect.Source: §A.7.2 · p. 25
• The memory object enforces authentication and freshness internally: signature verification blocks forged values, while a monotonic counter check rejects replayed values with stale sequence numbers.Because the sequence number is bound inside the signature, a Byzantine process cannot forge a fresh counter or alter a replayed one.Source: §A.7.2 · p. 25

A.7.3 Correctness for Idealized Signature-Aware Memory.

• Algorithm 8 implements a Byzantine-linearizable single-writer multi-reader register according to Definition 3.The proof establishes object-level linearizability, writer authenticity, and the required ordering properties.Source: §A.7.3 · p. 25

Object-Level Linearizability

• All completed reads and writes on Algorithm 8’s signature-aware memory object are linearizable.A successful write atomically updates the register, a read atomically reads it, and a write with an invalid signature has no effect.Source: Object-Level Linearizability · p. 25

Value Validity

• A non-NULL value returned by a correct read was written by the authorized writer, and Byzantine replayed writes cannot modify the register or prevent its sequence number from increasing.Successful writes require a valid signature and a sequence number greater than the stored sequence number.Source: Value Validity · p. 26

A.7.4 Byzantine Linearizability.

• Algorithm 8 implements a Byzantine-linearizable SWMR register, satisfying the required ordering properties through atomic register semantics.The ordering properties include No Future Read, No Read from Distant-Past, No New–Old Inversion, and Intra-Phase Reader Consistency.Source: §A.7.4 · p. 26

A.7.5 Resilience of the Idealized Signature-Aware Memory.

• With a correct writer and at least one correct reader, the construction tolerates up to n−1 Byzantine readers; correct reads return values written by the writer.The register state is determined exclusively by the writer’s writes, which induce the sequential order.Source: §A.7.5 · p. 26