KJB
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
fft_base_kernels.h
Go to the documentation of this file.
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 //
47 
48 
49 #ifndef __CL_FFT_BASE_KERNELS_
50 #define __CL_FFT_BASE_KERNELS_
51 
52 #ifdef KJB_HAVE_OPENCL
53 #include <string>
54 
55 #warning "[Code police] Do not put 'using namespace' in global scope of header."
56 using namespace std;
57 
58 static string baseKernels = string(
59  "#ifndef M_PI\n"
60  "#define M_PI 0x1.921fb54442d18p+1\n"
61  "#endif\n"
62  "#define complexMul(a,b) ((float2)(mad(-(a).y, (b).y, (a).x * (b).x), mad((a).y, (b).x, (a).x * (b).y)))\n"
63  "#define conj(a) ((float2)((a).x, -(a).y))\n"
64  "#define conjTransp(a) ((float2)(-(a).y, (a).x))\n"
65  "\n"
66  "#define fftKernel2(a,dir) \\\n"
67  "{ \\\n"
68  " float2 c = (a)[0]; \\\n"
69  " (a)[0] = c + (a)[1]; \\\n"
70  " (a)[1] = c - (a)[1]; \\\n"
71  "}\n"
72  "\n"
73  "#define fftKernel2S(d1,d2,dir) \\\n"
74  "{ \\\n"
75  " float2 c = (d1); \\\n"
76  " (d1) = c + (d2); \\\n"
77  " (d2) = c - (d2); \\\n"
78  "}\n"
79  "\n"
80  "#define fftKernel4(a,dir) \\\n"
81  "{ \\\n"
82  " fftKernel2S((a)[0], (a)[2], dir); \\\n"
83  " fftKernel2S((a)[1], (a)[3], dir); \\\n"
84  " fftKernel2S((a)[0], (a)[1], dir); \\\n"
85  " (a)[3] = (float2)(dir)*(conjTransp((a)[3])); \\\n"
86  " fftKernel2S((a)[2], (a)[3], dir); \\\n"
87  " float2 c = (a)[1]; \\\n"
88  " (a)[1] = (a)[2]; \\\n"
89  " (a)[2] = c; \\\n"
90  "}\n"
91  "\n"
92  "#define fftKernel4s(a0,a1,a2,a3,dir) \\\n"
93  "{ \\\n"
94  " fftKernel2S((a0), (a2), dir); \\\n"
95  " fftKernel2S((a1), (a3), dir); \\\n"
96  " fftKernel2S((a0), (a1), dir); \\\n"
97  " (a3) = (float2)(dir)*(conjTransp((a3))); \\\n"
98  " fftKernel2S((a2), (a3), dir); \\\n"
99  " float2 c = (a1); \\\n"
100  " (a1) = (a2); \\\n"
101  " (a2) = c; \\\n"
102  "}\n"
103  "\n"
104  "#define bitreverse8(a) \\\n"
105  "{ \\\n"
106  " float2 c; \\\n"
107  " c = (a)[1]; \\\n"
108  " (a)[1] = (a)[4]; \\\n"
109  " (a)[4] = c; \\\n"
110  " c = (a)[3]; \\\n"
111  " (a)[3] = (a)[6]; \\\n"
112  " (a)[6] = c; \\\n"
113  "}\n"
114  "\n"
115  "#define fftKernel8(a,dir) \\\n"
116  "{ \\\n"
117  " const float2 w1 = (float2)(0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f); \\\n"
118  " const float2 w3 = (float2)(-0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f); \\\n"
119  " float2 c; \\\n"
120  " fftKernel2S((a)[0], (a)[4], dir); \\\n"
121  " fftKernel2S((a)[1], (a)[5], dir); \\\n"
122  " fftKernel2S((a)[2], (a)[6], dir); \\\n"
123  " fftKernel2S((a)[3], (a)[7], dir); \\\n"
124  " (a)[5] = complexMul(w1, (a)[5]); \\\n"
125  " (a)[6] = (float2)(dir)*(conjTransp((a)[6])); \\\n"
126  " (a)[7] = complexMul(w3, (a)[7]); \\\n"
127  " fftKernel2S((a)[0], (a)[2], dir); \\\n"
128  " fftKernel2S((a)[1], (a)[3], dir); \\\n"
129  " fftKernel2S((a)[4], (a)[6], dir); \\\n"
130  " fftKernel2S((a)[5], (a)[7], dir); \\\n"
131  " (a)[3] = (float2)(dir)*(conjTransp((a)[3])); \\\n"
132  " (a)[7] = (float2)(dir)*(conjTransp((a)[7])); \\\n"
133  " fftKernel2S((a)[0], (a)[1], dir); \\\n"
134  " fftKernel2S((a)[2], (a)[3], dir); \\\n"
135  " fftKernel2S((a)[4], (a)[5], dir); \\\n"
136  " fftKernel2S((a)[6], (a)[7], dir); \\\n"
137  " bitreverse8((a)); \\\n"
138  "}\n"
139  "\n"
140  "#define bitreverse4x4(a) \\\n"
141  "{ \\\n"
142  " float2 c; \\\n"
143  " c = (a)[1]; (a)[1] = (a)[4]; (a)[4] = c; \\\n"
144  " c = (a)[2]; (a)[2] = (a)[8]; (a)[8] = c; \\\n"
145  " c = (a)[3]; (a)[3] = (a)[12]; (a)[12] = c; \\\n"
146  " c = (a)[6]; (a)[6] = (a)[9]; (a)[9] = c; \\\n"
147  " c = (a)[7]; (a)[7] = (a)[13]; (a)[13] = c; \\\n"
148  " c = (a)[11]; (a)[11] = (a)[14]; (a)[14] = c; \\\n"
149  "}\n"
150  "\n"
151  "#define fftKernel16(a,dir) \\\n"
152  "{ \\\n"
153  " const float w0 = 0x1.d906bcp-1f; \\\n"
154  " const float w1 = 0x1.87de2ap-2f; \\\n"
155  " const float w2 = 0x1.6a09e6p-1f; \\\n"
156  " fftKernel4s((a)[0], (a)[4], (a)[8], (a)[12], dir); \\\n"
157  " fftKernel4s((a)[1], (a)[5], (a)[9], (a)[13], dir); \\\n"
158  " fftKernel4s((a)[2], (a)[6], (a)[10], (a)[14], dir); \\\n"
159  " fftKernel4s((a)[3], (a)[7], (a)[11], (a)[15], dir); \\\n"
160  " (a)[5] = complexMul((a)[5], (float2)(w0, dir*w1)); \\\n"
161  " (a)[6] = complexMul((a)[6], (float2)(w2, dir*w2)); \\\n"
162  " (a)[7] = complexMul((a)[7], (float2)(w1, dir*w0)); \\\n"
163  " (a)[9] = complexMul((a)[9], (float2)(w2, dir*w2)); \\\n"
164  " (a)[10] = (float2)(dir)*(conjTransp((a)[10])); \\\n"
165  " (a)[11] = complexMul((a)[11], (float2)(-w2, dir*w2)); \\\n"
166  " (a)[13] = complexMul((a)[13], (float2)(w1, dir*w0)); \\\n"
167  " (a)[14] = complexMul((a)[14], (float2)(-w2, dir*w2)); \\\n"
168  " (a)[15] = complexMul((a)[15], (float2)(-w0, dir*-w1)); \\\n"
169  " fftKernel4((a), dir); \\\n"
170  " fftKernel4((a) + 4, dir); \\\n"
171  " fftKernel4((a) + 8, dir); \\\n"
172  " fftKernel4((a) + 12, dir); \\\n"
173  " bitreverse4x4((a)); \\\n"
174  "}\n"
175  "\n"
176  "#define bitreverse32(a) \\\n"
177  "{ \\\n"
178  " float2 c1, c2; \\\n"
179  " 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"
180  " 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"
181  " 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"
182  " 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"
183  " 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"
184  " 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"
185  "}\n"
186  "\n"
187  "#define fftKernel32(a,dir) \\\n"
188  "{ \\\n"
189  " fftKernel2S((a)[0], (a)[16], dir); \\\n"
190  " fftKernel2S((a)[1], (a)[17], dir); \\\n"
191  " fftKernel2S((a)[2], (a)[18], dir); \\\n"
192  " fftKernel2S((a)[3], (a)[19], dir); \\\n"
193  " fftKernel2S((a)[4], (a)[20], dir); \\\n"
194  " fftKernel2S((a)[5], (a)[21], dir); \\\n"
195  " fftKernel2S((a)[6], (a)[22], dir); \\\n"
196  " fftKernel2S((a)[7], (a)[23], dir); \\\n"
197  " fftKernel2S((a)[8], (a)[24], dir); \\\n"
198  " fftKernel2S((a)[9], (a)[25], dir); \\\n"
199  " fftKernel2S((a)[10], (a)[26], dir); \\\n"
200  " fftKernel2S((a)[11], (a)[27], dir); \\\n"
201  " fftKernel2S((a)[12], (a)[28], dir); \\\n"
202  " fftKernel2S((a)[13], (a)[29], dir); \\\n"
203  " fftKernel2S((a)[14], (a)[30], dir); \\\n"
204  " fftKernel2S((a)[15], (a)[31], dir); \\\n"
205  " (a)[17] = complexMul((a)[17], (float2)(0x1.f6297cp-1f, dir*0x1.8f8b84p-3f)); \\\n"
206  " (a)[18] = complexMul((a)[18], (float2)(0x1.d906bcp-1f, dir*0x1.87de2ap-2f)); \\\n"
207  " (a)[19] = complexMul((a)[19], (float2)(0x1.a9b662p-1f, dir*0x1.1c73b4p-1f)); \\\n"
208  " (a)[20] = complexMul((a)[20], (float2)(0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f)); \\\n"
209  " (a)[21] = complexMul((a)[21], (float2)(0x1.1c73b4p-1f, dir*0x1.a9b662p-1f)); \\\n"
210  " (a)[22] = complexMul((a)[22], (float2)(0x1.87de2ap-2f, dir*0x1.d906bcp-1f)); \\\n"
211  " (a)[23] = complexMul((a)[23], (float2)(0x1.8f8b84p-3f, dir*0x1.f6297cp-1f)); \\\n"
212  " (a)[24] = complexMul((a)[24], (float2)(0x0p+0f, dir*0x1p+0f)); \\\n"
213  " (a)[25] = complexMul((a)[25], (float2)(-0x1.8f8b84p-3f, dir*0x1.f6297cp-1f)); \\\n"
214  " (a)[26] = complexMul((a)[26], (float2)(-0x1.87de2ap-2f, dir*0x1.d906bcp-1f)); \\\n"
215  " (a)[27] = complexMul((a)[27], (float2)(-0x1.1c73b4p-1f, dir*0x1.a9b662p-1f)); \\\n"
216  " (a)[28] = complexMul((a)[28], (float2)(-0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f)); \\\n"
217  " (a)[29] = complexMul((a)[29], (float2)(-0x1.a9b662p-1f, dir*0x1.1c73b4p-1f)); \\\n"
218  " (a)[30] = complexMul((a)[30], (float2)(-0x1.d906bcp-1f, dir*0x1.87de2ap-2f)); \\\n"
219  " (a)[31] = complexMul((a)[31], (float2)(-0x1.f6297cp-1f, dir*0x1.8f8b84p-3f)); \\\n"
220  " fftKernel16((a), dir); \\\n"
221  " fftKernel16((a) + 16, dir); \\\n"
222  " bitreverse32((a)); \\\n"
223  "}\n\n"
224  );
225 
226 static string twistKernelInterleaved = string(
227  "__kernel void \\\n"
228  "clFFT_1DTwistInterleaved(__global float2 *in, unsigned int startRow, unsigned int numCols, unsigned int N, unsigned int numRowsToProcess, int dir) \\\n"
229  "{ \\\n"
230  " float2 a, w; \\\n"
231  " float ang; \\\n"
232  " unsigned int j; \\\n"
233  " unsigned int i = get_global_id(0); \\\n"
234  " unsigned int startIndex = i; \\\n"
235  " \\\n"
236  " if(i < numCols) \\\n"
237  " { \\\n"
238  " for(j = 0; j < numRowsToProcess; j++) \\\n"
239  " { \\\n"
240  " a = in[startIndex]; \\\n"
241  " ang = 2.0f * M_PI * dir * i * (startRow + j) / N; \\\n"
242  " w = (float2)(native_cos(ang), native_sin(ang)); \\\n"
243  " a = complexMul(a, w); \\\n"
244  " in[startIndex] = a; \\\n"
245  " startIndex += numCols; \\\n"
246  " } \\\n"
247  " } \\\n"
248  "} \\\n"
249  );
250 
251 static string twistKernelPlannar = string(
252  "__kernel void \\\n"
253  "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"
254  "{ \\\n"
255  " float2 a, w; \\\n"
256  " float ang; \\\n"
257  " unsigned int j; \\\n"
258  " unsigned int i = get_global_id(0); \\\n"
259  " unsigned int startIndex = i; \\\n"
260  " \\\n"
261  " if(i < numCols) \\\n"
262  " { \\\n"
263  " for(j = 0; j < numRowsToProcess; j++) \\\n"
264  " { \\\n"
265  " a = (float2)(in_real[startIndex], in_imag[startIndex]); \\\n"
266  " ang = 2.0f * M_PI * dir * i * (startRow + j) / N; \\\n"
267  " w = (float2)(native_cos(ang), native_sin(ang)); \\\n"
268  " a = complexMul(a, w); \\\n"
269  " in_real[startIndex] = a.x; \\\n"
270  " in_imag[startIndex] = a.y; \\\n"
271  " startIndex += numCols; \\\n"
272  " } \\\n"
273  " } \\\n"
274  "} \\\n"
275  );
276 
277 
278 
279 #endif /* KJB_HAVE_OPENCL */
280 #endif
Generated by   doxygen 1.8.6