| 65 | namespace cuda { |
| 66 | |
| 67 | static string getKernelString(const string& funcName, |
| 68 | const vector<Node*>& full_nodes, |
| 69 | const vector<Node_ids>& full_ids, |
| 70 | const vector<int>& output_ids, |
| 71 | const bool is_linear, const bool loop0, |
| 72 | const bool loop1, const bool loop2, |
| 73 | const bool loop3) { |
| 74 | const std::string includeFileStr(jit_cuh, jit_cuh_len); |
| 75 | |
| 76 | const std::string paramTStr = R"JIT( |
| 77 | template<typename T> |
| 78 | struct Param { |
| 79 | dim_t dims[4]; |
| 80 | dim_t strides[4]; |
| 81 | T *ptr; |
| 82 | }; |
| 83 | )JIT"; |
| 84 | |
| 85 | std::string typedefStr{"typedef unsigned int uint;\ntypedef "}; |
| 86 | typedefStr += getFullName<dim_t>(); |
| 87 | typedefStr += " dim_t;\n"; |
| 88 | |
| 89 | // Common CUDA code |
| 90 | // This part of the code does not change with the kernel. |
| 91 | |
| 92 | static const char* kernelVoid = "extern \"C\" __global__ void\n"; |
| 93 | static const char* dimParams = ""; |
| 94 | |
| 95 | static const char* blockStart = "{"; |
| 96 | static const char* blockEnd = "\n}\n"; |
| 97 | |
| 98 | static const char* linearInit = R"JIT( |
| 99 | int idx = blockIdx.x * blockDim.x + threadIdx.x; |
| 100 | const int idxEnd = outref.dims[0]; |
| 101 | if (idx < idxEnd) {)JIT"; |
| 102 | static const char* linearEnd = R"JIT( |
| 103 | })JIT"; |
| 104 | |
| 105 | static const char* linearLoop0Start = R"JIT( |
| 106 | const int idxID0Inc = gridDim.x*blockDim.x; |
| 107 | do {)JIT"; |
| 108 | static const char* linearLoop0End = R"JIT( |
| 109 | idx += idxID0Inc; |
| 110 | if (idx >= idxEnd) break; |
| 111 | } while (true);)JIT"; |
| 112 | |
| 113 | // /////////////////////////////////////////////// |
| 114 | // oInfo = output optimized information (dims, strides, offset). |
| 115 | // oInfo has removed dimensions, to optimized block scheduling |
| 116 | // iInfo = input internal information (dims, strides, offset) |
| 117 | // iInfo has the original dimensions, auto generated code |
| 118 | // |
| 119 | // Loop3 is fastest and becomes inside loop, since |
| 120 | // - #of loops is known upfront |
| 121 | // Loop1 is used for extra dynamic looping (writing into cache) |
| 122 | // Loop0 is used for extra dynamic looping (writing into cache), |
| 123 | // VECTORS ONLY!! |
| 124 | // All loops are conditional and idependent Format Loop1 & Loop3 |
no test coverage detected