| 61 | } |
| 62 | |
| 63 | __host__ __device__ |
| 64 | static Index3 fromLinear(Index linear, const Index3& dims, const Index3& stride) |
| 65 | { |
| 66 | //for (int i = 0; i < 3; ++i) outCoords[i] = (linear / stride[i]) % dims[i]; |
| 67 | //manual loop unrolling |
| 68 | Index3 coords = { |
| 69 | (linear / stride.get<0>()) % dims.get<0>(), |
| 70 | (linear / stride.get<1>()) % dims.get<1>(), |
| 71 | (linear / stride.get<2>()) % dims.get<2>() |
| 72 | }; |
| 73 | //printf("index: %d; stride: %d,%d,%d -> coords: %d,%d,%d\n", |
| 74 | // (int)linear, |
| 75 | // (int)stride.get<0>(), (int)stride.get<1>(), (int)stride.get<2>(), |
| 76 | // (int)coords.get<0>(), (int)coords.get<1>(), (int)coords.get<2>()); |
| 77 | return coords; |
| 78 | } |
| 79 | |
| 80 | /// Postfix increment |
| 81 | __host__ __device__ CUMAT_STRONG_INLINE self_type operator++(int) |
no outgoing calls
no test coverage detected