| 22 | }; |
| 23 | |
| 24 | DEVICE TileScheduler(Params params) { |
| 25 | static_assert(ClusterZ == 1 && "Cluster Z should be 1."); |
| 26 | current_idx = blockIdx.x + blockIdx.y * gridDim.x; |
| 27 | problem_m = ceil_div(params.M, BlockX); |
| 28 | problem_n = ceil_div(params.N, BlockY); |
| 29 | problem_m = ceil_div(problem_m, ClusterX) * ClusterX; |
| 30 | problem_n = ceil_div(problem_n, ClusterY) * ClusterY; |
| 31 | total_problem_mn = problem_m * problem_n; |
| 32 | } |
| 33 | |
| 34 | DEVICE TileInfo get_current_tile() { |
| 35 | return get_current_tile_for_idx(current_idx); |
nothing calls this directly
no test coverage detected