No commit message

No commit message

samples/convolution/conv.bat

@@ -2,5 +2,5 @@ java ^
  -Djava.library.path=../../com.amd.aparapi.jni/dist ^
  -Dcom.amd.aparapi.executionMode=%1 ^
  -classpath ../../com.amd.aparapi/dist/aparapi.jar;convolution.jar ^
- com.amd.aparapi.sample.convolution.Main
+ com.amd.aparapi.sample.convolution.Convolution
 

samples/convolution/pureJava.bat

+java ^
+ -Djava.library.path=../../com.amd.aparapi.jni/dist ^
+ -Dcom.amd.aparapi.executionMode=%1 ^
+ -classpath ../../com.amd.aparapi/dist/aparapi.jar;convolution.jar ^
+ com.amd.aparapi.sample.convolution.PureJava
+

samples/convolution/src/com/amd/aparapi/sample/convolution/ConvMatrix3x3Editor.java

+
+/*
+Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+following conditions are met:
+
+Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+disclaimer. 
+
+Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+disclaimer in the documentation and/or other materials provided with the distribution. 
+
+Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+derived from this software without specific prior written permission. 
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, 
+WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730 through
+774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000.  Further, pursuant to Section 740.6 of the EAR,
+you hereby certify that, except pursuant to a license granted by the United States Department of Commerce Bureau of 
+Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export Administration 
+Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in Country Groups D:1,
+E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2) export to Country Groups
+D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced direct product is subject
+to national security controls as identified on the Commerce Control List (currently found in Supplement 1 to Part 774
+of EAR).  For the most current Country Group listings, or for additional information about the EAR or your obligations
+under those regulations, please refer to the U.S. Bureau of Industry and Security's website at http://www.bis.doc.gov/. 
+
+*/
+
+package com.amd.aparapi.sample.convolution;
+
+import java.awt.BorderLayout;
+import java.awt.Component;
+import java.awt.GridLayout;
+import java.awt.event.ActionEvent;
+import java.awt.event.ActionListener;
+import java.util.Arrays;
+
+import javax.swing.BoxLayout;
+import javax.swing.JComboBox;
+import javax.swing.JPanel;
+import javax.swing.JSpinner;
+import javax.swing.SpinnerModel;
+import javax.swing.SpinnerNumberModel;
+import javax.swing.event.ChangeEvent;
+import javax.swing.event.ChangeListener;
+
+public class ConvMatrix3x3Editor{
+   Component component;
+
+   float[] default3x3;
+
+   float[] none3x3 = new float[] {
+         0,
+         0,
+         0,
+         0,
+         1,
+         0,
+         0,
+         0,
+         0
+   };
+
+   float[] blur3x3 = new float[] {
+         .1f,
+         .1f,
+         .1f,
+         .1f,
+         .1f,
+         .1f,
+         .1f,
+         .1f,
+         .1f
+   };
+
+   JSpinner[] spinners = new JSpinner[9];
+
+   protected void updated(float[] _convMatrix3x3) {
+
+   };
+
+   void set(float[] _to, float[] _from) {
+      for (int i = 0; i < 9; i++) {
+         _to[i] = _from[i];
+         spinners[i].setValue((Double) (double) _to[i]);
+
+      }
+      updated(_to);
+   }
+
+   ConvMatrix3x3Editor(final float[] _convMatrix3x3) {
+      default3x3 = Arrays.copyOf(_convMatrix3x3, _convMatrix3x3.length);
+      JPanel leftPanel = new JPanel();
+      JPanel controlPanel = new JPanel();
+      BoxLayout layout = new BoxLayout(controlPanel, BoxLayout.Y_AXIS);
+      controlPanel.setLayout(layout);
+      component = leftPanel;
+      JPanel grid3x3Panel = new JPanel();
+      controlPanel.add(grid3x3Panel);
+      grid3x3Panel.setLayout(new GridLayout(3, 3));
+      for (int i = 0; i < 9; i++) {
+         final int index = i;
+         SpinnerModel model = new SpinnerNumberModel(_convMatrix3x3[index], -50f, 50f, 1f);
+         JSpinner spinner = new JSpinner(model);
+         spinners[i] = spinner;
+         spinner.addChangeListener(new ChangeListener(){
+            public void stateChanged(ChangeEvent ce) {
+               JSpinner source = (JSpinner) ce.getSource();
+               double value = ((Double) source.getValue());
+               _convMatrix3x3[index] = (float) value;
+               updated(_convMatrix3x3);
+            }
+         });
+         grid3x3Panel.add(spinner);
+      }
+      String[] options = new String[] {
+            "DEFAULT",
+            "NONE",
+            "BLUR"
+      };
+      JComboBox combo = new JComboBox(options);
+      combo.addActionListener(new ActionListener(){
+
+         @Override public void actionPerformed(ActionEvent e) {
+            JComboBox cb = (JComboBox) e.getSource();
+            String value = (String) cb.getSelectedItem();
+            if (value.equals("DEFAULT")) {
+               set(_convMatrix3x3, default3x3);
+            } else if (value.equals("NONE")) {
+               set(_convMatrix3x3, none3x3);
+            } else if (value.equals("BLUR")) {
+               set(_convMatrix3x3, blur3x3);
+            }
+         }
+
+      });
+      controlPanel.add(combo);
+
+      leftPanel.add(controlPanel, BorderLayout.NORTH);
+   }
+}

samples/convolution/src/com/amd/aparapi/sample/convolution/Main.java → samples/convolution/src/com/amd/aparapi/sample/convolution/Convolution.java

@@ -38,197 +38,80 @@ under those regulations, please refer to the U.S. Bureau of Industry and Securit
 
 package com.amd.aparapi.sample.convolution;
 
-import java.awt.Dimension;
-import java.awt.Graphics;
-import java.awt.image.BufferedImage;
-import java.awt.image.DataBufferInt;
 import java.io.File;
-import java.io.IOException;
-
-import javax.imageio.ImageIO;
-import javax.swing.JComponent;
-import javax.swing.JFrame;
-import javax.swing.WindowConstants;
 
 import com.amd.aparapi.Kernel;
-import com.amd.aparapi.Range;
-
-/**
- * An example Aparapi application which demonstrates image manipulation via convolution filter
- * 
- * Converted to use int buffer and some performance tweaks by Gary Frost
- * http://processing.org/learning/pixels/
- * 
- * @author Gary Frost
- */
-public class Main{
-   // http://docs.gimp.org/en/plug-in-convmatrix.html
-
-   final static class ConvolutionFilter{
-      private float[] weights;
-
-      private int offset;
-
-      ConvolutionFilter(float _nw, float _n, float ne, float _w, float _o, float _e, float _sw, float _s, float _se, int _offset) {
-         weights = new float[] {
-               _nw,
-               _w,
-               ne,
-               _w,
-               _o,
-               _e,
-               _sw,
-               _s,
-               _se
-         };
-         offset = _offset;
-      }
-
-   }
-
-   private static final ConvolutionFilter NONE = new ConvolutionFilter(0f, 0f, 0f, 0f, 1f, 0f, 0f, 0f, 0f, 0);
 
-   private static final ConvolutionFilter BLUR = new ConvolutionFilter(1f, 1f, 1f, 1f, 1f, 1f, 1f, 1f, 1f, 0);
+public class Convolution{
 
-   private static final ConvolutionFilter EMBOSS = new ConvolutionFilter(-2f, -1f, 0f, -1f, 1f, 1f, 0f, 1f, 2f, 0);
+   final static class ImageConvolution extends Kernel{
+      private float convMatrix3x3[];
 
-   public static class ConvolutionKernel extends Kernel{
+      private int width, height;
 
-      private final float[] filter = new float[9];
+      private byte imageIn[], imageOut[];
 
-      private final int[] inputData;
+      public void processPixel(int x, int y, int w, int h) {
+         float accum = 0f;
+         int count = 0;
+         for (int dx = -3; dx < 6; dx += 3) {
+            for (int dy = -1; dy < 2; dy += 1) {
+               int rgb = 0xff & imageIn[((y + dy) * w) + (x + dx)];
 
-      private final int[] outputData;
-
-      private final int width;
-
-      private final int height;
-
-      private int offset;
-
-      public ConvolutionKernel(int _width, int _height, BufferedImage _inputImage, BufferedImage _outputImage) {
-         inputData = ((DataBufferInt) _inputImage.getRaster().getDataBuffer()).getData();
-         outputData = ((DataBufferInt) _outputImage.getRaster().getDataBuffer()).getData();
-         width = _width;
-         height = _height;
+               accum += rgb * convMatrix3x3[count++];
+            }
+         }
+         byte value = (byte) (max(0, min((int) accum, 255)));
+         imageOut[y * w + x] = value;
 
-         // setExplicit(true); // This gives us a performance boost
-         //  put(inputData); // Because we are using explicit buffer management we must put the imageData array
       }
 
       @Override public void run() {
+         int x = getGlobalId(0) % (width * 3);
+         int y = getGlobalId(0) / (width * 3);
 
-         int x = getGlobalId(0);
-         int y = getGlobalId(1);
-         int lx = getLocalId(0);
-         int ly = getLocalId(1);
-         int w = getGlobalSize(0);
-         int h = getGlobalSize(1);
-         // System.out.println(x+","+y+" "+lx+","+ly+" "+w+","+h);
-         if (x > 1 && x < (w - 1) && y > 1 && y < (h - 1)) {
-
-            int result = 0;
-            // We handle each color separately using rgbshift as an 8 bit mask for red, green, blue
-            for (int rgbShift = 0; rgbShift < 24; rgbShift += 8) { // 0,8,16
-               int channelAccum = 0;
-               float accum = 0;
-
-               for (int count = 0; count < 9; count++) {
-                  int dx = (count % 3) - 1; // 0,1,2 -> -1,0,1
-                  int dy = (count / 3) - 1; // 0,1,2 -> -1,0,1
-
-                  int rgb = (inputData[((y + dy) * w) + (x + dx)]);
-                  int channelValue = ((rgb >> rgbShift) & 0xff);
-                  accum += filter[count];
-                  channelAccum += channelValue * filter[count++];
-
-               }
-               channelAccum /= accum;
-               channelAccum += offset;
-               channelAccum = max(0, min(channelAccum, 0xff));
-               result |= (channelAccum << rgbShift);
-            }
-            outputData[y * w + x] = result;
+         if (x > 3 && x < (width * 3 - 3) && y > 1 && y < (height - 1)) {
+            processPixel(x, y, width * 3, height);
          }
-      }
 
-      public void convolve(ConvolutionFilter _filter) {
-         System.arraycopy(_filter.weights, 0, filter, 0, _filter.weights.length);
-         offset = _filter.offset;
-         put(filter);
-         execute(Range.create2D(width, height, 8, 8));
-         get(outputData);
       }
-   }
 
-   public static final int PAD = 1024;
+      public void applyConvolution(float[] _convMatrix3x3, byte[] _imageIn, byte[] _imageOut, int _width, int _height) {
+         imageIn = _imageIn;
+         imageOut = _imageOut;
+         width = _width;
+         height = _height;
+         convMatrix3x3 = _convMatrix3x3;
 
-   public static int padValue(int value) {
-      return (PAD - (value % PAD));
-   }
+         execute(3 * width * height);
+      }
 
-   public static int padTo(int value) {
-      return (value + padValue(value));
    }
 
-   public static void main(String[] _args) throws IOException, InterruptedException {
-
-      JFrame frame = new JFrame("Convolution");
-
-      BufferedImage testCard = ImageIO.read(new File("testcard.jpg"));
-
-      int imageHeight = testCard.getHeight();
-
-      int imageWidth = testCard.getWidth();
-
-      final int width = padTo(imageWidth);// now multiple of 64
-
-      final int height = padTo(imageHeight); // now multiple of 64
-
-      System.out.println("image width,height=" + width + "," + height);
-
-      final BufferedImage inputImage = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
-
-      inputImage.getGraphics().drawImage(testCard, padValue(imageWidth) / 2, padValue(imageHeight) / 2, null);
-      final BufferedImage outputImage = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
-      outputImage.getGraphics().drawImage(testCard, padValue(imageWidth) / 2, padValue(imageHeight) / 2, null);
-      final ConvolutionKernel lifeKernel = new ConvolutionKernel(width, height, inputImage, outputImage);
-
-      // Create a component for viewing the offsecreen image
-      @SuppressWarnings("serial") JComponent viewer = new JComponent(){
-         @Override public void paintComponent(Graphics g) {
-            //  if (lifeKernel.isExplicit()) {
-            //    lifeKernel.get(lifeKernel.inputData); // We only pull the imageData when we intend to use it.
-            //  }
-            g.drawImage(outputImage, 0, 0, width, height, 0, 0, width, height, this);
-         }
-      };
-
-      // Set the default size and add to the frames content pane
-      viewer.setPreferredSize(new Dimension(width, height));
-      frame.getContentPane().add(viewer);
-
-      // Swing housekeeping
-      frame.pack();
-      frame.setVisible(true);
-      frame.setDefaultCloseOperation(WindowConstants.EXIT_ON_CLOSE);
-
-      ConvolutionFilter filters[] = new ConvolutionFilter[] {
-            NONE,
-            BLUR,
-            EMBOSS,
+   public static void main(final String[] _args) {
+      File file = new File(_args.length == 1 ? _args[0] : "testcard.jpg");
+
+      final ImageConvolution convolution = new ImageConvolution();
+
+      float convMatrix3x3[] = new float[] {
+            0f,
+            -10f,
+            0f,
+            -10f,
+            40f,
+            -10f,
+            0f,
+            -10f,
+            0f,
       };
-      long start = System.nanoTime();
-      for (int i = 0; i < 100; i++) {
-         for (ConvolutionFilter filter : filters) {
 
-            lifeKernel.convolve(filter); // Work is performed here
-
-            viewer.repaint(); // Request a repaint of the viewer (causes paintComponent(Graphics) to be called later not synchronous
-            //Thread.sleep(1000);
+      new ConvolutionViewer(file, convMatrix3x3){
+         @Override protected void applyConvolution(float[] _convMatrix3x3, byte[] _inBytes, byte[] _outBytes, int _width,
+               int _height) {
+            convolution.applyConvolution(_convMatrix3x3, _inBytes, _outBytes, _width, _height);
          }
-      }
-      System.out.println((System.nanoTime() - start) / 1000000);
+      };
 
    }
+
 }

samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionOpenCL.java

+
+/*
+Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+following conditions are met:
+
+Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+disclaimer. 
+
+Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+disclaimer in the documentation and/or other materials provided with the distribution. 
+
+Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+derived from this software without specific prior written permission. 
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, 
+WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730 through
+774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000.  Further, pursuant to Section 740.6 of the EAR,
+you hereby certify that, except pursuant to a license granted by the United States Department of Commerce Bureau of 
+Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export Administration 
+Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in Country Groups D:1,
+E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2) export to Country Groups
+D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced direct product is subject
+to national security controls as identified on the Commerce Control List (currently found in Supplement 1 to Part 774
+of EAR).  For the most current Country Group listings, or for additional information about the EAR or your obligations
+under those regulations, please refer to the U.S. Bureau of Industry and Security's website at http://www.bis.doc.gov/. 
+
+*/
+
+package com.amd.aparapi.sample.convolution;
+
+import java.io.File;
+
+import com.amd.aparapi.Device;
+import com.amd.aparapi.OpenCL;
+import com.amd.aparapi.OpenCLDevice;
+import com.amd.aparapi.Range;
+
+public class ConvolutionOpenCL{
+
+   @OpenCL.Resource("com/amd/aparapi/sample/convolution/convolution.cl") interface Convolution extends OpenCL<Convolution>{
+      Convolution applyConvolution(//
+            Range range, //
+            @GlobalReadOnly("_convMatrix3x3") float[] _convMatrix3x3,//// only read from kernel 
+            @GlobalReadOnly("_imagIn") byte[] _imageIn,// only read from kernel (actually char[])
+            @GlobalWriteOnly("_imagOut") byte[] _imageOut, // only written to (never read) from kernel (actually char[])
+            @Arg("_width") int _width,// 
+            @Arg("_height") int _height);
+   }
+
+   public static void main(final String[] _args) {
+      File file = new File(_args.length == 1 ? _args[0] : "testcard.jpg");
+
+      final OpenCLDevice openclDevice = (OpenCLDevice) Device.best();
+
+      final Convolution convolution = openclDevice.bind(Convolution.class);
+      float convMatrix3x3[] = new float[] {
+            0f,
+            -10f,
+            0f,
+            -10f,
+            40f,
+            -10f,
+            0f,
+            -10f,
+            0f,
+      };
+
+      new ConvolutionViewer(file, convMatrix3x3){
+         Range range = null;
+
+         @Override protected void applyConvolution(float[] _convMatrix3x3, byte[] _inBytes, byte[] _outBytes, int _width,
+               int _height) {
+            if (range == null) {
+               range = openclDevice.createRange(_width * _height * 3);
+            }
+            convolution.applyConvolution(range, _convMatrix3x3, _inBytes, _outBytes, _width, _height);
+         }
+      };
+
+   }
+
+}

samples/convolution/src/com/amd/aparapi/sample/convolution/Convolve.java → samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionViewer.java

+
 /*
 Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
 All rights reserved.
@@ -38,189 +39,94 @@ under those regulations, please refer to the U.S. Bureau of Industry and Securit
 
 package com.amd.aparapi.sample.convolution;
 
-import java.awt.Dimension;
-import java.awt.Graphics;
+import java.awt.BorderLayout;
+import java.awt.Color;
+import java.awt.Graphics2D;
 import java.awt.image.BufferedImage;
-import java.awt.image.ConvolveOp;
-import java.awt.image.DataBufferInt;
+import java.awt.image.DataBufferByte;
 import java.io.File;
 import java.io.IOException;
 
 import javax.imageio.ImageIO;
-import javax.swing.JComponent;
+import javax.swing.ImageIcon;
 import javax.swing.JFrame;
+import javax.swing.JLabel;
 import javax.swing.WindowConstants;
 
-import com.amd.aparapi.Kernel;
-import com.amd.aparapi.Range;
-
-/**
- * An example Aparapi application which demonstrates image manipulation via convolution filter
- * 
- * Converted to use int buffer and some performance tweaks by Gary Frost
- * http://processing.org/learning/pixels/
- * 
- * @author Gary Frost
- */
-public class Convolve{
-   // http://docs.gimp.org/en/plug-in-convmatrix.html
-
-   final static class ConvolutionFilter{
-      private float[] weights;
-
-      private int offset;
-
-      ConvolutionFilter(float _nw, float _n, float ne, float _w, float _o, float _e, float _sw, float _s, float _se, int _offset) {
-         weights = new float[] {
-               _nw,
-               _w,
-               ne,
-               _w,
-               _o,
-               _e,
-               _sw,
-               _s,
-               _se
-         };
-         offset = _offset;
-      }
-
-   }
-
-   public static class ConvolutionKernel extends Kernel{
+@SuppressWarnings("serial") public abstract class ConvolutionViewer extends JFrame{
 
-      private final float[] filter = new float[9];
+   private int height;
 
-      private final int[] inputData;
+   private int width;
 
-      private final int[] outputData;
+   private BufferedImage outputImage;
 
-      private final int width;
-
-      private final int height;
-
-      private int offset;
-
-      public ConvolutionKernel(int _width, int _height, BufferedImage _inputImage, BufferedImage _outputImage) {
-         inputData = ((DataBufferInt) _inputImage.getRaster().getDataBuffer()).getData();
-         outputData = ((DataBufferInt) _outputImage.getRaster().getDataBuffer()).getData();
-         width = _width;
-         height = _height;
-
-         // setExplicit(true); // This gives us a performance boost
-         //  put(inputData); // Because we are using explicit buffer management we must put the imageData array
-      }
+   private BufferedImage inputImage;
 
-      @Override public void run() {
-
-         int x = getGlobalId(0);
-         int y = getGlobalId(1);
-         int lx = getLocalId(0);
-         int ly = getLocalId(1);
-         int w = getGlobalSize(0);
-         int h = getGlobalSize(1);
-         // System.out.println(x+","+y+" "+lx+","+ly+" "+w+","+h);
-         if (x > 1 && x < (w - 1) && y > 1 && y < (h - 1)) {
-
-            int result = 0;
-            // We handle each color separately using rgbshift as an 8 bit mask for red, green, blue
-            for (int rgbShift = 0; rgbShift < 24; rgbShift += 8) { // 0,8,16
-               int channelAccum = 0;
-               float accum = 0;
-
-               for (int count = 0; count < 9; count++) {
-                  int dx = (count % 3) - 1; // 0,1,2 -> -1,0,1
-                  int dy = (count / 3) - 1; // 0,1,2 -> -1,0,1
-
-                  int rgb = (inputData[((y + dy) * w) + (x + dx)]);
-                  int channelValue = ((rgb >> rgbShift) & 0xff);
-                  accum += filter[count];
-                  channelAccum += channelValue * filter[count++];
-
-               }
-               channelAccum /= accum;
-               channelAccum += offset;
-               channelAccum = max(0, min(channelAccum, 0xff));
-               result |= (channelAccum << rgbShift);
-            }
-            outputData[y * w + x] = result;
-         }
-      }
+   private byte[] inBytes;
 
-      public void convolve(ConvolutionFilter _filter) {
-         System.arraycopy(_filter.weights, 0, filter, 0, _filter.weights.length);
-         offset = _filter.offset;
-         put(filter);
-         execute(Range.create2D(width, height, 8, 8));
-         get(outputData);
-      }
-   }
+   private byte[] outBytes;
 
-   public static final int PAD = 1024;
+   private Graphics2D gc;
 
-   public static int padValue(int value) {
-      return (PAD - (value % PAD));
-   }
+   private float[] convMatrix3x3;
 
-   public static int padTo(int value) {
-      return (value + padValue(value));
-   }
+   public ConvolutionViewer(File _file, float[] _convMatrix3x3) {
 
-   public static void main(String[] _args) throws IOException, InterruptedException {
+      JFrame frame = new JFrame("Convolution Viewer");
 
-      JFrame frame = new JFrame("Convolution");
+      convMatrix3x3 = _convMatrix3x3;
+      try {
+         inputImage = ImageIO.read(_file);
 
-      BufferedImage testCard = ImageIO.read(new File(_args[0]));
+         // System.out.println(inputImage);
 
-      java.awt.image.Kernel conv = new java.awt.image.Kernel(3, 3, new float[] {
-            Float.parseFloat(_args[1]),
-            Float.parseFloat(_args[2]),
-            Float.parseFloat(_args[3]),
-            Float.parseFloat(_args[4]),
-            Float.parseFloat(_args[5]),
-            Float.parseFloat(_args[6]),
-            Float.parseFloat(_args[7]),
-            Float.parseFloat(_args[8]),
-            Float.parseFloat(_args[9])
-      });
-      final ConvolveOp convOp = new ConvolveOp(conv, ConvolveOp.EDGE_NO_OP, null);
+         height = inputImage.getHeight();
 
-      int imageHeight = testCard.getHeight();
+         width = inputImage.getWidth();
 
-      int imageWidth = testCard.getWidth();
+         outputImage = new BufferedImage(width, height, inputImage.getType());
 
-      final int width = padTo(imageWidth);// now multiple of 64
+         gc = outputImage.createGraphics();
 
-      final int height = padTo(imageHeight); // now multiple of 64
+         inBytes = ((DataBufferByte) inputImage.getRaster().getDataBuffer()).getData();
+         outBytes = ((DataBufferByte) outputImage.getRaster().getDataBuffer()).getData();
 
-      final BufferedImage inputImage = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
+         final JLabel imageLabel = new JLabel();
+         imageLabel.setIcon(new ImageIcon(outputImage));
 
-      inputImage.getGraphics().drawImage(testCard, padValue(imageWidth) / 2, padValue(imageHeight) / 2, null);
+         ConvMatrix3x3Editor editor = new ConvMatrix3x3Editor(_convMatrix3x3){
+            @Override protected void updated(float[] _convMatrix3x3) {
+               convMatrix3x3 = _convMatrix3x3;
+               long start = System.currentTimeMillis();
 
-      final BufferedImage outputImage = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
-      
-      convOp.filter(inputImage, outputImage);
+               applyConvolution(convMatrix3x3, inBytes, outBytes, width, height);
+               long end = System.currentTimeMillis();
+               gc.setColor(Color.BLACK);
+               gc.fillRect(0, 0, 50, 40);
+               gc.setColor(Color.YELLOW);
+               gc.drawString("" + (end - start) + "ms", 10, 20);
 
-     // outputImage.getGraphics().drawImage(testCard, padValue(imageWidth) / 2, padValue(imageHeight) / 2, null);
-      final ConvolutionKernel lifeKernel = new ConvolutionKernel(width, height, inputImage, outputImage);
+               imageLabel.repaint();
+            }
+         };
+         frame.getContentPane().add(editor.component, BorderLayout.WEST);
 
-      // Create a component for viewing the offsecreen image
-      @SuppressWarnings("serial") JComponent viewer = new JComponent(){
-         @Override public void paintComponent(Graphics g) {
+         frame.getContentPane().add(imageLabel, BorderLayout.CENTER);
+         frame.pack();
+         frame.setVisible(true);
+         frame.setDefaultCloseOperation(WindowConstants.EXIT_ON_CLOSE);
 
-            g.drawImage(outputImage, 0, 0, width, height, 0, 0, width, height, this);
-         }
-      };
+         applyConvolution(convMatrix3x3, inBytes, outBytes, width, height);
 
-      viewer.setPreferredSize(new Dimension(width, height));
-      frame.getContentPane().add(viewer);
+         imageLabel.repaint();
+      } catch (IOException e1) {
+         // TODO Auto-generated catch block
+         e1.printStackTrace();
+      }
 
-      // Swing housekeeping
-      frame.pack();
-      frame.setVisible(true);
-      frame.setDefaultCloseOperation(WindowConstants.EXIT_ON_CLOSE);
+   }
 
-      viewer.repaint();
+   abstract protected void applyConvolution(float[] convMatrix3x3, byte[] _inBytes, byte[] _outBytes, int _width, int _height);
 
-   }
 }

samples/convolution/src/com/amd/aparapi/sample/convolution/PureJava.java

+/*
+Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+following conditions are met:
+
+Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+disclaimer. 
+
+Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+disclaimer in the documentation and/or other materials provided with the distribution. 
+
+Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+derived from this software without specific prior written permission. 
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, 
+WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730 through
+774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000.  Further, pursuant to Section 740.6 of the EAR,
+you hereby certify that, except pursuant to a license granted by the United States Department of Commerce Bureau of 
+Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export Administration 
+Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in Country Groups D:1,
+E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2) export to Country Groups
+D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced direct product is subject
+to national security controls as identified on the Commerce Control List (currently found in Supplement 1 to Part 774
+of EAR).  For the most current Country Group listings, or for additional information about the EAR or your obligations
+under those regulations, please refer to the U.S. Bureau of Industry and Security's website at http://www.bis.doc.gov/. 
+
+*/
+
+package com.amd.aparapi.sample.convolution;
+
+import java.io.File;
+
+import com.amd.aparapi.Kernel;
+
+public class PureJava{
+
+   final static class ImageConvolution extends Kernel{
+      private float convMatrix3x3[];
+
+      private int width, height;
+
+      private byte imageIn[], imageOut[];
+
+      public void processPixel(int x, int y, int w, int h) {
+         float accum = 0f;
+         int count = 0;
+         for (int dx = -3; dx < 6; dx += 3) {
+            for (int dy = -1; dy < 2; dy += 1) {
+               int rgb = 0xff & imageIn[((y + dy) * w) + (x + dx)];
+
+               accum += rgb * convMatrix3x3[count++];
+            }
+         }
+         byte value = (byte) (max(0, min((int) accum, 255)));
+         imageOut[y * w + x] = value;
+
+      }
+
+      @Override public void run() {
+         int x = getGlobalId(0) % (width * 3);
+         int y = getGlobalId(0) / (width * 3);
+
+         if (x > 3 && x < (width * 3 - 3) && y > 1 && y < (height - 1)) {
+            processPixel(x, y, width * 3, height);
+         }
+
+      }
+
+      public void applyConvolution(float[] _convMatrix3x3, byte[] _imageIn, byte[] _imageOut, int _width, int _height) {
+         imageIn = _imageIn;
+         imageOut = _imageOut;
+         width = _width;
+         height = _height;
+         convMatrix3x3 = _convMatrix3x3;
+
+         execute(3 * width * height);
+      }
+
+   }
+
+   public static void main(final String[] _args) {
+      File file = new File(_args.length == 1 ? _args[0] : "testcard.jpg");
+
+      final ImageConvolution convolution = new ImageConvolution();
+
+      float convMatrix3x3[] = new float[] {
+            0f,
+            -10f,
+            0f,
+            -10f,
+            40f,
+            -10f,
+            0f,
+            -10f,
+            0f,
+      };
+
+      new ConvolutionViewer(file, convMatrix3x3){
+         @Override protected void applyConvolution(float[] _convMatrix3x3, byte[] _inBytes, byte[] _outBytes, int _width,
+               int _height) {
+            convolution.applyConvolution(_convMatrix3x3, _inBytes, _outBytes, _width, _height);
+         }
+      };
+
+   }
+
+}

samples/convolution/src/com/amd/aparapi/sample/convolution/convolution.cl

+void processPixel(__global float* _convMatrix3x3, __global char* _imageIn, __global char* _imageOut, int _width, int _height, int _x, int _y){
+   float accum = 0.0f;
+   int count = 0;
+   for (int dx = -3; dx<6; dx+=3){
+      for (int dy = -1; dy<2; dy++){
+         int rgb = 0xff & _imageIn[(((_y + dy) * _width) + (_x + dx))];
+         accum = accum + ((float)rgb * _convMatrix3x3[count++]);
+      }
+   }
+   char value = (char )max(0, min((int)accum, 255));
+   _imageOut[(_y * _width) + _x]  = value;
+   return;
+}
+
+__kernel void applyConvolution(
+   __global float *_convMatrix3x3,  // only read from kernel
+   __global char  *_imageIn, // only read from kernel
+   __global char  *_imageOut, // only written to (never read) from kernel
+   int _width,
+   int _height
+){
+ int x = get_global_id(0) % (_width * 3);
+ int y = get_global_id(0) / (_width * 3);
+ if (x>3 && x<((_width * 3) - 3) && y>1 && y<(_height - 1)){
+    processPixel(_convMatrix3x3, _imageIn, _imageOut, _width*3, _height, x, y);
+ }
+}
+