fft_execute.cpp 17.1 KB
Newer Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
/***************************************************************************
 *   Copyright (C) 2012 by Oliver Bock,Heinz-Bernd Eggenstein              *
 *   oliver.bock[AT]aei.mpg.de                                             *
 *   heinz-bernd.eggenstein[AT]aei.mpg.de                                  *
 *                                                                         *
 *   This file is part of libclfft (originally for Einstein@Home)          *
 *   Derived from clFFT,  (C) Apple, see notice below.                     *
 *                                                                         *
 *                                                                         *
 *   libclfft  is distributed in the hope that it will be useful,          *
 *   but WITHOUT ANY WARRANTY; without even the implied warranty of        *
 *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See              *
 *   notice below for more details.                                        *
 *                                                                         *
 ***************************************************************************/
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
//
// File:       fft_execute.cpp
//
// Version:    <1.0>
//
// Disclaimer: IMPORTANT:  This Apple software is supplied to you by Apple Inc. ("Apple")
//             in consideration of your agreement to the following terms, and your use,
//             installation, modification or redistribution of this Apple software
//             constitutes acceptance of these terms.  If you do not agree with these
//             terms, please do not use, install, modify or redistribute this Apple
//             software.¬
//
//             In consideration of your agreement to abide by the following terms, and
//             subject to these terms, Apple grants you a personal, non - exclusive
//             license, under Apple's copyrights in this original Apple software ( the
//             "Apple Software" ), to use, reproduce, modify and redistribute the Apple
//             Software, with or without modifications, in source and / or binary forms;
//             provided that if you redistribute the Apple Software in its entirety and
//             without modifications, you must retain this notice and the following text
//             and disclaimers in all such redistributions of the Apple Software. Neither
//             the name, trademarks, service marks or logos of Apple Inc. may be used to
//             endorse or promote products derived from the Apple Software without specific
//             prior written permission from Apple.  Except as expressly stated in this
//             notice, no other rights or licenses, express or implied, are granted by
//             Apple herein, including but not limited to any patent rights that may be
//             infringed by your derivative works or by other works in which the Apple
//             Software may be incorporated.
//
//             The Apple Software is provided by Apple on an "AS IS" basis.  APPLE MAKES NO
//             WARRANTIES, EXPRESS OR IMPLIED, INCLUDING WITHOUT LIMITATION THE IMPLIED
//             WARRANTIES OF NON - INFRINGEMENT, MERCHANTABILITY AND FITNESS FOR A
//             PARTICULAR PURPOSE, REGARDING THE APPLE SOFTWARE OR ITS USE AND OPERATION
//             ALONE OR IN COMBINATION WITH YOUR PRODUCTS.
//
//             IN NO EVENT SHALL APPLE BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL OR
//             CONSEQUENTIAL DAMAGES ( INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
//             SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
//             INTERRUPTION ) ARISING IN ANY WAY OUT OF THE USE, REPRODUCTION, MODIFICATION
//             AND / OR DISTRIBUTION OF THE APPLE SOFTWARE, HOWEVER CAUSED AND WHETHER
//             UNDER THEORY OF CONTRACT, TORT ( INCLUDING NEGLIGENCE ), STRICT LIABILITY OR
//             OTHERWISE, EVEN IF APPLE HAS BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Copyright ( C ) 2008 Apple Inc. All Rights Reserved.
//
////////////////////////////////////////////////////////////////////////////////////////////////////


#include "fft_internal.h"
#include "clFFT.h"
#include <stdlib.h>
#include <stdio.h>
#include <math.h>

#define max(a,b) (((a)>(b)) ? (a) : (b))
#define min(a,b) (((a)<(b)) ? (a) : (b))

static cl_int
allocateTemporaryBufferInterleaved(cl_fft_plan *plan, cl_uint batchSize)
{
Oliver Bock's avatar
Oliver Bock committed
75
76
77
78
79
80
81
82
83
84
85
86
    cl_int err = CL_SUCCESS;
    if(plan->temp_buffer_needed && plan->last_batch_size != batchSize)
    {
        plan->last_batch_size = batchSize;
        size_t tmpLength = plan->n.x * plan->n.y * plan->n.z * batchSize * 2 * sizeof(cl_float);

        if(plan->tempmemobj)
            clReleaseMemObject(plan->tempmemobj);

        plan->tempmemobj = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, NULL, &err);
    }
    return err;
87
88
89
90
91
}

static cl_int
allocateTemporaryBufferPlannar(cl_fft_plan *plan, cl_uint batchSize)
{
Oliver Bock's avatar
Oliver Bock committed
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
    cl_int err = CL_SUCCESS;
    cl_int terr;
    if(plan->temp_buffer_needed && plan->last_batch_size != batchSize)
    {
        plan->last_batch_size = batchSize;
        size_t tmpLength = plan->n.x * plan->n.y * plan->n.z * batchSize * sizeof(cl_float);

        if(plan->tempmemobj_real)
            clReleaseMemObject(plan->tempmemobj_real);

        if(plan->tempmemobj_imag)
            clReleaseMemObject(plan->tempmemobj_imag);

        plan->tempmemobj_real = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, NULL, &err);
        plan->tempmemobj_imag = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, NULL, &terr);
        err |= terr;
    }
    return err;
110
111
112
113
114
}

void
getKernelWorkDimensions(cl_fft_plan *plan, cl_fft_kernel_info *kernelInfo, cl_int *batchSize, size_t *gWorkItems, size_t *lWorkItems)
{
Oliver Bock's avatar
Oliver Bock committed
115
116
    *lWorkItems = kernelInfo->num_workitems_per_workgroup;
    int numWorkGroups = kernelInfo->num_workgroups;
117
    int numXFormsPerWG = kernelInfo->num_xforms_per_workgroup;
Oliver Bock's avatar
Oliver Bock committed
118

Oliver Bock's avatar
Oliver Bock committed
119
120
121
    switch(kernelInfo->dir)
    {
        case cl_fft_kernel_x:
122
123
124
            *batchSize *= (plan->n.y * plan->n.z);
            numWorkGroups = (*batchSize % numXFormsPerWG) ? (*batchSize/numXFormsPerWG + 1) : (*batchSize/numXFormsPerWG);
            numWorkGroups *= kernelInfo->num_workgroups;
Oliver Bock's avatar
Oliver Bock committed
125
126
127
128
129
130
131
132
133
134
135
            break;
        case cl_fft_kernel_y:
            *batchSize *= plan->n.z;
            numWorkGroups *= *batchSize;
            break;
        case cl_fft_kernel_z:
            numWorkGroups *= *batchSize;
            break;
    }

    *gWorkItems = numWorkGroups * *lWorkItems;
136
137
}

Oliver Bock's avatar
Oliver Bock committed
138
139
cl_int
clFFT_ExecuteInterleaved( cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
Oliver Bock's avatar
Oliver Bock committed
140
141
                         cl_mem data_in, cl_mem data_out,
                         cl_int num_events, cl_event *event_list, cl_event *event )
Oliver Bock's avatar
Oliver Bock committed
142
{
Oliver Bock's avatar
Oliver Bock committed
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
    int s;
    cl_fft_plan *plan = (cl_fft_plan *) Plan;
    if(plan->format != clFFT_InterleavedComplexFormat)
        return CL_INVALID_VALUE;

    cl_int err;
    size_t gWorkItems, lWorkItems;
    int inPlaceDone = -1;

    cl_int isInPlace = data_in == data_out ? 1 : 0;

    if((err = allocateTemporaryBufferInterleaved(plan, batchSize)) != CL_SUCCESS)
        return err;

    cl_mem memObj[3];
    memObj[0] = data_in;
    memObj[1] = data_out;
    memObj[2] = plan->tempmemobj;
    cl_fft_kernel_info *kernelInfo = plan->kernel_info;
    int numKernels = plan->num_kernels;

    int numKernelsOdd = numKernels & 1;
    int currRead  = 0;
    int currWrite = 1;

    // at least one external dram shuffle (transpose) required
    if(plan->temp_buffer_needed)
    {
        // in-place transform
        if(isInPlace)
        {
            inPlaceDone = 0;
            currRead  = 1;
            currWrite = 2;
        }
        else
        {
            currWrite = (numKernels & 1) ? 1 : 2;
        }

        while(kernelInfo)
        {
            if( isInPlace && numKernelsOdd && !inPlaceDone && kernelInfo->in_place_possible)
            {
                currWrite = currRead;
                inPlaceDone = 1;
            }

            s = batchSize;
            getKernelWorkDimensions(plan, kernelInfo, &s, &gWorkItems, &lWorkItems);
            err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj[currRead]);
            err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj[currWrite]);
            err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_int), &dir);
            err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_int), &s);
197
198
            err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_mem), &(plan->cossin_LUT_d2));
199
            
Oliver Bock's avatar
Oliver Bock committed
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
            err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL, NULL);
            if(err)
                return err;

            currRead  = (currWrite == 1) ? 1 : 2;
            currWrite = (currWrite == 1) ? 2 : 1;

            kernelInfo = kernelInfo->next;
        }
    }
    // no dram shuffle (transpose required) transform
    // all kernels can execute in-place.
    else {

        while(kernelInfo)
        {
            s = batchSize;
            getKernelWorkDimensions(plan, kernelInfo, &s, &gWorkItems, &lWorkItems);
            err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj[currRead]);
            err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj[currWrite]);
            err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_int), &dir);
            err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_int), &s);
222
223
            err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_mem), &(plan->cossin_LUT_d2));
Oliver Bock's avatar
Oliver Bock committed
224
225
226
227
228
229
230
231
232
233
234
235
236

            err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL, NULL);
            if(err)
                return err;

            currRead  = 1;
            currWrite = 1;

            kernelInfo = kernelInfo->next;
        }
    }

    return err;
237
238
}

Oliver Bock's avatar
Oliver Bock committed
239
240
cl_int
clFFT_ExecutePlannar( cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
Oliver Bock's avatar
Oliver Bock committed
241
242
                      cl_mem data_in_real, cl_mem data_in_imag, cl_mem data_out_real, cl_mem data_out_imag,
                      cl_int num_events, cl_event *event_list, cl_event *event)
Oliver Bock's avatar
Oliver Bock committed
243
{
Oliver Bock's avatar
Oliver Bock committed
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
    int s;
    cl_fft_plan *plan = (cl_fft_plan *) Plan;

    if(plan->format != clFFT_SplitComplexFormat)
        return CL_INVALID_VALUE;

    cl_int err;
    size_t gWorkItems, lWorkItems;
    int inPlaceDone = -1;

    cl_int isInPlace = ((data_in_real == data_out_real) && (data_in_imag == data_out_imag)) ? 1 : 0;

    if((err = allocateTemporaryBufferPlannar(plan, batchSize)) != CL_SUCCESS)
        return err;

    cl_mem memObj_real[3];
    cl_mem memObj_imag[3];
    memObj_real[0] = data_in_real;
    memObj_real[1] = data_out_real;
    memObj_real[2] = plan->tempmemobj_real;
    memObj_imag[0] = data_in_imag;
    memObj_imag[1] = data_out_imag;
    memObj_imag[2] = plan->tempmemobj_imag;

    cl_fft_kernel_info *kernelInfo = plan->kernel_info;
    int numKernels = plan->num_kernels;

    int numKernelsOdd = numKernels & 1;
    int currRead  = 0;
    int currWrite = 1;

    // at least one external dram shuffle (transpose) required
    if(plan->temp_buffer_needed)
    {
        // in-place transform
        if(isInPlace)
        {
            inPlaceDone = 0;
            currRead  = 1;
            currWrite = 2;
        }
        else
        {
            currWrite = (numKernels & 1) ? 1 : 2;
        }

        while(kernelInfo)
        {
            if( isInPlace && numKernelsOdd && !inPlaceDone && kernelInfo->in_place_possible)
            {
                currWrite = currRead;
                inPlaceDone = 1;
            }

            s = batchSize;
            getKernelWorkDimensions(plan, kernelInfo, &s, &gWorkItems, &lWorkItems);
            err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj_real[currRead]);
            err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj_imag[currRead]);
            err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_mem), &memObj_real[currWrite]);
            err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_mem), &memObj_imag[currWrite]);
            err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_int), &dir);
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_int), &s);
306
307
            err |= clSetKernelArg(kernelInfo->kernel, 6, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 7, sizeof(cl_mem), &(plan->cossin_LUT_d2));
Oliver Bock's avatar
Oliver Bock committed
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331

            err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL, NULL);
            if(err)
                return err;

            currRead  = (currWrite == 1) ? 1 : 2;
            currWrite = (currWrite == 1) ? 2 : 1;

            kernelInfo = kernelInfo->next;
        }
    }
    // no dram shuffle (transpose required) transform
    else {

        while(kernelInfo)
        {
            s = batchSize;
            getKernelWorkDimensions(plan, kernelInfo, &s, &gWorkItems, &lWorkItems);
            err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj_real[currRead]);
            err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj_imag[currRead]);
            err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_mem), &memObj_real[currWrite]);
            err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_mem), &memObj_imag[currWrite]);
            err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_int), &dir);
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_int), &s);
332
333
            err |= clSetKernelArg(kernelInfo->kernel, 6, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 7, sizeof(cl_mem), &(plan->cossin_LUT_d2));
Oliver Bock's avatar
Oliver Bock committed
334
335
336
337
338
339
340
341
342
343
344
345
346

            err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL, NULL);
            if(err)
                return err;

            currRead  = 1;
            currWrite = 1;

            kernelInfo = kernelInfo->next;
        }
    }

    return err;
347
348
}

Oliver Bock's avatar
Oliver Bock committed
349
350
cl_int
clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
Oliver Bock's avatar
Oliver Bock committed
351
                         size_t numRows, size_t numCols, size_t startRow, size_t rowsToProcess, clFFT_Direction dir)
352
{
Oliver Bock's avatar
Oliver Bock committed
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
    cl_fft_plan *plan = (cl_fft_plan *) Plan;

    unsigned int N = numRows*numCols;
    unsigned int nCols = numCols;
    unsigned int sRow = startRow;
    unsigned int rToProcess = rowsToProcess;
    int d = dir;
    int err = 0;

    cl_device_id device_id;
    err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, NULL);
    if(err)
        return err;

    size_t gSize;
    err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, NULL);
    if(err)
        return err;

    gSize = min(128, gSize);
    size_t numGlobalThreads[1] = { max(numCols / gSize, 1)*gSize };
    size_t numLocalThreads[1]  = { gSize };

    err |= clSetKernelArg(plan->twist_kernel, 0, sizeof(cl_mem), &array);
    err |= clSetKernelArg(plan->twist_kernel, 1, sizeof(unsigned int), &sRow);
    err |= clSetKernelArg(plan->twist_kernel, 2, sizeof(unsigned int), &nCols);
    err |= clSetKernelArg(plan->twist_kernel, 3, sizeof(unsigned int), &N);
    err |= clSetKernelArg(plan->twist_kernel, 4, sizeof(unsigned int), &rToProcess);
    err |= clSetKernelArg(plan->twist_kernel, 5, sizeof(int), &d);

    err |= clEnqueueNDRangeKernel(queue, plan->twist_kernel, 1, NULL, numGlobalThreads, numLocalThreads, 0, NULL, NULL);

    return err;
386
387
}

Oliver Bock's avatar
Oliver Bock committed
388
389
cl_int
clFFT_1DTwistPlannar(clFFT_Plan Plan, cl_command_queue queue, cl_mem array_real, cl_mem array_imag,
Oliver Bock's avatar
Oliver Bock committed
390
                     size_t numRows, size_t numCols, size_t startRow, size_t rowsToProcess, clFFT_Direction dir)
391
{
Oliver Bock's avatar
Oliver Bock committed
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
    cl_fft_plan *plan = (cl_fft_plan *) Plan;

    unsigned int N = numRows*numCols;
    unsigned int nCols = numCols;
    unsigned int sRow = startRow;
    unsigned int rToProcess = rowsToProcess;
    int d = dir;
    int err = 0;

    cl_device_id device_id;
    err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, NULL);
    if(err)
        return err;

    size_t gSize;
    err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, NULL);
    if(err)
        return err;

    gSize = min(128, gSize);
    size_t numGlobalThreads[1] = { max(numCols / gSize, 1)*gSize };
    size_t numLocalThreads[1]  = { gSize };

    err |= clSetKernelArg(plan->twist_kernel, 0, sizeof(cl_mem), &array_real);
    err |= clSetKernelArg(plan->twist_kernel, 1, sizeof(cl_mem), &array_imag);
    err |= clSetKernelArg(plan->twist_kernel, 2, sizeof(unsigned int), &sRow);
    err |= clSetKernelArg(plan->twist_kernel, 3, sizeof(unsigned int), &nCols);
    err |= clSetKernelArg(plan->twist_kernel, 4, sizeof(unsigned int), &N);
    err |= clSetKernelArg(plan->twist_kernel, 5, sizeof(unsigned int), &rToProcess);
    err |= clSetKernelArg(plan->twist_kernel, 6, sizeof(int), &d);

    err |= clEnqueueNDRangeKernel(queue, plan->twist_kernel, 1, NULL, numGlobalThreads, numLocalThreads, 0, NULL, NULL);

    return err;
426
427
}