Bug fixes to real-to-complex FFT

platforms/opencl/src/OpenCLFFT3D.cpp

@@ -313,16 +313,24 @@ cl::Kernel OpenCLFFT3D::createKernel(int xsize, int ysize, int zsize, int& threa
         bool outputIsPacked = (inputIsReal && axis == 2 && forward);
         string outputSuffix = (outputIsReal ? ".x" : "");
         if (loopRequired) {
+            if (outputIsPacked)
+                source<<"if (x < XSIZE/2+1)\n";
             source<<"for (int z = get_local_id(0); z < ZSIZE; z += get_local_size(0))\n";
+            if (outputIsPacked)
+                source<<"out[y*(ZSIZE*(XSIZE/2+1))+z*(XSIZE/2+1)+x] = data"<<(stage%2)<<"[z]"<<outputSuffix<<";\n";
+            else
                 source<<"out[y*(ZSIZE*XSIZE)+z*XSIZE+x] = data"<<(stage%2)<<"[z]"<<outputSuffix<<";\n";
         }
         else {
-            source<<"if (index < XSIZE*YSIZE)\n";
-            if (outputIsPacked)
+            if (outputIsPacked) {
+                source<<"if (index < XSIZE*YSIZE && x < XSIZE/2+1)\n";
                 source<<"out[y*(ZSIZE*(XSIZE/2+1))+(get_local_id(0)%ZSIZE)*(XSIZE/2+1)+x] = data"<<(stage%2)<<"[get_local_id(0)]"<<outputSuffix<<";\n";
-            else
+            }
+            else {
+                source<<"if (index < XSIZE*YSIZE)\n";
                 source<<"out[y*(ZSIZE*XSIZE)+(get_local_id(0)%ZSIZE)*XSIZE+x] = data"<<(stage%2)<<"[get_local_id(0)]"<<outputSuffix<<";\n";
             }
+        }
         map<string, string> replacements;
         replacements["XSIZE"] = context.intToString(xsize);
         replacements["YSIZE"] = context.intToString(ysize);

platforms/opencl/src/kernels/fft.cl

@@ -30,8 +30,7 @@ __kernel void execFFT(__global const INPUT_TYPE* restrict in, __global OUTPUT_TY
         int x = index/YSIZE;
         int y = index-x*YSIZE;
 #if OUTPUT_IS_PACKED
-        if (x >= XSIZE/2+1)
-            continue;
+        if (x < XSIZE/2+1) {
 #endif
 #if LOOP_REQUIRED
         for (int z = get_local_id(0); z < ZSIZE; z += get_local_size(0))
@@ -53,6 +52,9 @@ __kernel void execFFT(__global const INPUT_TYPE* restrict in, __global OUTPUT_TY
     #endif
 #endif
         barrier(CLK_LOCAL_MEM_FENCE);
+#if OUTPUT_IS_PACKED
+        }
+#endif
         COMPUTE_FFT
     }
 }