Skip to content
GitLab
Commit bb903c74 authored by Gaurav Khanna's avatar Gaurav Khanna Committed by Oliver Bock
Browse files

Initial version 

* Based on Apple's OpenCL FFT implementation
* Derivative work done by Gaurav Khanna
parents
Expand all Hide whitespace changes
Inline Side-by-side
build 0 → 100755
View file @ bb903c74
#! /bin/csh -f
set NV = ~/NVIDIA_GPU_Computing_SDK
g++ -g -o fft -I$NV/OpenCL/common/inc\
main.cpp fft_setup.cpp fft_execute.cpp fft_kernelstring.cpp\
-lOpenCL
clFFT.h 0 → 100644
View file @ bb903c74
#ifndef __CLFFT_H
#define __CLFFT_H
#ifdef __cplusplus
extern "C" {
#endif
#include <CL/cl.h>
#include <stdio.h>
// XForm type
typedef enum
{
clFFT_Forward = -1,
clFFT_Inverse = 1
}clFFT_Direction;
// XForm dimension
typedef enum
{
clFFT_1D = 0,
clFFT_2D = 1,
clFFT_3D = 3
}clFFT_Dimension;
// XForm Data type
typedef enum
{
clFFT_SplitComplexFormat = 0,
clFFT_InterleavedComplexFormat = 1
}clFFT_DataFormat;
typedef struct
{
unsigned int x;
unsigned int y;
unsigned int z;
}clFFT_Dim3;
typedef struct
{
float *real;
float *imag;
} clFFT_SplitComplex;
typedef struct
{
float real;
float imag;
}clFFT_Complex;
typedef void* clFFT_Plan;
clFFT_Plan clFFT_CreatePlan( cl_context context, clFFT_Dim3 n, clFFT_Dimension dim, clFFT_DataFormat dataFormat, cl_int *error_code );
void clFFT_DestroyPlan( clFFT_Plan plan );
cl_int clFFT_ExecuteInterleaved( cl_command_queue queue, clFFT_Plan plan, cl_int batchSize, clFFT_Direction dir,
cl_mem data_in, cl_mem data_out,
cl_int num_events, cl_event *event_list, cl_event *event );
cl_int clFFT_ExecutePlannar( cl_command_queue queue, clFFT_Plan plan, cl_int batchSize, clFFT_Direction dir,
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 );
cl_int clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
size_t numRows, size_t numCols, size_t startRow, size_t rowsToProcess, clFFT_Direction dir);
cl_int clFFT_1DTwistPlannar(clFFT_Plan Plan, cl_command_queue queue, cl_mem array_real, cl_mem array_imag,
size_t numRows, size_t numCols, size_t startRow, size_t rowsToProcess, clFFT_Direction dir);
void clFFT_DumpPlan( clFFT_Plan plan, FILE *file);
#ifdef __cplusplus
}
#endif
#endif
fft_base_kernels.h 0 → 100644
View file @ bb903c74
#ifndef __CL_FFT_BASE_KERNELS_
#define __CL_FFT_BASE_KERNELS_
#include <string.h>
using namespace std;
static string baseKernels = string
(
"#ifndef M_PI\n"
"#define M_PI 0x1.921fb54442d18p+1\n"
"#endif\n"
"#define complexMul(a,b) ((float2)(mad(-(a).y, (b).y, (a).x * (b).x), mad((a).y, (b).x, (a).x * (b).y)))\n"
"#define conj(a) ((float2)((a).x, -(a).y))\n"
"#define conjTransp(a) ((float2)(-(a).y, (a).x))\n"
"\n"
"#define fftKernel2(a,dir) \\\n"
"{ \\\n"
" float2 c = (a)[0]; \\\n"
" (a)[0] = c + (a)[1]; \\\n"
" (a)[1] = c - (a)[1]; \\\n"
"}\n"
"\n"
"#define fftKernel2S(d1,d2,dir) \\\n"
"{ \\\n"
" float2 c = (d1); \\\n"
" (d1) = c + (d2); \\\n"
" (d2) = c - (d2); \\\n"
"}\n"
"\n"
"#define fftKernel4(a,dir) \\\n"
"{ \\\n"
" fftKernel2S((a)[0], (a)[2], dir); \\\n"
" fftKernel2S((a)[1], (a)[3], dir); \\\n"
" fftKernel2S((a)[0], (a)[1], dir); \\\n"
" (a)[3] = (float2)(dir)*(conjTransp((a)[3])); \\\n"
" fftKernel2S((a)[2], (a)[3], dir); \\\n"
" float2 c = (a)[1]; \\\n"
" (a)[1] = (a)[2]; \\\n"
" (a)[2] = c; \\\n"
"}\n"
"\n"
"#define fftKernel4s(a0,a1,a2,a3,dir) \\\n"
"{ \\\n"
" fftKernel2S((a0), (a2), dir); \\\n"
" fftKernel2S((a1), (a3), dir); \\\n"
" fftKernel2S((a0), (a1), dir); \\\n"
" (a3) = (float2)(dir)*(conjTransp((a3))); \\\n"
" fftKernel2S((a2), (a3), dir); \\\n"
" float2 c = (a1); \\\n"
" (a1) = (a2); \\\n"
" (a2) = c; \\\n"
"}\n"
"\n"
"#define bitreverse8(a) \\\n"
"{ \\\n"
" float2 c; \\\n"
" c = (a)[1]; \\\n"
" (a)[1] = (a)[4]; \\\n"
" (a)[4] = c; \\\n"
" c = (a)[3]; \\\n"
" (a)[3] = (a)[6]; \\\n"
" (a)[6] = c; \\\n"
"}\n"
"\n"
"#define fftKernel8(a,dir) \\\n"
"{ \\\n"
" const float2 w1 = (float2)(0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f); \\\n"
" const float2 w3 = (float2)(-0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f); \\\n"
" float2 c; \\\n"
" fftKernel2S((a)[0], (a)[4], dir); \\\n"
" fftKernel2S((a)[1], (a)[5], dir); \\\n"
" fftKernel2S((a)[2], (a)[6], dir); \\\n"
" fftKernel2S((a)[3], (a)[7], dir); \\\n"
" (a)[5] = complexMul(w1, (a)[5]); \\\n"
" (a)[6] = (float2)(dir)*(conjTransp((a)[6])); \\\n"
" (a)[7] = complexMul(w3, (a)[7]); \\\n"
" fftKernel2S((a)[0], (a)[2], dir); \\\n"
" fftKernel2S((a)[1], (a)[3], dir); \\\n"
" fftKernel2S((a)[4], (a)[6], dir); \\\n"
" fftKernel2S((a)[5], (a)[7], dir); \\\n"
" (a)[3] = (float2)(dir)*(conjTransp((a)[3])); \\\n"
" (a)[7] = (float2)(dir)*(conjTransp((a)[7])); \\\n"
" fftKernel2S((a)[0], (a)[1], dir); \\\n"
" fftKernel2S((a)[2], (a)[3], dir); \\\n"
" fftKernel2S((a)[4], (a)[5], dir); \\\n"
" fftKernel2S((a)[6], (a)[7], dir); \\\n"
" bitreverse8((a)); \\\n"
"}\n"
"\n"
"#define bitreverse4x4(a) \\\n"
"{ \\\n"
" float2 c; \\\n"
" c = (a)[1]; (a)[1] = (a)[4]; (a)[4] = c; \\\n"
" c = (a)[2]; (a)[2] = (a)[8]; (a)[8] = c; \\\n"
" c = (a)[3]; (a)[3] = (a)[12]; (a)[12] = c; \\\n"
" c = (a)[6]; (a)[6] = (a)[9]; (a)[9] = c; \\\n"
" c = (a)[7]; (a)[7] = (a)[13]; (a)[13] = c; \\\n"
" c = (a)[11]; (a)[11] = (a)[14]; (a)[14] = c; \\\n"
"}\n"
"\n"
"#define fftKernel16(a,dir) \\\n"
"{ \\\n"
" const float w0 = 0x1.d906bcp-1f; \\\n"
" const float w1 = 0x1.87de2ap-2f; \\\n"
" const float w2 = 0x1.6a09e6p-1f; \\\n"
" fftKernel4s((a)[0], (a)[4], (a)[8], (a)[12], dir); \\\n"
" fftKernel4s((a)[1], (a)[5], (a)[9], (a)[13], dir); \\\n"
" fftKernel4s((a)[2], (a)[6], (a)[10], (a)[14], dir); \\\n"
" fftKernel4s((a)[3], (a)[7], (a)[11], (a)[15], dir); \\\n"
" (a)[5] = complexMul((a)[5], (float2)(w0, dir*w1)); \\\n"
" (a)[6] = complexMul((a)[6], (float2)(w2, dir*w2)); \\\n"
" (a)[7] = complexMul((a)[7], (float2)(w1, dir*w0)); \\\n"
" (a)[9] = complexMul((a)[9], (float2)(w2, dir*w2)); \\\n"
" (a)[10] = (float2)(dir)*(conjTransp((a)[10])); \\\n"
" (a)[11] = complexMul((a)[11], (float2)(-w2, dir*w2)); \\\n"
" (a)[13] = complexMul((a)[13], (float2)(w1, dir*w0)); \\\n"
" (a)[14] = complexMul((a)[14], (float2)(-w2, dir*w2)); \\\n"
" (a)[15] = complexMul((a)[15], (float2)(-w0, dir*-w1)); \\\n"
" fftKernel4((a), dir); \\\n"
" fftKernel4((a) + 4, dir); \\\n"
" fftKernel4((a) + 8, dir); \\\n"
" fftKernel4((a) + 12, dir); \\\n"
" bitreverse4x4((a)); \\\n"
"}\n"
"\n"
"#define bitreverse32(a) \\\n"
"{ \\\n"
" float2 c1, c2; \\\n"
" 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"
" 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"
" 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"
" 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"
" 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"
" 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"
"}\n"
"\n"
"#define fftKernel32(a,dir) \\\n"
"{ \\\n"
" fftKernel2S((a)[0], (a)[16], dir); \\\n"
" fftKernel2S((a)[1], (a)[17], dir); \\\n"
" fftKernel2S((a)[2], (a)[18], dir); \\\n"
" fftKernel2S((a)[3], (a)[19], dir); \\\n"
" fftKernel2S((a)[4], (a)[20], dir); \\\n"
" fftKernel2S((a)[5], (a)[21], dir); \\\n"
" fftKernel2S((a)[6], (a)[22], dir); \\\n"
" fftKernel2S((a)[7], (a)[23], dir); \\\n"
" fftKernel2S((a)[8], (a)[24], dir); \\\n"
" fftKernel2S((a)[9], (a)[25], dir); \\\n"
" fftKernel2S((a)[10], (a)[26], dir); \\\n"
" fftKernel2S((a)[11], (a)[27], dir); \\\n"
" fftKernel2S((a)[12], (a)[28], dir); \\\n"
" fftKernel2S((a)[13], (a)[29], dir); \\\n"
" fftKernel2S((a)[14], (a)[30], dir); \\\n"
" fftKernel2S((a)[15], (a)[31], dir); \\\n"
" (a)[17] = complexMul((a)[17], (float2)(0x1.f6297cp-1f, dir*0x1.8f8b84p-3f)); \\\n"
" (a)[18] = complexMul((a)[18], (float2)(0x1.d906bcp-1f, dir*0x1.87de2ap-2f)); \\\n"
" (a)[19] = complexMul((a)[19], (float2)(0x1.a9b662p-1f, dir*0x1.1c73b4p-1f)); \\\n"
" (a)[20] = complexMul((a)[20], (float2)(0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f)); \\\n"
" (a)[21] = complexMul((a)[21], (float2)(0x1.1c73b4p-1f, dir*0x1.a9b662p-1f)); \\\n"
" (a)[22] = complexMul((a)[22], (float2)(0x1.87de2ap-2f, dir*0x1.d906bcp-1f)); \\\n"
" (a)[23] = complexMul((a)[23], (float2)(0x1.8f8b84p-3f, dir*0x1.f6297cp-1f)); \\\n"
" (a)[24] = complexMul((a)[24], (float2)(0x0p+0f, dir*0x1p+0f)); \\\n"
" (a)[25] = complexMul((a)[25], (float2)(-0x1.8f8b84p-3f, dir*0x1.f6297cp-1f)); \\\n"
" (a)[26] = complexMul((a)[26], (float2)(-0x1.87de2ap-2f, dir*0x1.d906bcp-1f)); \\\n"
" (a)[27] = complexMul((a)[27], (float2)(-0x1.1c73b4p-1f, dir*0x1.a9b662p-1f)); \\\n"
" (a)[28] = complexMul((a)[28], (float2)(-0x1.6a09e6p-1f, dir*0x1.6a09e6p-1f)); \\\n"
" (a)[29] = complexMul((a)[29], (float2)(-0x1.a9b662p-1f, dir*0x1.1c73b4p-1f)); \\\n"
" (a)[30] = complexMul((a)[30], (float2)(-0x1.d906bcp-1f, dir*0x1.87de2ap-2f)); \\\n"
" (a)[31] = complexMul((a)[31], (float2)(-0x1.f6297cp-1f, dir*0x1.8f8b84p-3f)); \\\n"
" fftKernel16((a), dir); \\\n"
" fftKernel16((a) + 16, dir); \\\n"
" bitreverse32((a)); \\\n"
"}\n\n"
);
static string twistKernelInterleaved = string
(
"__kernel void \\\n"
"clFFT_1DTwistInterleaved(__global float2 *in, unsigned int startRow, unsigned int numCols, unsigned int N, unsigned int numRowsToProcess, int dir) \\\n"
"{ \\\n"
" float2 a, w; \\\n"
" float ang; \\\n"
" unsigned int j; \\\n"
" unsigned int i = get_global_id(0); \\\n"
" unsigned int startIndex = i; \\\n"
" \\\n"
" if(i < numCols) \\\n"
" { \\\n"
" for(j = 0; j < numRowsToProcess; j++) \\\n"
" { \\\n"
" a = in[startIndex]; \\\n"
" ang = 2.0f * M_PI * dir * i * (startRow + j) / N; \\\n"
" w = (float2)(native_cos(ang), native_sin(ang)); \\\n"
" a = complexMul(a, w); \\\n"
" in[startIndex] = a; \\\n"
" startIndex += numCols; \\\n"
" } \\\n"
" } \\\n"
"} \\\n"
);
static string twistKernelPlannar = string
(
"__kernel void \\\n"
"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"
"{ \\\n"
" float2 a, w; \\\n"
" float ang; \\\n"
" unsigned int j; \\\n"
" unsigned int i = get_global_id(0); \\\n"
" unsigned int startIndex = i; \\\n"
" \\\n"
" if(i < numCols) \\\n"
" { \\\n"
" for(j = 0; j < numRowsToProcess; j++) \\\n"
" { \\\n"
" a = (float2)(in_real[startIndex], in_imag[startIndex]); \\\n"
" ang = 2.0f * M_PI * dir * i * (startRow + j) / N; \\\n"
" w = (float2)(native_cos(ang), native_sin(ang)); \\\n"
" a = complexMul(a, w); \\\n"
" in_real[startIndex] = a.x; \\\n"
" in_imag[startIndex] = a.y; \\\n"
" startIndex += numCols; \\\n"
" } \\\n"
" } \\\n"
"} \\\n"
);
#endif
fft_execute.cpp 0 → 100644
View file @ bb903c74
#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)
{
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;
}
static cl_int
allocateTemporaryBufferPlannar(cl_fft_plan *plan, cl_uint batchSize)
{
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;
}
void
getKernelWorkDimensions(cl_fft_plan *plan, cl_fft_kernel_info *kernelInfo, cl_int *batchSize, size_t *gWorkItems, size_t *lWorkItems)
{
*lWorkItems = kernelInfo->num_workitems_per_workgroup;
int numWorkGroups = kernelInfo->num_workgroups;
int numXFormsPerWG = kernelInfo->num_xforms_per_workgroup;
switch(kernelInfo->dir)
{
case cl_fft_kernel_x:
*batchSize *= (plan->n.y * plan->n.z);
numWorkGroups = (*batchSize % numXFormsPerWG) ? (*batchSize/numXFormsPerWG + 1) : (*batchSize/numXFormsPerWG);
numWorkGroups *= kernelInfo->num_workgroups;
break;
case cl_fft_kernel_y:
*batchSize *= plan->n.z;
numWorkGroups *= *batchSize;
break;
case cl_fft_kernel_z:
numWorkGroups *= *batchSize;
break;
}
*gWorkItems = numWorkGroups * *lWorkItems;
}
cl_int
clFFT_ExecuteInterleaved( cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
cl_mem data_in, cl_mem data_out,
cl_int num_events, cl_event *event_list, cl_event *event )
{
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;
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);
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);
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;
}
cl_int
clFFT_ExecutePlannar( cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
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)
{
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;
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);
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);
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;
}
cl_int
clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
size_t numRows, size_t numCols, size_t startRow, size_t rowsToProcess, clFFT_Direction dir)
{
puts ("X");
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 };