| 63 | using jit::BufferNode; |
| 64 | |
| 65 | string getKernelString(const string& funcName, const vector<Node*>& full_nodes, |
| 66 | const vector<Node_ids>& full_ids, |
| 67 | const vector<int>& output_ids, const bool is_linear, |
| 68 | const bool loop0, const bool loop1, const bool loop3) { |
| 69 | // Common OpenCL code |
| 70 | // This part of the code does not change with the kernel. |
| 71 | |
| 72 | static const char* kernelVoid = R"JIT( |
| 73 | __kernel void )JIT"; |
| 74 | static const char* dimParams = "KParam oInfo"; |
| 75 | static const char* blockStart = "{"; |
| 76 | static const char* blockEnd = "\n}\n"; |
| 77 | |
| 78 | static const char* linearInit = R"JIT( |
| 79 | int idx = get_global_id(0); |
| 80 | const int idxEnd = oInfo.dims[0]; |
| 81 | if (idx < idxEnd) { |
| 82 | )JIT"; |
| 83 | static const char* linearEnd = R"JIT( |
| 84 | })JIT"; |
| 85 | |
| 86 | static const char* linearLoop0Start = R"JIT( |
| 87 | const int idxID0Inc = get_global_size(0); |
| 88 | do {)JIT"; |
| 89 | static const char* linearLoop0End = R"JIT( |
| 90 | idx += idxID0Inc; |
| 91 | if (idx >= idxEnd) break; |
| 92 | } while (true);)JIT"; |
| 93 | |
| 94 | // /////////////////////////////////////////////// |
| 95 | // oInfo = output optimized information (dims, strides, offset). |
| 96 | // oInfo has removed dimensions, to optimized block scheduling |
| 97 | // iInfo = input internal information (dims, strides, offset) |
| 98 | // iInfo has the original dimensions, auto generated code |
| 99 | // |
| 100 | // Loop3 is fastest and becomes inside loop, since |
| 101 | // - #of loops is known upfront |
| 102 | // Loop1 is used for extra dynamic looping (writing into cache) |
| 103 | // All loops are conditional and idependent |
| 104 | // Format Loop1 & Loop3 |
| 105 | // //////////////////////////// |
| 106 | // *stridedLoopNInit // Always |
| 107 | // *stridedLoop1Init // Conditional |
| 108 | // *stridedLoop2Init // Conditional |
| 109 | // *stridedLoop3Init // Conditional |
| 110 | // *stridedLoop1Start // Conditional |
| 111 | // *stridedLoop3Start // Conditional |
| 112 | // auto generated code // Always |
| 113 | // *stridedLoop3End // Conditional |
| 114 | // *stridedLoop1End // Conditional |
| 115 | // *StridedEnd // Always |
| 116 | // |
| 117 | // format loop0 (Vector only) |
| 118 | // ////////////////////////// |
| 119 | // *stridedLoop0Init // Always |
| 120 | // *stridedLoop0Start // Always |
| 121 | // auto generated code // Always |
| 122 | // *stridedLoop0End // Always |
no test coverage detected