(
source,
dest,
source_idx,
dest_idx,
base_idx,
aCount,
bCount,
aStart,
bStart,
kStart,
diag,
step_count,
first,
last,
i_buf,
j_buf,
)
| 255 | last[0] = mid |
| 256 | |
| 257 | def serial_merge( |
| 258 | source, |
| 259 | dest, |
| 260 | source_idx, |
| 261 | dest_idx, |
| 262 | base_idx, |
| 263 | aCount, |
| 264 | bCount, |
| 265 | aStart, |
| 266 | bStart, |
| 267 | kStart, |
| 268 | diag, |
| 269 | step_count, |
| 270 | first, |
| 271 | last, |
| 272 | i_buf, |
| 273 | j_buf, |
| 274 | ): |
| 275 | i_val = aStart + first[0] |
| 276 | j_val = bStart + diag - last[0] |
| 277 | if is_webgpu: |
| 278 | i_buf[0] = cast(i_val, target_dtype) |
| 279 | j_buf[0] = cast(j_val, target_dtype) |
| 280 | else: |
| 281 | i_buf[0] = i_val |
| 282 | j_buf[0] = j_val |
| 283 | |
| 284 | with T.serial(0, tvm.te.min(aCount + bCount - diag, step_count)) as count: |
| 285 | i_idx = base_idx + i_buf[0] |
| 286 | j_idx = base_idx + j_buf[0] |
| 287 | k_idx = base_idx + (kStart + diag + count) |
| 288 | |
| 289 | with T.If(tvm.tirx.all(i_buf[0] < aStart + aCount, j_buf[0] < bStart + bCount)): |
| 290 | with T.Then(): |
| 291 | with T.If(compare(source[i_idx], source[j_idx])): |
| 292 | with T.Then(): |
| 293 | dest[k_idx] = source[i_idx] |
| 294 | if values is not None: |
| 295 | dest_idx[k_idx] = source_idx[i_idx] |
| 296 | i_buf[0] = i_buf[0] + 1 |
| 297 | with T.Else(): |
| 298 | dest[k_idx] = source[j_idx] |
| 299 | if values is not None: |
| 300 | dest_idx[k_idx] = source_idx[j_idx] |
| 301 | j_buf[0] = j_buf[0] + 1 |
| 302 | with T.Else(): |
| 303 | with T.If(i_buf[0] < aStart + aCount): |
| 304 | with T.Then(): |
| 305 | dest[k_idx] = source[i_idx] |
| 306 | if values is not None: |
| 307 | dest_idx[k_idx] = source_idx[i_idx] |
| 308 | i_buf[0] = i_buf[0] + 1 |
| 309 | with T.Else(): |
| 310 | dest[k_idx] = source[j_idx] |
| 311 | if values is not None: |
| 312 | dest_idx[k_idx] = source_idx[j_idx] |
| 313 | j_buf[0] = j_buf[0] + 1 |
| 314 |
no test coverage detected
searching dependent graphs…