(
source,
dest,
source_idx,
dest_idx,
base_idx,
start_pos,
middle,
end,
bx,
tx,
step_count,
even,
)
| 381 | ) |
| 382 | |
| 383 | def dual_mergepath( |
| 384 | source, |
| 385 | dest, |
| 386 | source_idx, |
| 387 | dest_idx, |
| 388 | base_idx, |
| 389 | start_pos, |
| 390 | middle, |
| 391 | end, |
| 392 | bx, |
| 393 | tx, |
| 394 | step_count, |
| 395 | even, |
| 396 | ): |
| 397 | outer_first_buf = T.decl_buffer([1], target_dtype, scope="local") |
| 398 | outer_last_buf = T.decl_buffer([1], target_dtype, scope="local") |
| 399 | first_buf = T.decl_buffer([1], target_dtype, scope="local") |
| 400 | last_buf = T.decl_buffer([1], target_dtype, scope="local") |
| 401 | i_buf_buf = T.decl_buffer([1], target_dtype, scope="local") |
| 402 | j_buf_buf = T.decl_buffer([1], target_dtype, scope="local") |
| 403 | outer_first = T.buffer_proxy(outer_first_buf) |
| 404 | outer_last = T.buffer_proxy(outer_last_buf) |
| 405 | first = T.buffer_proxy(first_buf) |
| 406 | last = T.buffer_proxy(last_buf) |
| 407 | i_buf = T.buffer_proxy(i_buf_buf) |
| 408 | j_buf = T.buffer_proxy(j_buf_buf) |
| 409 | |
| 410 | diag = bx * step_count |
| 411 | with T.If(even): |
| 412 | with T.Then(): |
| 413 | get_merge_begin( |
| 414 | source, |
| 415 | base_idx, |
| 416 | middle - start_pos, |
| 417 | end - middle, |
| 418 | start_pos, |
| 419 | middle, |
| 420 | diag, |
| 421 | outer_first, |
| 422 | outer_last, |
| 423 | ) |
| 424 | aStart = start_pos + outer_first[0] |
| 425 | bStart = middle + diag - outer_last[0] |
| 426 | aCount = tvm.te.min(middle - aStart, step_count) |
| 427 | bCount = tvm.te.min(end - bStart, step_count) |
| 428 | inner_diag = tx * thread_work |
| 429 | get_merge_begin( |
| 430 | source, base_idx, aCount, bCount, aStart, bStart, inner_diag, first, last |
| 431 | ) |
| 432 | serial_merge( |
| 433 | source, |
| 434 | dest, |
| 435 | source_idx, |
| 436 | dest_idx, |
| 437 | base_idx, |
| 438 | aCount, |
| 439 | bCount, |
| 440 | aStart, |
no test coverage detected
searching dependent graphs…