Commit cb388aca authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed bugs in CPU version of FFT

parent 52a6a366
...@@ -210,8 +210,7 @@ cl::Kernel OpenCLFFT3D::createKernel(int xsize, int ysize, int zsize) { ...@@ -210,8 +210,7 @@ cl::Kernel OpenCLFFT3D::createKernel(int xsize, int ysize, int zsize) {
replacements["ZSIZE"] = OpenCLExpressionUtilities::intToString(zsize); replacements["ZSIZE"] = OpenCLExpressionUtilities::intToString(zsize);
replacements["M_PI"] = OpenCLExpressionUtilities::doubleToString(M_PI); replacements["M_PI"] = OpenCLExpressionUtilities::doubleToString(M_PI);
replacements["COMPUTE_FFT"] = source.str(); replacements["COMPUTE_FFT"] = source.str();
if (loopRequired) replacements["LOOP_REQUIRED"] = (loopRequired ? "1" : "0");
replacements["LOOP_REQUIRED"] = "1";
cl::Program program = context.createProgram(context.replaceStrings(OpenCLKernelSources::fft, replacements)); cl::Program program = context.createProgram(context.replaceStrings(OpenCLKernelSources::fft, replacements));
cl::Kernel kernel(program, "execFFT"); cl::Kernel kernel(program, "execFFT");
kernel.setArg(3, zsize*sizeof(mm_float2), NULL); kernel.setArg(3, zsize*sizeof(mm_float2), NULL);
......
...@@ -14,7 +14,7 @@ __kernel void execFFT(__global const float2* restrict in, __global float2* restr ...@@ -14,7 +14,7 @@ __kernel void execFFT(__global const float2* restrict in, __global float2* restr
for (int index = get_group_id(0); index < XSIZE*YSIZE; index += get_num_groups(0)) { for (int index = get_group_id(0); index < XSIZE*YSIZE; index += get_num_groups(0)) {
int x = index/YSIZE; int x = index/YSIZE;
int y = index-x*YSIZE; int y = index-x*YSIZE;
#ifdef LOOP_REQUIRED #if LOOP_REQUIRED
for (int z = get_local_id(0); z < ZSIZE; z += get_local_size(0)) for (int z = get_local_id(0); z < ZSIZE; z += get_local_size(0))
data0[z] = in[x*(YSIZE*ZSIZE)+y*ZSIZE+z]; data0[z] = in[x*(YSIZE*ZSIZE)+y*ZSIZE+z];
#else #else
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment