| 401 | } |
| 402 | |
| 403 | DEVICE uint32_t elect_one_sync() { |
| 404 | uint32_t pred = 0; |
| 405 | uint32_t laneid = 0; |
| 406 | asm volatile( |
| 407 | "{\n" |
| 408 | ".reg .b32 %rx;\n" |
| 409 | ".reg .pred %px;\n" |
| 410 | " elect.sync %rx|%px, %2;\n" |
| 411 | "@%px mov.s32 %1, 1;\n" |
| 412 | " mov.s32 %0, %rx;\n" |
| 413 | "}\n" |
| 414 | : "+r"(laneid), "+r"(pred) |
| 415 | : "r"(0xFFFFFFFF)); |
| 416 | return pred; |
| 417 | } |
| 418 | |
| 419 | DEVICE uint32_t block_rank_in_cluster() { |
| 420 | uint32_t rank; |
no outgoing calls
no test coverage detected