|
2
|
1
|
|
|
2 //
|
|
|
3 // File: fft_base_kernels.h
|
|
|
4 //
|
|
|
5 // Version: <1.0>
|
|
|
6 //
|
|
|
7 // Disclaimer: IMPORTANT: This Apple software is supplied to you by Apple Inc. ("Apple")
|
|
|
8 // in consideration of your agreement to the following terms, and your use,
|
|
|
9 // installation, modification or redistribution of this Apple software
|
|
|
10 // constitutes acceptance of these terms. If you do not agree with these
|
|
|
11 // terms, please do not use, install, modify or redistribute this Apple
|
|
|
12 // software.
|
|
|
13 //
|
|
|
14 // In consideration of your agreement to abide by the following terms, and
|
|
|
15 // subject to these terms, Apple grants you a personal, non - exclusive
|
|
|
16 // license, under Apple's copyrights in this original Apple software ( the
|
|
|
17 // "Apple Software" ), to use, reproduce, modify and redistribute the Apple
|
|
|
18 // Software, with or without modifications, in source and / or binary forms;
|
|
|
19 // provided that if you redistribute the Apple Software in its entirety and
|
|
|
20 // without modifications, you must retain this notice and the following text
|
|
|
21 // and disclaimers in all such redistributions of the Apple Software. Neither
|
|
|
22 // the name, trademarks, service marks or logos of Apple Inc. may be used to
|
|
|
23 // endorse or promote products derived from the Apple Software without specific
|
|
|
24 // prior written permission from Apple. Except as expressly stated in this
|
|
|
25 // notice, no other rights or licenses, express or implied, are granted by
|
|
|
26 // Apple herein, including but not limited to any patent rights that may be
|
|
|
27 // infringed by your derivative works or by other works in which the Apple
|
|
|
28 // Software may be incorporated.
|
|
|
29 //
|
|
|
30 // The Apple Software is provided by Apple on an "AS IS" basis. APPLE MAKES NO
|
|
|
31 // WARRANTIES, EXPRESS OR IMPLIED, INCLUDING WITHOUT LIMITATION THE IMPLIED
|
|
|
32 // WARRANTIES OF NON - INFRINGEMENT, MERCHANTABILITY AND FITNESS FOR A
|
|
|
33 // PARTICULAR PURPOSE, REGARDING THE APPLE SOFTWARE OR ITS USE AND OPERATION
|
|
|
34 // ALONE OR IN COMBINATION WITH YOUR PRODUCTS.
|
|
|
35 //
|
|
|
36 // IN NO EVENT SHALL APPLE BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL OR
|
|
|
37 // CONSEQUENTIAL DAMAGES ( INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
|
|
|
38 // SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
|
|
|
39 // INTERRUPTION ) ARISING IN ANY WAY OUT OF THE USE, REPRODUCTION, MODIFICATION
|
|
|
40 // AND / OR DISTRIBUTION OF THE APPLE SOFTWARE, HOWEVER CAUSED AND WHETHER
|
|
|
41 // UNDER THEORY OF CONTRACT, TORT ( INCLUDING NEGLIGENCE ), STRICT LIABILITY OR
|
|
|
42 // OTHERWISE, EVEN IF APPLE HAS BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
|
|
43 //
|
|
|
44 // Copyright ( C ) 2008 Apple Inc. All Rights Reserved.
|
|
|
45 //
|
|
|
46 ////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
47
|
|
|
48
|
|
|
49 #ifndef __CL_FFT_BASE_KERNELS_
|
|
|
50 #define __CL_FFT_BASE_KERNELS_
|
|
|
51
|
|
|
52 #include <string>
|
|
|
53
|
|
|
54 using namespace std;
|
|
|
55
|
|
|
56 static string baseKernels = string(
|
|
|
57 "#ifndef M_PI\n"
|
|
7
|
58 "#define M_PI 0x1.921fb54442d18p+1\n"
|
|
|
59 "#endif\n"
|
|
|
60 "#define complexMul(a,b) ((float2)(mad(-(a).y, (b).y, (a).x * (b).x), mad((a).y, (b).x, (a).x * (b).y)))\n"
|
|
|
61 "#define conj(a) ((float2)((a).x, -(a).y))\n"
|
|
|
62 "#define conjTransp(a) ((float2)(-(a).y, (a).x))\n"
|
|
|
63 "\n"
|
|
|
64 "#define fftKernel2(a,dir) \\\n"
|
|
|
65 "{ \\\n"
|
|
|
66 " float2 c = (a)[0]; \\\n"
|
|
|
67 " (a)[0] = c + (a)[1]; \\\n"
|
|
|
68 " (a)[1] = c - (a)[1]; \\\n"
|
|
|
69 "}\n"
|
|
|
70 "\n"
|
|
|
71 "#define fftKernel2S(d1,d2,dir) \\\n"
|
|
|
72 "{ \\\n"
|
|
|
73 " float2 c = (d1); \\\n"
|
|
|
74 " (d1) = c + (d2); \\\n"
|
|
|
75 " (d2) = c - (d2); \\\n"
|
|
|
76 "}\n"
|
|
|
77 "\n"
|
|
|
78 "#define fftKernel4(a,dir) \\\n"
|
|
|
79 "{ \\\n"
|
|
|
80 " fftKernel2S((a)[0], (a)[2], dir); \\\n"
|
|
|
81 " fftKernel2S((a)[1], (a)[3], dir); \\\n"
|
|
|
82 " fftKernel2S((a)[0], (a)[1], dir); \\\n"
|
|
|
83 " (a)[3] = (float2)(dir)*(conjTransp((a)[3])); \\\n"
|
|
|
84 " fftKernel2S((a)[2], (a)[3], dir); \\\n"
|
|
|
85 " float2 c = (a)[1]; \\\n"
|
|
|
86 " (a)[1] = (a)[2]; \\\n"
|
|
|
87 " (a)[2] = c; \\\n"
|
|
|
88 "}\n"
|
|
|
89 "\n"
|
|
|
90 "#define fftKernel4s(a0,a1,a2,a3,dir) \\\n"
|
|
|
91 "{ \\\n"
|
|
|
92 " fftKernel2S((a0), (a2), dir); \\\n"
|
|
|
93 " fftKernel2S((a1), (a3), dir); \\\n"
|
|
|
94 " fftKernel2S((a0), (a1), dir); \\\n"
|
|
|
95 " (a3) = (float2)(dir)*(conjTransp((a3))); \\\n"
|
|
|
96 " fftKernel2S((a2), (a3), dir); \\\n"
|
|
|
97 " float2 c = (a1); \\\n"
|
|
|
98 " (a1) = (a2); \\\n"
|
|
|
99 " (a2) = c; \\\n"
|
|
|
100 "}\n"
|
|
|
101 "\n"
|
|
|
102 "#define bitreverse8(a) \\\n"
|
|
|
103 "{ \\\n"
|
|
|
104 " float2 c; \\\n"
|
|
|
105 " c = (a)[1]; \\\n"
|
|
|
106 " (a)[1] = (a)[4]; \\\n"
|
|
|
107 " (a)[4] = c; \\\n"
|
|
|
108 " c = (a)[3]; \\\n"
|
|
|
109 " (a)[3] = (a)[6]; \\\n"
|
|
|
110 " (a)[6] = c; \\\n"
|
|
|
111 "}\n"
|
|
|
112 "\n"
|
|
|
113 "#define fftKernel8(a,dir) \\\n"
|
|
|
114 "{ \\\n"
|
|
|
115 " const float2 w1 = (float2)(0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f); \\\n"
|
|
|
116 " const float2 w3 = (float2)(-0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f); \\\n"
|
|
|
117 " float2 c; \\\n"
|
|
|
118 " fftKernel2S((a)[0], (a)[4], dir); \\\n"
|
|
|
119 " fftKernel2S((a)[1], (a)[5], dir); \\\n"
|
|
|
120 " fftKernel2S((a)[2], (a)[6], dir); \\\n"
|
|
|
121 " fftKernel2S((a)[3], (a)[7], dir); \\\n"
|
|
|
122 " (a)[5] = complexMul(w1, (a)[5]); \\\n"
|
|
|
123 " (a)[6] = (float2)(dir)*(conjTransp((a)[6])); \\\n"
|
|
|
124 " (a)[7] = complexMul(w3, (a)[7]); \\\n"
|
|
|
125 " fftKernel2S((a)[0], (a)[2], dir); \\\n"
|
|
|
126 " fftKernel2S((a)[1], (a)[3], dir); \\\n"
|
|
|
127 " fftKernel2S((a)[4], (a)[6], dir); \\\n"
|
|
|
128 " fftKernel2S((a)[5], (a)[7], dir); \\\n"
|
|
|
129 " (a)[3] = (float2)(dir)*(conjTransp((a)[3])); \\\n"
|
|
|
130 " (a)[7] = (float2)(dir)*(conjTransp((a)[7])); \\\n"
|
|
|
131 " fftKernel2S((a)[0], (a)[1], dir); \\\n"
|
|
|
132 " fftKernel2S((a)[2], (a)[3], dir); \\\n"
|
|
|
133 " fftKernel2S((a)[4], (a)[5], dir); \\\n"
|
|
|
134 " fftKernel2S((a)[6], (a)[7], dir); \\\n"
|
|
|
135 " bitreverse8((a)); \\\n"
|
|
|
136 "}\n"
|
|
|
137 "\n"
|
|
|
138 "#define bitreverse4x4(a) \\\n"
|
|
|
139 "{ \\\n"
|
|
|
140 " float2 c; \\\n"
|
|
|
141 " c = (a)[1]; (a)[1] = (a)[4]; (a)[4] = c; \\\n"
|
|
|
142 " c = (a)[2]; (a)[2] = (a)[8]; (a)[8] = c; \\\n"
|
|
|
143 " c = (a)[3]; (a)[3] = (a)[12]; (a)[12] = c; \\\n"
|
|
|
144 " c = (a)[6]; (a)[6] = (a)[9]; (a)[9] = c; \\\n"
|
|
|
145 " c = (a)[7]; (a)[7] = (a)[13]; (a)[13] = c; \\\n"
|
|
|
146 " c = (a)[11]; (a)[11] = (a)[14]; (a)[14] = c; \\\n"
|
|
|
147 "}\n"
|
|
|
148 "\n"
|
|
|
149 "#define fftKernel16(a,dir) \\\n"
|
|
|
150 "{ \\\n"
|
|
|
151 " const float w0 = 0x1.d906bcp-1f; \\\n"
|
|
|
152 " const float w1 = 0x1.87de2ap-2f; \\\n"
|
|
|
153 " const float w2 = 0x1.6a09e6p-1f; \\\n"
|
|
|
154 " fftKernel4s((a)[0], (a)[4], (a)[8], (a)[12], dir); \\\n"
|
|
|
155 " fftKernel4s((a)[1], (a)[5], (a)[9], (a)[13], dir); \\\n"
|
|
|
156 " fftKernel4s((a)[2], (a)[6], (a)[10], (a)[14], dir); \\\n"
|
|
|
157 " fftKernel4s((a)[3], (a)[7], (a)[11], (a)[15], dir); \\\n"
|
|
|
158 " (a)[5] = complexMul((a)[5], (float2)(w0, dir*w1)); \\\n"
|
|
|
159 " (a)[6] = complexMul((a)[6], (float2)(w2, dir*w2)); \\\n"
|
|
|
160 " (a)[7] = complexMul((a)[7], (float2)(w1, dir*w0)); \\\n"
|
|
|
161 " (a)[9] = complexMul((a)[9], (float2)(w2, dir*w2)); \\\n"
|
|
|
162 " (a)[10] = (float2)(dir)*(conjTransp((a)[10])); \\\n"
|
|
|
163 " (a)[11] = complexMul((a)[11], (float2)(-w2, dir*w2)); \\\n"
|
|
|
164 " (a)[13] = complexMul((a)[13], (float2)(w1, dir*w0)); \\\n"
|
|
|
165 " (a)[14] = complexMul((a)[14], (float2)(-w2, dir*w2)); \\\n"
|
|
|
166 " (a)[15] = complexMul((a)[15], (float2)(-w0, dir*-w1)); \\\n"
|
|
|
167 " fftKernel4((a), dir); \\\n"
|
|
|
168 " fftKernel4((a) + 4, dir); \\\n"
|
|
|
169 " fftKernel4((a) + 8, dir); \\\n"
|
|
|
170 " fftKernel4((a) + 12, dir); \\\n"
|
|
|
171 " bitreverse4x4((a)); \\\n"
|
|
|
172 "}\n"
|
|
|
173 "\n"
|
|
|
174 "#define bitreverse32(a) \\\n"
|
|
|
175 "{ \\\n"
|
|
|
176 " float2 c1, c2; \\\n"
|
|
|
177 " c1 = (a)[2]; (a)[2] = (a)[1]; c2 = (a)[4]; (a)[4] = c1; c1 = (a)[8]; (a)[8] = c2; c2 = (a)[16]; (a)[16] = c1; (a)[1] = c2; \\\n"
|
|
|
178 " c1 = (a)[6]; (a)[6] = (a)[3]; c2 = (a)[12]; (a)[12] = c1; c1 = (a)[24]; (a)[24] = c2; c2 = (a)[17]; (a)[17] = c1; (a)[3] = c2; \\\n"
|
|
|
179 " c1 = (a)[10]; (a)[10] = (a)[5]; c2 = (a)[20]; (a)[20] = c1; c1 = (a)[9]; (a)[9] = c2; c2 = (a)[18]; (a)[18] = c1; (a)[5] = c2; \\\n"
|
|
|
180 " c1 = (a)[14]; (a)[14] = (a)[7]; c2 = (a)[28]; (a)[28] = c1; c1 = (a)[25]; (a)[25] = c2; c2 = (a)[19]; (a)[19] = c1; (a)[7] = c2; \\\n"
|
|
|
181 " c1 = (a)[22]; (a)[22] = (a)[11]; c2 = (a)[13]; (a)[13] = c1; c1 = (a)[26]; (a)[26] = c2; c2 = (a)[21]; (a)[21] = c1; (a)[11] = c2; \\\n"
|
|
|
182 " c1 = (a)[30]; (a)[30] = (a)[15]; c2 = (a)[29]; (a)[29] = c1; c1 = (a)[27]; (a)[27] = c2; c2 = (a)[23]; (a)[23] = c1; (a)[15] = c2; \\\n"
|
|
|
183 "}\n"
|
|
|
184 "\n"
|
|
|
185 "#define fftKernel32(a,dir) \\\n"
|
|
|
186 "{ \\\n"
|
|
|
187 " fftKernel2S((a)[0], (a)[16], dir); \\\n"
|
|
|
188 " fftKernel2S((a)[1], (a)[17], dir); \\\n"
|
|
|
189 " fftKernel2S((a)[2], (a)[18], dir); \\\n"
|
|
|
190 " fftKernel2S((a)[3], (a)[19], dir); \\\n"
|
|
|
191 " fftKernel2S((a)[4], (a)[20], dir); \\\n"
|
|
|
192 " fftKernel2S((a)[5], (a)[21], dir); \\\n"
|
|
|
193 " fftKernel2S((a)[6], (a)[22], dir); \\\n"
|
|
|
194 " fftKernel2S((a)[7], (a)[23], dir); \\\n"
|
|
|
195 " fftKernel2S((a)[8], (a)[24], dir); \\\n"
|
|
|
196 " fftKernel2S((a)[9], (a)[25], dir); \\\n"
|
|
|
197 " fftKernel2S((a)[10], (a)[26], dir); \\\n"
|
|
|
198 " fftKernel2S((a)[11], (a)[27], dir); \\\n"
|
|
|
199 " fftKernel2S((a)[12], (a)[28], dir); \\\n"
|
|
|
200 " fftKernel2S((a)[13], (a)[29], dir); \\\n"
|
|
|
201 " fftKernel2S((a)[14], (a)[30], dir); \\\n"
|
|
|
202 " fftKernel2S((a)[15], (a)[31], dir); \\\n"
|
|
|
203 " (a)[17] = complexMul((a)[17], (float2)(0x1.f6297cp-1f, dir*0x1.8f8b84p-3f)); \\\n"
|
|
|
204 " (a)[18] = complexMul((a)[18], (float2)(0x1.d906bcp-1f, dir*0x1.87de2ap-2f)); \\\n"
|
|
|
205 " (a)[19] = complexMul((a)[19], (float2)(0x1.a9b662p-1f, dir*0x1.1c73b4p-1f)); \\\n"
|
|
|
206 " (a)[20] = complexMul((a)[20], (float2)(0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f)); \\\n"
|
|
|
207 " (a)[21] = complexMul((a)[21], (float2)(0x1.1c73b4p-1f, dir*0x1.a9b662p-1f)); \\\n"
|
|
|
208 " (a)[22] = complexMul((a)[22], (float2)(0x1.87de2ap-2f, dir*0x1.d906bcp-1f)); \\\n"
|
|
|
209 " (a)[23] = complexMul((a)[23], (float2)(0x1.8f8b84p-3f, dir*0x1.f6297cp-1f)); \\\n"
|
|
|
210 " (a)[24] = complexMul((a)[24], (float2)(0x0p+0f, dir*0x1p+0f)); \\\n"
|
|
|
211 " (a)[25] = complexMul((a)[25], (float2)(-0x1.8f8b84p-3f, dir*0x1.f6297cp-1f)); \\\n"
|
|
|
212 " (a)[26] = complexMul((a)[26], (float2)(-0x1.87de2ap-2f, dir*0x1.d906bcp-1f)); \\\n"
|
|
|
213 " (a)[27] = complexMul((a)[27], (float2)(-0x1.1c73b4p-1f, dir*0x1.a9b662p-1f)); \\\n"
|
|
|
214 " (a)[28] = complexMul((a)[28], (float2)(-0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f)); \\\n"
|
|
|
215 " (a)[29] = complexMul((a)[29], (float2)(-0x1.a9b662p-1f, dir*0x1.1c73b4p-1f)); \\\n"
|
|
|
216 " (a)[30] = complexMul((a)[30], (float2)(-0x1.d906bcp-1f, dir*0x1.87de2ap-2f)); \\\n"
|
|
|
217 " (a)[31] = complexMul((a)[31], (float2)(-0x1.f6297cp-1f, dir*0x1.8f8b84p-3f)); \\\n"
|
|
|
218 " fftKernel16((a), dir); \\\n"
|
|
|
219 " fftKernel16((a) + 16, dir); \\\n"
|
|
|
220 " bitreverse32((a)); \\\n"
|
|
|
221 "}\n\n"
|
|
|
222 );
|
|
2
|
223
|
|
|
224 static string twistKernelInterleaved = string(
|
|
7
|
225 "__kernel void \\\n"
|
|
|
226 "clFFT_1DTwistInterleaved(__global float2 *in, unsigned int startRow, unsigned int numCols, unsigned int N, unsigned int numRowsToProcess, int dir) \\\n"
|
|
|
227 "{ \\\n"
|
|
|
228 " float2 a, w; \\\n"
|
|
|
229 " float ang; \\\n"
|
|
|
230 " unsigned int j; \\\n"
|
|
|
231 " unsigned int i = get_global_id(0); \\\n"
|
|
|
232 " unsigned int startIndex = i; \\\n"
|
|
|
233 " \\\n"
|
|
|
234 " if(i < numCols) \\\n"
|
|
|
235 " { \\\n"
|
|
|
236 " for(j = 0; j < numRowsToProcess; j++) \\\n"
|
|
|
237 " { \\\n"
|
|
|
238 " a = in[startIndex]; \\\n"
|
|
|
239 " ang = 2.0f * M_PI * dir * i * (startRow + j) / N; \\\n"
|
|
|
240 " w = (float2)(native_cos(ang), native_sin(ang)); \\\n"
|
|
|
241 " a = complexMul(a, w); \\\n"
|
|
|
242 " in[startIndex] = a; \\\n"
|
|
|
243 " startIndex += numCols; \\\n"
|
|
|
244 " } \\\n"
|
|
|
245 " } \\\n"
|
|
|
246 "} \\\n"
|
|
|
247 );
|
|
2
|
248
|
|
|
249 static string twistKernelPlannar = string(
|
|
7
|
250 "__kernel void \\\n"
|
|
|
251 "clFFT_1DTwistSplit(__global float *in_real, __global float *in_imag , unsigned int startRow, unsigned int numCols, unsigned int N, unsigned int numRowsToProcess, int dir) \\\n"
|
|
|
252 "{ \\\n"
|
|
|
253 " float2 a, w; \\\n"
|
|
|
254 " float ang; \\\n"
|
|
|
255 " unsigned int j; \\\n"
|
|
|
256 " unsigned int i = get_global_id(0); \\\n"
|
|
|
257 " unsigned int startIndex = i; \\\n"
|
|
|
258 " \\\n"
|
|
|
259 " if(i < numCols) \\\n"
|
|
|
260 " { \\\n"
|
|
|
261 " for(j = 0; j < numRowsToProcess; j++) \\\n"
|
|
|
262 " { \\\n"
|
|
|
263 " a = (float2)(in_real[startIndex], in_imag[startIndex]); \\\n"
|
|
|
264 " ang = 2.0f * M_PI * dir * i * (startRow + j) / N; \\\n"
|
|
|
265 " w = (float2)(native_cos(ang), native_sin(ang)); \\\n"
|
|
|
266 " a = complexMul(a, w); \\\n"
|
|
|
267 " in_real[startIndex] = a.x; \\\n"
|
|
|
268 " in_imag[startIndex] = a.y; \\\n"
|
|
|
269 " startIndex += numCols; \\\n"
|
|
|
270 " } \\\n"
|
|
|
271 " } \\\n"
|
|
|
272 "} \\\n"
|
|
|
273 );
|
|
2
|
274
|
|
|
275
|
|
|
276
|
|
|
277 #endif
|
