Returns the relative dim3 block rank local to the cluster.
| 433 | |
| 434 | // Returns the relative dim3 block rank local to the cluster. |
| 435 | DEVICE dim3 block_id_in_cluster() { |
| 436 | uint32_t x, y, z; |
| 437 | asm volatile("mov.u32 %0, %cluster_ctaid.x;\n" : "=r"(x) :); |
| 438 | asm volatile("mov.u32 %0, %cluster_ctaid.y;\n" : "=r"(y) :); |
| 439 | asm volatile("mov.u32 %0, %cluster_ctaid.z;\n" : "=r"(z) :); |
| 440 | return {x, y, z}; |
| 441 | } |
| 442 | |
| 443 | // Returns the warp id in block |
| 444 | DEVICE int warp_id_in_block() { return threadIdx.x / WARP_SIZE; } |
no outgoing calls
no test coverage detected