diff --git a/samples/convolution/conv.bat b/samples/convolution/conv.bat index f6a9b3b067f3fddfafde221edaf5308b35930f70..bb5dc207389aba0489b5d20481e8c3ed354c214f 100644 --- a/samples/convolution/conv.bat +++ b/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 diff --git a/samples/convolution/pureJava.bat b/samples/convolution/pureJava.bat new file mode 100644 index 0000000000000000000000000000000000000000..cd41ba09b1f7f0ed13baafb1a1529bc816c88933 --- /dev/null +++ b/samples/convolution/pureJava.bat @@ -0,0 +1,6 @@ +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 + diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvMatrix3x3Editor.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvMatrix3x3Editor.java new file mode 100644 index 0000000000000000000000000000000000000000..05f3e47c6ef44fd02250ec281fa04f11cb209484 --- /dev/null +++ b/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvMatrix3x3Editor.java @@ -0,0 +1,152 @@ + +/* +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); + } +} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/Convolution.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/Convolution.java new file mode 100644 index 0000000000000000000000000000000000000000..9cef6e81e96488a523a5df0907723ab01a67ad1c --- /dev/null +++ b/samples/convolution/src/com/amd/aparapi/sample/convolution/Convolution.java @@ -0,0 +1,117 @@ +/* +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 Convolution{ + + 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); + } + }; + + } + +} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionOpenCL.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionOpenCL.java new file mode 100644 index 0000000000000000000000000000000000000000..54b096e67852c2258d125e66910e82da1deb4831 --- /dev/null +++ b/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionOpenCL.java @@ -0,0 +1,93 @@ + +/* +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); + } + }; + + } + +} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionViewer.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionViewer.java new file mode 100644 index 0000000000000000000000000000000000000000..679189b4fec8c880187add3202979515d63f3808 --- /dev/null +++ b/samples/convolution/src/com/amd/aparapi/sample/convolution/ConvolutionViewer.java @@ -0,0 +1,132 @@ + +/* +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.Color; +import java.awt.Graphics2D; +import java.awt.image.BufferedImage; +import java.awt.image.DataBufferByte; +import java.io.File; +import java.io.IOException; + +import javax.imageio.ImageIO; +import javax.swing.ImageIcon; +import javax.swing.JFrame; +import javax.swing.JLabel; +import javax.swing.WindowConstants; + +@SuppressWarnings("serial") public abstract class ConvolutionViewer extends JFrame{ + + private int height; + + private int width; + + private BufferedImage outputImage; + + private BufferedImage inputImage; + + private byte[] inBytes; + + private byte[] outBytes; + + private Graphics2D gc; + + private float[] convMatrix3x3; + + public ConvolutionViewer(File _file, float[] _convMatrix3x3) { + + JFrame frame = new JFrame("Convolution Viewer"); + + convMatrix3x3 = _convMatrix3x3; + try { + inputImage = ImageIO.read(_file); + + // System.out.println(inputImage); + + height = inputImage.getHeight(); + + width = inputImage.getWidth(); + + outputImage = new BufferedImage(width, height, inputImage.getType()); + + gc = outputImage.createGraphics(); + + inBytes = ((DataBufferByte) inputImage.getRaster().getDataBuffer()).getData(); + outBytes = ((DataBufferByte) outputImage.getRaster().getDataBuffer()).getData(); + + final JLabel imageLabel = new JLabel(); + imageLabel.setIcon(new ImageIcon(outputImage)); + + ConvMatrix3x3Editor editor = new ConvMatrix3x3Editor(_convMatrix3x3){ + @Override protected void updated(float[] _convMatrix3x3) { + convMatrix3x3 = _convMatrix3x3; + long start = System.currentTimeMillis(); + + 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); + + imageLabel.repaint(); + } + }; + frame.getContentPane().add(editor.component, BorderLayout.WEST); + + frame.getContentPane().add(imageLabel, BorderLayout.CENTER); + frame.pack(); + frame.setVisible(true); + frame.setDefaultCloseOperation(WindowConstants.EXIT_ON_CLOSE); + + applyConvolution(convMatrix3x3, inBytes, outBytes, width, height); + + imageLabel.repaint(); + } catch (IOException e1) { + // TODO Auto-generated catch block + e1.printStackTrace(); + } + + } + + abstract protected void applyConvolution(float[] convMatrix3x3, byte[] _inBytes, byte[] _outBytes, int _width, int _height); + +} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/Convolve.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/Convolve.java deleted file mode 100644 index b3826a22a393809d4053bbd38d58e6d6cd003b42..0000000000000000000000000000000000000000 --- a/samples/convolution/src/com/amd/aparapi/sample/convolution/Convolve.java +++ /dev/null @@ -1,226 +0,0 @@ -/* -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.Dimension; -import java.awt.Graphics; -import java.awt.image.BufferedImage; -import java.awt.image.ConvolveOp; -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 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{ - - private final float[] filter = new float[9]; - - private final int[] inputData; - - 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; - - // 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); - 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; - } - } - - 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 static int padValue(int value) { - return (PAD - (value % PAD)); - } - - 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(_args[0])); - - 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); - - 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 - - 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); - - convOp.filter(inputImage, outputImage); - - // 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) { - - g.drawImage(outputImage, 0, 0, width, height, 0, 0, width, height, this); - } - }; - - viewer.setPreferredSize(new Dimension(width, height)); - frame.getContentPane().add(viewer); - - // Swing housekeeping - frame.pack(); - frame.setVisible(true); - frame.setDefaultCloseOperation(WindowConstants.EXIT_ON_CLOSE); - - viewer.repaint(); - - } -} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/Main.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/Main.java deleted file mode 100644 index 9ad0f872008613953edd2be753a4d05ceb8677f3..0000000000000000000000000000000000000000 --- a/samples/convolution/src/com/amd/aparapi/sample/convolution/Main.java +++ /dev/null @@ -1,234 +0,0 @@ -/* -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.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); - - private static final ConvolutionFilter EMBOSS = new ConvolutionFilter(-2f, -1f, 0f, -1f, 1f, 1f, 0f, 1f, 2f, 0); - - public static class ConvolutionKernel extends Kernel{ - - private final float[] filter = new float[9]; - - private final int[] inputData; - - 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; - - // 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); - 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; - } - } - - 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 static int padValue(int value) { - return (PAD - (value % PAD)); - } - - 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, - }; - 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); - } - } - System.out.println((System.nanoTime() - start) / 1000000); - - } -} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/PureJava.java b/samples/convolution/src/com/amd/aparapi/sample/convolution/PureJava.java new file mode 100644 index 0000000000000000000000000000000000000000..5c0158df3e986a555bf57c960f29738bf8f3554c --- /dev/null +++ b/samples/convolution/src/com/amd/aparapi/sample/convolution/PureJava.java @@ -0,0 +1,117 @@ +/* +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); + } + }; + + } + +} diff --git a/samples/convolution/src/com/amd/aparapi/sample/convolution/convolution.cl b/samples/convolution/src/com/amd/aparapi/sample/convolution/convolution.cl new file mode 100644 index 0000000000000000000000000000000000000000..e1acfa8f5ad707775443cc179e187c9591aca099 --- /dev/null +++ b/samples/convolution/src/com/amd/aparapi/sample/convolution/convolution.cl @@ -0,0 +1,28 @@ +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); + } +} +