GCC Code Coverage Report

Directory: brotensor/
Coverage: low: ≥ 0% medium: ≥ 75.0% high: ≥ 90.0%
Coverage Exec / Excl / Total
Lines: 33.6% 75 / 0 / 223
Functions: 78.6% 11 / 0 / 14
Branches: 9.0% 80 / 0 / 888
src/jit/ptx_emit_detail.h
Line Branch Exec Source
1 #pragma once
2
3 // Shared PTX-emission helpers for the trace JIT's two CUDA emitters.
4 //
5 // One rule governs every function here: an emitter helper may write to the
6 // stream, so its result must be bound to a local before it is used inside
7 // another stream expression. `o << "add.u64 " << d << ", " << p << ", " <<
8 // offset(dt)` interleaves the two instructions and produces PTX that does not
9 // parse.
10
11 #include "ptx_emit.h"
12
13 #include <cstdint>
14 #include <cstring>
15 #include <iomanip>
16 #include <ostream>
17 #include <sstream>
18 #include <stdexcept>
19 #include <string>
20 #include <unordered_map>
21 #include <vector>
22
23 namespace brotensor::jit::ptx::detail {
24
25 44 inline std::string fhex(float f) {
26 44 std::uint32_t bits = 0;
27 44 std::memcpy(&bits, &f, sizeof(float));
28 44 std::stringstream ss;
29
8/16
✓ Branch 0 taken 44 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 44 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 44 times.
✗ Branch 5 not taken.
✓ Branch 6 taken 44 times.
✗ Branch 7 not taken.
✓ Branch 8 taken 44 times.
✗ Branch 9 not taken.
✓ Branch 10 taken 44 times.
✗ Branch 11 not taken.
✓ Branch 12 taken 44 times.
✗ Branch 13 not taken.
✓ Branch 14 taken 44 times.
✗ Branch 15 not taken.
 44 ss << "0f" << std::uppercase << std::hex << std::setfill('0') << std::setw(8) << bits;
30
1/2
✓ Branch 0 taken 44 times.
✗ Branch 1 not taken.
 44 return ss.str();
31 44 }
32
33 // Virtual-register allocator. PTX register banks are declared up front, so a
34 // body is emitted into its own stream and `declarations()` is prepended once
35 // the high-water marks are known.
36 16 struct Alloc {
37 16 int f = 0, r = 0, rd = 0, h = 0, p = 0;
38
1/2
✓ Branch 0 taken 672 times.
✗ Branch 1 not taken.
 672 std::string F() { return "%f" + std::to_string(f++); }
39
1/2
✓ Branch 0 taken 546 times.
✗ Branch 1 not taken.
 546 std::string R() { return "%r" + std::to_string(r++); }
40
1/2
✓ Branch 0 taken 284 times.
✗ Branch 1 not taken.
 284 std::string RD() { return "%rd" + std::to_string(rd++); }
41 ✗ std::string H() { return "%h" + std::to_string(h++); }
42
1/2
✓ Branch 0 taken 81 times.
✗ Branch 1 not taken.
 81 std::string P() { return "%p" + std::to_string(p++); }
43 };
44
45 16 inline std::string declarations(const Alloc& a) {
46 16 std::stringstream ss;
47
4/8
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 16 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 16 times.
✗ Branch 5 not taken.
✓ Branch 6 taken 16 times.
✗ Branch 7 not taken.
 16 if (a.p) ss << " .reg .pred %p<" << a.p << ">;\n";
48
4/8
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 16 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 16 times.
✗ Branch 5 not taken.
✓ Branch 6 taken 16 times.
✗ Branch 7 not taken.
 16 if (a.r) ss << " .reg .b32 %r<" << a.r << ">;\n";
49
4/8
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 16 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 16 times.
✗ Branch 5 not taken.
✓ Branch 6 taken 16 times.
✗ Branch 7 not taken.
 16 if (a.rd) ss << " .reg .b64 %rd<" << a.rd << ">;\n";
50
1/8
✗ Branch 0 not taken.
✓ Branch 1 taken 16 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
 16 if (a.h) ss << " .reg .b16 %h<" << a.h << ">;\n";
51
4/8
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 16 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 16 times.
✗ Branch 5 not taken.
✓ Branch 6 taken 16 times.
✗ Branch 7 not taken.
 16 if (a.f) ss << " .reg .f32 %f<" << a.f << ">;\n";
52
1/2
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
 16 return ss.str();
53 16 }
54
55 ✗ inline bool is_pow2(int v) { return v > 0 && (v & (v - 1)) == 0; }
56
57 // ── one lane of the op chain ────────────────────────────────────────────────
58
59 ✗ inline void emit_sigmoid(std::ostream& o, Alloc& a, const std::string& x,
60 const std::string& d) {
61 ✗ const std::string t0 = a.F(), t1 = a.F(), t2 = a.F();
62 ✗ o << " neg.f32 " << t0 << ", " << x << ";\n";
63 ✗ o << " mul.f32 " << t0 << ", " << t0 << ", 0f3FB8AA3B;\n"; // * log2(e)
64 ✗ o << " ex2.approx.f32 " << t1 << ", " << t0 << ";\n";
65 ✗ o << " add.f32 " << t2 << ", " << t1 << ", 0f3F800000;\n";
66 ✗ o << " rcp.approx.f32 " << d << ", " << t2 << ";\n";
67 ✗ }
68
69 40 inline void emit_lane_op(std::ostream& o, Alloc& a, const TraceNode& n,
70 const std::vector<std::string>& in, const std::string& d) {
71
1/22
✓ Branch 0 taken 40 times.
✗ Branch 1 not taken.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
✗ Branch 10 not taken.
✗ Branch 11 not taken.
✗ Branch 12 not taken.
✗ Branch 13 not taken.
✗ Branch 14 not taken.
✗ Branch 15 not taken.
✗ Branch 16 not taken.
✗ Branch 17 not taken.
✗ Branch 18 not taken.
✗ Branch 19 not taken.
✗ Branch 20 not taken.
✗ Branch 21 not taken.
 40 switch (n.op) {
72 case TraceOpKind::Add:
73 40 o << " add.f32 " << d << ", " << in[0] << ", " << in[1] << ";\n";
74 40 break;
75 case TraceOpKind::Sub:
76 ✗ o << " sub.f32 " << d << ", " << in[0] << ", " << in[1] << ";\n";
77 ✗ break;
78 case TraceOpKind::Mul:
79 ✗ o << " mul.f32 " << d << ", " << in[0] << ", " << in[1] << ";\n";
80 ✗ break;
81 case TraceOpKind::Div:
82 ✗ o << " div.rn.f32 " << d << ", " << in[0] << ", " << in[1] << ";\n";
83 ✗ break;
84 case TraceOpKind::AddScalar:
85 ✗ o << " add.f32 " << d << ", " << in[0] << ", " << fhex(n.scalar) << ";\n";
86 ✗ break;
87 case TraceOpKind::SubScalar:
88 ✗ o << " sub.f32 " << d << ", " << in[0] << ", " << fhex(n.scalar) << ";\n";
89 ✗ break;
90 case TraceOpKind::MulScalar:
91 ✗ o << " mul.f32 " << d << ", " << in[0] << ", " << fhex(n.scalar) << ";\n";
92 ✗ break;
93 case TraceOpKind::DivScalar:
94 ✗ o << " div.rn.f32 " << d << ", " << in[0] << ", " << fhex(n.scalar) << ";\n";
95 ✗ break;
96 case TraceOpKind::ScalarSub:
97 ✗ o << " sub.f32 " << d << ", " << fhex(n.scalar) << ", " << in[0] << ";\n";
98 ✗ break;
99 case TraceOpKind::ScalarDiv:
100 ✗ o << " div.rn.f32 " << d << ", " << fhex(n.scalar) << ", " << in[0] << ";\n";
101 ✗ break;
102 case TraceOpKind::FMA:
103 ✗ o << " fma.rn.f32 " << d << ", " << in[0] << ", " << in[1] << ", " << in[2]
104 ✗ << ";\n";
105 ✗ break;
106 case TraceOpKind::ReLU:
107 ✗ o << " max.f32 " << d << ", " << in[0] << ", 0f00000000;\n";
108 ✗ break;
109 case TraceOpKind::Clamp: {
110 ✗ const std::string t = a.F();
111 ✗ o << " max.f32 " << t << ", " << in[0] << ", " << fhex(n.scalar) << ";\n";
112 ✗ o << " min.f32 " << d << ", " << t << ", " << fhex(n.scalar2) << ";\n";
113 break;
114 ✗ }
115 case TraceOpKind::Where: {
116 ✗ const std::string p = a.P();
117 ✗ o << " setp.ne.f32 " << p << ", " << in[0] << ", 0f00000000;\n";
118 ✗ o << " selp.f32 " << d << ", " << in[1] << ", " << in[2] << ", " << p << ";\n";
119 break;
120 ✗ }
121 case TraceOpKind::Cast: {
122 ✗ if (n.dtype == Dtype::FP16) {
123 ✗ const std::string h = a.H();
124 ✗ o << " cvt.rn.f16.f32 " << h << ", " << in[0] << ";\n";
125 ✗ o << " cvt.f32.f16 " << d << ", " << h << ";\n";
126 ✗ } else if (n.dtype == Dtype::BF16) {
127 ✗ const std::string h = a.H(), t = a.R();
128 ✗ o << " cvt.rn.bf16.f32 " << h << ", " << in[0] << ";\n";
129 ✗ o << " cvt.u32.u16 " << t << ", " << h << ";\n";
130 ✗ o << " shl.b32 " << t << ", " << t << ", 16;\n";
131 ✗ o << " mov.b32 " << d << ", " << t << ";\n";
132 ✗ } else {
133 ✗ o << " mov.f32 " << d << ", " << in[0] << ";\n";
134 }
135 ✗ break;
136 }
137 case TraceOpKind::Sigmoid:
138 ✗ emit_sigmoid(o, a, in[0], d);
139 ✗ break;
140 case TraceOpKind::SiLU: {
141 ✗ const std::string s = a.F();
142 ✗ emit_sigmoid(o, a, in[0], s);
143 ✗ o << " mul.f32 " << d << ", " << in[0] << ", " << s << ";\n";
144 break;
145 ✗ }
146 case TraceOpKind::Tanh:
147 ✗ o << " tanh.approx.f32 " << d << ", " << in[0] << ";\n";
148 ✗ break;
149 case TraceOpKind::Copy:
150 // The value already sits in a register; the node exists so the
151 // store below lands in the caller's buffer.
152 ✗ o << " mov.f32 " << d << ", " << in[0] << ";\n";
153 ✗ break;
154 case TraceOpKind::GELU: {
155 // 0.5 x (1 + tanh(sqrt(2/pi) (x + 0.044715 x^3))) — the tanh
156 // approximation nn.GELU(approximate="tanh") uses.
157 ✗ const std::string x2 = a.F(), inner = a.F(), poly = a.F();
158 ✗ const std::string arg = a.F(), th = a.F(), t1 = a.F(), hx = a.F();
159 ✗ o << " mul.f32 " << x2 << ", " << in[0] << ", " << in[0] << ";\n";
160 ✗ o << " fma.rn.f32 " << inner << ", " << x2 << ", " << fhex(0.044715f) << ", "
161 ✗ << fhex(1.0f) << ";\n";
162 ✗ o << " mul.f32 " << poly << ", " << inner << ", " << in[0] << ";\n";
163 ✗ o << " mul.f32 " << arg << ", " << poly << ", " << fhex(0.7978845608f) << ";\n";
164 ✗ o << " tanh.approx.f32 " << th << ", " << arg << ";\n";
165 ✗ o << " add.f32 " << t1 << ", " << th << ", " << fhex(1.0f) << ";\n";
166 ✗ o << " mul.f32 " << hx << ", " << in[0] << ", " << fhex(0.5f) << ";\n";
167 ✗ o << " mul.f32 " << d << ", " << hx << ", " << t1 << ";\n";
168 break;
169 ✗ }
170 case TraceOpKind::Modulate: {
171 // x * (1 + scale) + shift, with scale/shift normally (1, D) rows
172 // the loader has already resolved to this lane's column.
173 ✗ const std::string s1 = a.F();
174 ✗ o << " add.f32 " << s1 << ", " << in[1] << ", " << fhex(1.0f) << ";\n";
175 ✗ o << " fma.rn.f32 " << d << ", " << in[0] << ", " << s1 << ", " << in[2] << ";\n";
176 break;
177 ✗ }
178 default:
179 ✗ throw std::runtime_error(std::string("brotensor::jit: op '") + op_kind_name(n.op) +
180 "' has no elementwise PTX form");
181 }
182 40 }
183
184 // ── typed load / store ──────────────────────────────────────────────────────
185
186 // Loads `lanes` consecutive elements at [addr] into fresh FP32 registers.
187 // `lanes` is 1, 2, 4 or 8 and the access is one instruction.
188 88 inline std::vector<std::string> emit_load(std::ostream& o, Alloc& a, Dtype dt, int lanes,
189 const std::string& addr) {
190 88 std::vector<std::string> out(static_cast<std::size_t>(lanes));
191
3/4
✓ Branch 0 taken 220 times.
✓ Branch 1 taken 88 times.
✓ Branch 2 taken 220 times.
✗ Branch 3 not taken.
 308 for (int i = 0; i < lanes; ++i) out[static_cast<std::size_t>(i)] = a.F();
192
193
1/2
✗ Branch 0 not taken.
✓ Branch 1 taken 88 times.
 88 if (dt == Dtype::FP32) {
194
2/2
✓ Branch 0 taken 44 times.
✓ Branch 1 taken 44 times.
 88 if (lanes == 1) {
195
5/10
✓ Branch 0 taken 44 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 44 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 44 times.
✗ Branch 5 not taken.
✓ Branch 6 taken 44 times.
✗ Branch 7 not taken.
✓ Branch 8 taken 44 times.
✗ Branch 9 not taken.
 44 o << " ld.global.f32 " << out[0] << ", [" << addr << "];\n";
196 44 } else {
197
3/6
✓ Branch 0 taken 44 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 44 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 44 times.
✗ Branch 5 not taken.
 44 o << " ld.global.v" << lanes << ".f32 {";
198
2/2
✓ Branch 0 taken 44 times.
✓ Branch 1 taken 176 times.
 220 for (int i = 0; i < lanes; ++i)
199
2/4
✓ Branch 0 taken 176 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 176 times.
✗ Branch 3 not taken.
 176 o << (i ? ", " : "") << out[static_cast<std::size_t>(i)];
200
3/6
✓ Branch 0 taken 44 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 44 times.
✗ Branch 3 not taken.
✓ Branch 4 taken 44 times.
✗ Branch 5 not taken.
 44 o << "}, [" << addr << "];\n";
201 }
202 88 return out;
203 }
204
205 // BF16 / FP16: move 16-bit words as packed b32 and widen in registers.
206 ✗ if (lanes == 1) {
207 ✗ const std::string h = a.H();
208 ✗ o << " ld.global.b16 " << h << ", [" << addr << "];\n";
209 ✗ if (dt == Dtype::BF16) {
210 ✗ const std::string t = a.R();
211 ✗ o << " cvt.u32.u16 " << t << ", " << h << ";\n";
212 ✗ o << " shl.b32 " << t << ", " << t << ", 16;\n";
213 ✗ o << " mov.b32 " << out[0] << ", " << t << ";\n";
214 ✗ } else {
215 ✗ o << " cvt.f32.f16 " << out[0] << ", " << h << ";\n";
216 }
217 ✗ return out;
218 ✗ }
219
220 ✗ const int words = lanes / 2;
221 ✗ std::vector<std::string> w(static_cast<std::size_t>(words));
222 ✗ for (int i = 0; i < words; ++i) w[static_cast<std::size_t>(i)] = a.R();
223 ✗ if (words == 1) {
224 ✗ o << " ld.global.u32 " << w[0] << ", [" << addr << "];\n";
225 ✗ } else {
226 ✗ o << " ld.global.v" << words << ".u32 {";
227 ✗ for (int i = 0; i < words; ++i) o << (i ? ", " : "") << w[static_cast<std::size_t>(i)];
228 ✗ o << "}, [" << addr << "];\n";
229 }
230 ✗ for (int i = 0; i < words; ++i) {
231 ✗ const std::string& pk = w[static_cast<std::size_t>(i)];
232 ✗ const std::string& lo = out[static_cast<std::size_t>(2 * i)];
233 ✗ const std::string& hi = out[static_cast<std::size_t>(2 * i + 1)];
234 ✗ if (dt == Dtype::BF16) {
235 ✗ const std::string t = a.R();
236 ✗ o << " shl.b32 " << t << ", " << pk << ", 16;\n";
237 ✗ o << " mov.b32 " << lo << ", " << t << ";\n";
238 ✗ const std::string t2 = a.R();
239 ✗ o << " and.b32 " << t2 << ", " << pk << ", -65536;\n";
240 ✗ o << " mov.b32 " << hi << ", " << t2 << ";\n";
241 ✗ } else {
242 ✗ const std::string h0 = a.H(), h1 = a.H();
243 ✗ o << " mov.b32 {" << h0 << ", " << h1 << "}, " << pk << ";\n";
244 ✗ o << " cvt.f32.f16 " << lo << ", " << h0 << ";\n";
245 ✗ o << " cvt.f32.f16 " << hi << ", " << h1 << ";\n";
246 ✗ }
247 ✗ }
248 ✗ return out;
249 88 }
250
251 20 inline void emit_store(std::ostream& o, Alloc& a, Dtype dt, int lanes,
252 const std::string& addr, const std::vector<std::string>& vals) {
253
1/2
✗ Branch 0 not taken.
✓ Branch 1 taken 20 times.
 20 if (dt == Dtype::FP32) {
254
2/2
✓ Branch 0 taken 10 times.
✓ Branch 1 taken 10 times.
 20 if (lanes == 1) {
255 10 o << " st.global.f32 [" << addr << "], " << vals[0] << ";\n";
256 10 } else {
257 10 o << " st.global.v" << lanes << ".f32 [" << addr << "], {";
258
2/2
✓ Branch 0 taken 40 times.
✓ Branch 1 taken 10 times.
 50 for (int i = 0; i < lanes; ++i)
259 40 o << (i ? ", " : "") << vals[static_cast<std::size_t>(i)];
260 10 o << "};\n";
261 }
262 20 return;
263 }
264
265 ✗ const char* cvt = (dt == Dtype::BF16) ? "cvt.rn.bf16.f32" : "cvt.rn.f16.f32";
266 ✗ if (lanes == 1) {
267 ✗ const std::string h = a.H();
268 ✗ o << " " << cvt << " " << h << ", " << vals[0] << ";\n";
269 ✗ o << " st.global.b16 [" << addr << "], " << h << ";\n";
270 return;
271 ✗ }
272
273 ✗ const int words = lanes / 2;
274 ✗ std::vector<std::string> w(static_cast<std::size_t>(words));
275 ✗ for (int i = 0; i < words; ++i) {
276 ✗ const std::string h0 = a.H(), h1 = a.H();
277 ✗ o << " " << cvt << " " << h0 << ", " << vals[static_cast<std::size_t>(2 * i)] << ";\n";
278 ✗ o << " " << cvt << " " << h1 << ", " << vals[static_cast<std::size_t>(2 * i + 1)]
279 ✗ << ";\n";
280 ✗ w[static_cast<std::size_t>(i)] = a.R();
281 ✗ o << " mov.b32 " << w[static_cast<std::size_t>(i)] << ", {" << h0 << ", " << h1
282 ✗ << "};\n";
283 ✗ }
284 ✗ if (words == 1) {
285 ✗ o << " st.global.u32 [" << addr << "], " << w[0] << ";\n";
286 ✗ } else {
287 ✗ o << " st.global.v" << words << ".u32 [" << addr << "], {";
288 ✗ for (int i = 0; i < words; ++i) o << (i ? ", " : "") << w[static_cast<std::size_t>(i)];
289 ✗ o << "};\n";
290 }
291 20 }
292
293 // ── the op chain ────────────────────────────────────────────────────────────
294
295 using ValueMap = std::unordered_map<int, std::vector<std::string>>;
296
297 // Evaluates every node in `order` (default: the whole DAG, in creation order,
298 // which is topological) that does not already have a value. Nodes seeded by
299 // the caller — a loaded input, or the row-norm emitter's normalised value —
300 // are left alone.
301 52 inline ValueMap emit_chain(std::ostream& o, Alloc& a, const TraceDAG& dag, int lanes,
302 ValueMap vals, const std::vector<int>* order = nullptr) {
303 52 std::vector<int> ids;
304
1/2
✓ Branch 0 taken 52 times.
✗ Branch 1 not taken.
 52 if (order) {
305
1/2
✓ Branch 0 taken 52 times.
✗ Branch 1 not taken.
 52 ids = *order;
306 52 } else {
307 ✗ ids.reserve(dag.node_count());
308 ✗ for (const auto& n : dag.nodes()) ids.push_back(n.id);
309 }
310
2/2
✓ Branch 0 taken 16 times.
✓ Branch 1 taken 52 times.
 68 for (int id : ids) {
311
1/2
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
 16 const TraceNode& n = dag.node(id);
312
1/2
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
 16 if (n.op == TraceOpKind::Input) continue;
313
2/4
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 16 times.
✗ Branch 3 not taken.
 16 if (vals.count(n.id)) continue;
314
315
1/2
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
 16 std::vector<std::string> dst(static_cast<std::size_t>(lanes));
316
2/2
✓ Branch 0 taken 40 times.
✓ Branch 1 taken 16 times.
 56 for (int j = 0; j < lanes; ++j) {
317 40 std::vector<std::string> in;
318
1/2
✓ Branch 0 taken 40 times.
✗ Branch 1 not taken.
 40 in.reserve(n.inputs.size());
319
2/2
✓ Branch 0 taken 40 times.
✓ Branch 1 taken 80 times.
 120 for (int src : n.inputs) {
320
1/2
✓ Branch 0 taken 80 times.
✗ Branch 1 not taken.
 80 auto it = vals.find(src);
321
2/4
✓ Branch 0 taken 80 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 80 times.
✗ Branch 3 not taken.
 80 if (it == vals.end()) {
322 ✗ throw std::runtime_error("brotensor::jit: trace node " +
323 ✗ std::to_string(n.id) +
324 " reads an operand that was never materialised");
325 }
326
2/4
✓ Branch 0 taken 80 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 80 times.
✗ Branch 3 not taken.
 80 in.push_back(it->second[static_cast<std::size_t>(j)]);
327 }
328
1/2
✓ Branch 0 taken 40 times.
✗ Branch 1 not taken.
 40 dst[static_cast<std::size_t>(j)] = a.F();
329
1/2
✓ Branch 0 taken 40 times.
✗ Branch 1 not taken.
 40 emit_lane_op(o, a, n, in, dst[static_cast<std::size_t>(j)]);
330 40 }
331
1/2
✓ Branch 0 taken 16 times.
✗ Branch 1 not taken.
 16 vals[n.id] = std::move(dst);
332 16 }
333 52 return vals;
334 52 }
335
336 } // namespace brotensor::jit::ptx::detail
337