/opencl/oclfft : revision 1

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

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"

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

);

223

224

static string twistKernelInterleaved = string(

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

);

248

249

static string twistKernelPlannar = string(

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

);

274

275

276

277

#endif