| 143 | } |
| 144 | |
| 145 | void GenerateKernel(std::string &str) |
| 146 | { |
| 147 | std::string rType = RegBaseType<PR>(1); |
| 148 | std::string r2Type = RegBaseType<PR>(2); |
| 149 | |
| 150 | bool inIlvd; // Input is interleaved format |
| 151 | bool outIlvd; // Output is interleaved format |
| 152 | inIlvd = ( (params.fft_inputLayout == CLFFT_COMPLEX_INTERLEAVED) || |
| 153 | (params.fft_inputLayout == CLFFT_HERMITIAN_INTERLEAVED) ) ? true : false; |
| 154 | outIlvd = ( (params.fft_outputLayout == CLFFT_COMPLEX_INTERLEAVED) || |
| 155 | (params.fft_outputLayout == CLFFT_HERMITIAN_INTERLEAVED) ) ? true : false; |
| 156 | |
| 157 | |
| 158 | |
| 159 | // Pragma |
| 160 | str += ClPragma<PR>(); |
| 161 | |
| 162 | std::string sfx = FloatSuffix<PR>(); |
| 163 | |
| 164 | //If pre-callback is set for the plan |
| 165 | if (params.fft_hasPreCallback && h2c) |
| 166 | { |
| 167 | //Insert callback function code at the beginning |
| 168 | str += params.fft_preCallback.funcstring; |
| 169 | str += "\n\n"; |
| 170 | } |
| 171 | |
| 172 | //if postcallback is set |
| 173 | if (params.fft_hasPostCallback) |
| 174 | { |
| 175 | //Insert callback function code at the beginning |
| 176 | str += params.fft_postCallback.funcstring; |
| 177 | str += "\n\n"; |
| 178 | } |
| 179 | |
| 180 | // Copy kernel begin |
| 181 | str += "__kernel void "; |
| 182 | |
| 183 | // Function name |
| 184 | if(general) |
| 185 | str += "copy_general"; |
| 186 | else |
| 187 | { |
| 188 | if(h2c) str += "copy_h2c"; |
| 189 | else str += "copy_c2h"; |
| 190 | } |
| 191 | |
| 192 | str += "("; |
| 193 | |
| 194 | if(inIlvd) |
| 195 | { |
| 196 | str += "__global const "; str += r2Type; str += " * restrict gbIn, "; |
| 197 | } |
| 198 | else |
| 199 | { |
| 200 | str += "__global const "; str += rType; str += " * restrict gbInRe, "; |
| 201 | str += "__global const "; str += rType; str += " * restrict gbInIm, "; |
| 202 | } |
no test coverage detected