| 91 | )JIT"); |
| 92 | |
| 93 | string getKernelString(const string& funcName, |
| 94 | const nonstd::span<Node* const> full_nodes, |
| 95 | nonstd::span<const Node_ids> full_ids, |
| 96 | const nonstd::span<int const> output_ids, |
| 97 | const bool is_linear, const bool loop0, const bool loop1, |
| 98 | const bool loop3) { |
| 99 | // Common OpenCL code |
| 100 | // This part of the code does not change with the kernel. |
| 101 | |
| 102 | static const char* kernelVoid = R"JIT( |
| 103 | __kernel void )JIT"; |
| 104 | static const char* dimParams = "KParam oInfo"; |
| 105 | static const char* blockStart = "{"; |
| 106 | static const char* blockEnd = "\n}\n"; |
| 107 | |
| 108 | static const char* linearInit = R"JIT( |
| 109 | int idx = get_global_id(0); |
| 110 | const int idxEnd = oInfo.dims[0]; |
| 111 | if (idx < idxEnd) { |
| 112 | )JIT"; |
| 113 | static const char* linearEnd = R"JIT( |
| 114 | })JIT"; |
| 115 | |
| 116 | static const char* linearLoop0Start = R"JIT( |
| 117 | const int idxID0Inc = get_global_size(0); |
| 118 | do {)JIT"; |
| 119 | static const char* linearLoop0End = R"JIT( |
| 120 | idx += idxID0Inc; |
| 121 | if (idx >= idxEnd) break; |
| 122 | } while (true);)JIT"; |
| 123 | |
| 124 | // /////////////////////////////////////////////// |
| 125 | // oInfo = output optimized information (dims, strides, offset). |
| 126 | // oInfo has removed dimensions, to optimized block scheduling |
| 127 | // iInfo = input internal information (dims, strides, offset) |
| 128 | // iInfo has the original dimensions, auto generated code |
| 129 | // |
| 130 | // Loop3 is fastest and becomes inside loop, since |
| 131 | // - #of loops is known upfront |
| 132 | // Loop1 is used for extra dynamic looping (writing into cache) |
| 133 | // All loops are conditional and idependent |
| 134 | // Format Loop1 & Loop3 |
| 135 | // //////////////////////////// |
| 136 | // *stridedLoopNInit // Always |
| 137 | // *stridedLoop1Init // Conditional |
| 138 | // *stridedLoop2Init // Conditional |
| 139 | // *stridedLoop3Init // Conditional |
| 140 | // *stridedLoop1Start // Conditional |
| 141 | // *stridedLoop3Start // Conditional |
| 142 | // auto generated code // Always |
| 143 | // *stridedLoop3End // Conditional |
| 144 | // *stridedLoop1End // Conditional |
| 145 | // *StridedEnd // Always |
| 146 | // |
| 147 | // format loop0 (Vector only) |
| 148 | // ////////////////////////// |
| 149 | // *stridedLoop0Init // Always |
| 150 | // *stridedLoop0Start // Always |
no test coverage detected