Search across packages, versions, files, declarations, and comments. Use multiple terms to narrow results (AND), and add @vertex, @fragment, or @compute to filter by shader stage.
1// Common routines for timestamp normalization.
2//
3// This is split out into its own file so that the tests in `tests` can include
4// it without including the normal endpoints and interface definitions.
6/// 64-bit unsigned integer type.
7///
8/// We cannot rely on native 64-bit integers, so we define our own 64-bit
9/// integer type as two 32-bit integers.
10struct Uint64 {
11 /// Least significant word.
12 low: u32,
13 /// Most significant word.
14 high: u32,
17/// 96-bit unsigned integer type.
18struct Uint96 {
19 /// Least significant word.
20 low: u32,
21 /// Middle word.
22 mid: u32,
23 /// Most significant word.
24 high: u32,
27/// Truncates a 96-bit number to a 64-bit number by discarding the upper 32 bits.
28fn truncate_u96_to_u64(x: Uint96) -> Uint64 {
29 return Uint64(
30 x.low,
31 x.mid,
32 );
35/// Returns the lower 16 bits of a 32-bit integer.
36fn low(a: u32) -> u32 {
37 return a & 0xFFFF;
40/// Returns the upper 16 bits of a 32-bit integer.
41fn high(a: u32) -> u32 {
42 return a >> 16;
45/// Combines two 16bit words into a single 32bit word.
46/// `w1` is the upper 16 bits and `w0` is the lower 16 bits.
47///
48/// The high 16 bits of each argument are discarded.
49fn u32_from_u16s(w1: u32, w0: u32) -> u32 {
50 return low(w1) << 16 | low(w0);
53// Multiplies a 64-bit number by a 32-bit number and outputs a 96-bit result.
54//
55// The number of digits (bits) needed to represent the result of a multiplication
56// is the sum of the number of input digits (bits). Since we are multiplying a
57// 64-bit number by a 32-bit number, we need 96 bits to represent the result.
58fn u64_mul_u32(a: Uint64, b: u32) -> Uint96 {
59 // Does not use any 64-bit operations and we don't have access to `mul(u32, u32) -> u64`
60 // operations, so we operate entirely on `mul(u16, u16) -> u32`.
62 // This implements standard "long multiplication" algorithm using 16-bit words.
63 // Each element in this diagram is a 16-bit word.
64 //
65 // a3 a2 a1 a0
66 // * b1 b0
67 // ----------------------------
68 // i0 = p00 p00
69 // i1 = p10 p10
70 // i2 = p20 p20
71 // i3 = p30 p30
72 // i4 = p01 p01
73 // i5 = p11 p11
74 // i6 = p21 p21
75 // i7 = p31 p31
76 // ----------------------------
77 // r6 r5 r4 r3 r2 r1 r0
79 // Decompose the 64-bit number into four 16-bit words.
80 let a0 = low(a.low);
81 let a1 = high(a.low);
82 let a2 = low(a.high);
83 let a3 = high(a.high);
85 // Decompose the 32-bit number into two 16-bit words.
86 let b0 = low(b);
87 let b1 = high(b);
89 // Each line represents one row in the diagram above.
90 let i0 = a0 * b0;
91 let i1 = a1 * b0;
92 let i2 = a2 * b0;
93 let i3 = a3 * b0;
94 let i4 = a0 * b1;
95 let i5 = a1 * b1;
96 let i6 = a2 * b1;
97 let i7 = a3 * b1;
99 // Each line represents one column in the diagram above.
100 //
101 // The high 16 bits of each column are the carry to the next column.
102 let r0 = low(i0);
103 let r1 = high(i0) + low(i1) + low(i4) + high(r0);
104 let r2 = high(i1) + low(i2) + high(i4) + low(i5) + high(r1);
105 let r3 = high(i2) + low(i3) + high(i5) + low(i6) + high(r2);
106 let r4 = high(i3) + high(i6) + low(i7) + high(r3);
107 let r5 = high(i7) + high(r4);
108 // The r5 carry will always be zero.
110 let out0 = u32_from_u16s(r1, r0);
111 let out1 = u32_from_u16s(r3, r2);
112 let out2 = u32_from_u16s(r5, r4);
114 return Uint96(out0, out1, out2);
117// Shifts a 96-bit number right by a given number of bits.
118//
119// The shift is in the range [0, 32].
120fn shift_right_96(x: Uint96, shift: u32) -> Uint96 {
121 // Shift wraps around at 32, which breaks the algorithm when
122 // either shift is 32 or inv_shift is 32.
123 if (shift == 0) {
124 return x;
125 }
126 if (shift == 32) {
127 return Uint96(x.mid, x.high, 0);
128 }
130 let inv_shift = 32 - shift;
132 let carry2 = x.high << inv_shift;
133 let carry1 = x.mid << inv_shift;
135 var out: Uint96;
137 out.high = x.high >> shift;
138 out.mid = (x.mid >> shift) | carry2;
139 out.low = (x.low >> shift) | carry1;
141 return out;
Contents