run() method (and every method that can be called directly or indirectly from the run() method)
- * can be converted into OpenCL.* Below is an example Kernel that calculates the square of a set of input values. @@ -154,12 +141,12 @@ to national security controls as identified on the Commerce Control List (curren * A different approach to creating kernels that avoids extending Kernel is to write an anonymous inner class: *
*
*- * + ** final int[] values = new int[1024]; * // fill the values array * final int[] squares = new int[values.length]; * final Range range = Range.create(values.length); - * + *
* Kernel kernel = new Kernel(){ * public void run() { * int gid = getGlobalID(); @@ -170,559 +157,568 @@ to national security controls as identified on the Commerce Control List (curren * for (int i=0; i< values.length; i++){ * System.out.printf("%4d %4d %8d\n", i, values[i], squares[i]); * } - * + *
*
* - * @author gfrost AMD Javalabs + * @author gfrost AMD Javalabs * @version Alpha, 21/09/2010 */ public abstract class Kernel implements Cloneable { - private static Logger logger = Logger.getLogger(Config.getLoggerName()); - - /** - * We can use this Annotation to 'tag' intended local buffers. - * - * So we can either annotate the buffer - *
- * @Local int[] buffer = new int[1024];
- *
- * Or use a special suffix
- *
- * int[] buffer_$local$ = new int[1024];
- *
- *
- * @see #LOCAL_SUFFIX
- *
- *
- */
- @Retention(RetentionPolicy.RUNTIME)
- public @interface Local {
-
- }
-
- /**
- * We can use this Annotation to 'tag' intended constant buffers.
- *
- * So we can either annotate the buffer
- *
- * @Constant int[] buffer = new int[1024];
- *
- * Or use a special suffix
- *
- * int[] buffer_$constant$ = new int[1024];
- *
- *
- * @see #LOCAL_SUFFIX
- *
- *
- */
- @Retention(RetentionPolicy.RUNTIME)
- public @interface Constant {
-
- }
-
- /**
- *
- * We can use this Annotation to 'tag' __private (unshared) array fields. Data in the __private address space in OpenCL is accessible only from
- * the current kernel instance.
- *
- * To so mark a field with a buffer size of 99, we can either annotate the buffer
- *
- * @PrivateMemorySpace(99) int[] buffer = new int[99];
- *
- * Or use a special suffix
- *
- * int[] buffer_$private$99 = new int[99];
- *
- *
- * Note that any code which must be runnable in {@link EXECUTION_MODE#JTP} will fail to work correctly if it uses such an
- * array, as the array will be shared by all threads. The solution is to create a {@link NoCL} method called at the start of {@link #run()} which sets
- * the field to an array returned from a static ThreadLocal
MedianKernel7x7 in the samples for an example.
- *
- * @see #PRIVATE_SUFFIX
- */
- @Retention(RetentionPolicy.RUNTIME)
- @Target({ElementType.FIELD})
- public @interface PrivateMemorySpace {
- /** Size of the array used as __private buffer. */
- int value();
- }
-
- /**
- * Annotation which can be applied to either a getter (with usual java bean naming convention relative to an instance field), or to any method
- * with void return type, which prevents both the method body and any calls to the method being emitted in the generated OpenCL. (In the case of a getter, the
- * underlying field is used in place of the NoCL getter method.) This allows for code specialization within a java/JTP execution path, for example to
- * allow logging/breakpointing when debugging, or to apply ThreadLocal processing (see {@link PrivateMemorySpace}) in java to simulate OpenCL __private
- * memory.
- */
- @Retention(RetentionPolicy.RUNTIME)
- @Target({ElementType.METHOD, ElementType.FIELD})
- public @interface NoCL {
- // empty
- }
-
- /**
- * We can use this suffix to 'tag' intended local buffers.
- *
- *
- * So either name the buffer
- *
- * int[] buffer_$local$ = new int[1024];
- *
- * Or use the Annotation form
- *
- * @Local int[] buffer = new int[1024];
- *
- */
- public final static String LOCAL_SUFFIX = "_$local$";
-
- /**
- * We can use this suffix to 'tag' intended constant buffers.
- *
- *
- * So either name the buffer
- *
- * int[] buffer_$constant$ = new int[1024];
- *
- * Or use the Annotation form
- *
- * @Constant int[] buffer = new int[1024];
- *
- */
- public final static String CONSTANT_SUFFIX = "_$constant$";
-
- /**
- * We can use this suffix to 'tag' __private buffers.
- *
- * So either name the buffer - *
- * int[] buffer_$private$32 = new int[32];
- *
- * Or use the Annotation form
- *
- * @PrivateMemorySpace(32) int[] buffer = new int[32];
- *
- *
- * @see PrivateMemorySpace for a more detailed usage summary
- */
- public final static String PRIVATE_SUFFIX = "_$private$";
-
- /**
- * This annotation is for internal use only
- */
- @Retention(RetentionPolicy.RUNTIME)
- protected @interface OpenCLDelegate {
-
- }
-
- /**
- * This annotation is for internal use only
- */
- @Retention(RetentionPolicy.RUNTIME)
- protected @interface OpenCLMapping {
- String mapTo() default "";
-
- boolean atomic32() default false;
-
- boolean atomic64() default false;
- }
-
- public abstract class Entry {
- public abstract void run();
-
- public Kernel execute(Range _range) {
- return (Kernel.this.execute("foo", _range, 1));
- }
- }
-
- /**
- * @deprecated It is no longer recommended that {@code EXECUTION_MODE}s are used, as a more sophisticated {@link com.aparapi.device.Device}
- * preference mechanism is in place, see {@link com.aparapi.internal.kernel.KernelManager}. Though {@link #setExecutionMode(EXECUTION_MODE)}
- * is still honored, the default EXECUTION_MODE is now {@link EXECUTION_MODE#AUTO}, which indicates that the KernelManager
- * will determine execution behaviours.
- *
- * - * The execution mode ENUM enumerates the possible modes of executing a kernel. - * One can request a mode of execution using the values below, and query a kernel after it first executes to - * determine how it executed. - * - *
- * Aparapi supports 5 execution modes. Default is GPU. - *
| Enum value | Execution |
|---|---|
GPU | Execute using OpenCL on first available GPU device |
ACC | Execute using OpenCL on first available Accelerator device |
CPU | Execute using OpenCL on first available CPU device |
JTP | Execute using a Java Thread Pool (one thread spawned per available core) |
SEQ | Execute using a single loop. This is useful for debugging but will be less - * performant than the other modes |
- * To request that a kernel is executed in a specific mode, call Kernel.setExecutionMode(EXECUTION_MODE) before the
- * kernel first executes.
- *
- *
- *- * int[] values = new int[1024]; - * // fill values array - * SquareKernel kernel = new SquareKernel(values); - * kernel.setExecutionMode(Kernel.EXECUTION_MODE.JTP); - * kernel.execute(values.length); - *
-<<<<<<< HEAD:src/main/java/com/codegen/Kernel.java
- * Alternatively, the property com.codegen.executionMode can be set to one of JTP,GPU,ACC,CPU,SEQ
- * when an application is launched.
- *
- * java -classpath ....;codegen.jar -Dcom.codegen.executionMode=GPU MyApplication -======= - * Alternatively, the propertycom.amd.codegen.executionModecan be set to one ofJTP,GPU,ACC,CPU,SEQ- * when an application is launched. - *- * java -classpath ....;codegen.jar -Dcom.amd.codegen.executionMode=GPU MyApplication ->>>>>>> b118aad... added method to set execution mode without any fallback:com.amd.codegen/src/java/com/amd/codegen/Kernel.java - *- * Generally setting the execution mode is not recommended (it is best to let Aparapi decide automatically) but the option - * provides a way to compare a kernel's performance under multiple execution modes. - * - * @author gfrost AMD Javalabs - * @version Alpha, 21/09/2010 - */ - @Deprecated - public static enum EXECUTION_MODE { - /** - * - */ - AUTO, - /** - * A dummy value to indicate an unknown state. - */ - NONE, - /** - * The value representing execution on a GPU device via OpenCL. - */ - GPU, - /** - * The value representing execution on a CPU device via OpenCL. - *
- * Note not all OpenCL implementations support OpenCL compute on the CPU. - */ - CPU, - /** - * The value representing execution on a Java Thread Pool. - *
- * By default one Java thread is started for each available core and each core will execute
globalSize/coreswork items. - * This creates a total ofglobalSize%coresthreads to complete the work. - * Choose suitable values forglobalSizeto minimize the number of threads that are spawned. - */ - JTP, - /** - * The value representing execution sequentially in a single loop. - *- * This is meant to be used for debugging a kernel. - */ - SEQ, - /** - * The value representing execution on an accelerator device (Xeon Phi) via OpenCL. - */ - ACC; - - /** - * @deprecated See {@link EXECUTION_MODE}. - */ - @Deprecated - static LinkedHashSet
getDefaultExecutionModes() { - LinkedHashSet defaultExecutionModes = new LinkedHashSet (); - - if (OpenCLLoader.isOpenCLAvailable()) { - defaultExecutionModes.add(GPU); - defaultExecutionModes.add(JTP); - } else { - defaultExecutionModes.add(JTP); - } - - final String executionMode = Config.executionMode; - - if (executionMode != null) { - try { - LinkedHashSet requestedExecutionModes; - requestedExecutionModes = EXECUTION_MODE.getExecutionModeFromString(executionMode); - logger.fine("requested execution mode ="); - for (final EXECUTION_MODE mode : requestedExecutionModes) { - logger.fine(" " + mode); - } - if ((OpenCLLoader.isOpenCLAvailable() && EXECUTION_MODE.anyOpenCL(requestedExecutionModes)) - || !EXECUTION_MODE.anyOpenCL(requestedExecutionModes)) { - defaultExecutionModes = requestedExecutionModes; - } - } catch (final Throwable t) { - // we will take the default - } - } - - logger.info("default execution modes = " + defaultExecutionModes); + final static AtomicInteger serials = new AtomicInteger(); + final int serial = serials.getAndIncrement(); - for (final EXECUTION_MODE e : defaultExecutionModes) { - logger.info("SETTING DEFAULT MODE: " + e.toString()); - } - - return (defaultExecutionModes); - } - - static LinkedHashSet getExecutionModeFromString(String executionMode) { - final LinkedHashSet executionModes = new LinkedHashSet (); - for (final String mode : executionMode.split(",")) { - executionModes.add(valueOf(mode.toUpperCase())); - } - return executionModes; - } + @Override + public int hashCode() { + return serial; + } - static EXECUTION_MODE getFallbackExecutionMode() { - final EXECUTION_MODE defaultFallbackExecutionMode = JTP; - logger.info("fallback execution mode = " + defaultFallbackExecutionMode); - return (defaultFallbackExecutionMode); - } + @Override + public boolean equals(Object obj) { + return this==obj; + } - static boolean anyOpenCL(LinkedHashSet _executionModes) { - for (final EXECUTION_MODE mode : _executionModes) { - if ((mode == GPU) || (mode == ACC) || (mode == CPU)) { - return true; - } - } - return false; - } + private static final Logger logger = Logger.getLogger(Config.getLoggerName()); - public boolean isOpenCL() { - return (this == GPU) || (this == ACC) || (this == CPU); - } - }; + public Kernel clone(Range range) { + Kernel k = clone(); + k.kernelState.range = range; + return k; + } - private KernelRunner kernelRunner = null; + /** + * We can use this Annotation to 'tag' intended local buffers. + * + * So we can either annotate the buffer + *
+ * Or use a special suffix + *+ * @Local int[] buffer = new int[1024]; + *+ * + * @see #LOCAL_SUFFIX + */ + @Retention(RetentionPolicy.RUNTIME) + public @interface Local { - private boolean autoCleanUpArrays = false; + } - private KernelState kernelState = new KernelState(); + /** + * We can use this Annotation to 'tag' intended constant buffers. + *+ * int[] buffer_$local$ = new int[1024]; + *+ * So we can either annotate the buffer + *
+ * Or use a special suffix + *+ * @Constant int[] buffer = new int[1024]; + *+ * + * @see #LOCAL_SUFFIX + */ + @Retention(RetentionPolicy.RUNTIME) + public @interface Constant { - /** - * This class is for internal Kernel state management+ * int[] buffer_$constant$ = new int[1024]; + *- * NOT INTENDED FOR USE BY USERS - */ - public final class KernelState { + } - private int[] globalIds = new int[] {0, 0, 0}; + /** + * We can use this Annotation to 'tag' __private (unshared) array fields. Data in the __private address space in OpenCL is accessible only from + * the current kernel instance. + *
+ * To so mark a field with a buffer size of 99, we can either annotate the buffer + *
+ * Or use a special suffix + *+ * @PrivateMemorySpace(99) int[] buffer = new int[99]; + *+ *+ * int[] buffer_$private$99 = new int[99]; + *+ *
Note that any code which must be runnable in {@link EXECUTION_MODE#JTP} will fail to work correctly if it uses such an + * array, as the array will be shared by all threads. The solution is to create a {@link NoCL} method called at the start of {@link #run()} which sets + * the field to an array returned from a static
. Please seeThreadLocalMedianKernel7x7in the samples for an example. + * + * @see #PRIVATE_SUFFIX + */ + @Retention(RetentionPolicy.RUNTIME) + @Target({ElementType.FIELD}) + public @interface PrivateMemorySpace { + /** + * Size of the array used as __private buffer. + */ + int value(); + } - private int[] localIds = new int[] {0, 0, 0}; + /** + * Annotation which can be applied to either a getter (with usual java bean naming convention relative to an instance field), or to any method + * with void return type, which prevents both the method body and any calls to the method being emitted in the generated OpenCL. (In the case of a getter, the + * underlying field is used in place of the NoCL getter method.) This allows for code specialization within a java/JTP execution path, for example to + * allow logging/breakpointing when debugging, or to apply ThreadLocal processing (see {@link PrivateMemorySpace}) in java to simulate OpenCL __private + * memory. + */ + @Retention(RetentionPolicy.RUNTIME) + @Target({ElementType.METHOD, ElementType.FIELD}) + public @interface NoCL { + // empty + } - private int[] groupIds = new int[] {0, 0, 0}; + /** + * We can use this suffix to 'tag' intended local buffers. + *+ *
+ * So either name the buffer + *
+ * Or use the Annotation form + *+ * int[] buffer_$local$ = new int[1024]; + *+ */ + public final static String LOCAL_SUFFIX = "_$local$"; - private Range range; + /** + * We can use this suffix to 'tag' intended constant buffers. + *+ * @Local int[] buffer = new int[1024]; + *+ *
+ * So either name the buffer + *
+ * Or use the Annotation form + *+ * int[] buffer_$constant$ = new int[1024]; + *+ */ + public final static String CONSTANT_SUFFIX = "_$constant$"; - private int passId; + /** + * We can use this suffix to 'tag' __private buffers. + *+ * @Constant int[] buffer = new int[1024]; + *+ *
So either name the buffer + *
+ * Or use the Annotation form + *+ * int[] buffer_$private$32 = new int[32]; + *+ * + * @see PrivateMemorySpace for a more detailed usage summary + */ + public final static String PRIVATE_SUFFIX = "_$private$"; - private volatile CyclicBarrier localBarrier; + /** + * This annotation is for internal use only + */ + @Retention(RetentionPolicy.RUNTIME) + private @interface OpenCLDelegate { - private boolean localBarrierDisabled; + } - /** - * Default constructor - */ - protected KernelState() { + /** + * This annotation is for internal use only + */ + @Retention(RetentionPolicy.RUNTIME) + @interface OpenCLMapping { + String mapTo() default ""; - } + boolean atomic32() default false; - /** - * Copy constructor - */ - protected KernelState(KernelState kernelState) { - globalIds = kernelState.getGlobalIds(); - localIds = kernelState.getLocalIds(); - groupIds = kernelState.getGroupIds(); - range = kernelState.getRange(); - passId = kernelState.getPassId(); - localBarrier = kernelState.getLocalBarrier(); - } + boolean atomic64() default false; + } - /** - * @return the globalIds - */ - public int[] getGlobalIds() { - return globalIds; - } + public abstract class Entry { + public abstract void run(); - /** - * @param globalIds the globalIds to set - */ - public void setGlobalIds(int[] globalIds) { - this.globalIds = globalIds; - } + public Kernel execute(Range _range) { + return (Kernel.this.execute("foo", _range, 1)); + } + } - /** - * Set a specific index value - * - * @param _index - * @param value - */ - public void setGlobalId(int _index, int value) { - globalIds[_index] = value; - } + /** + * @author gfrost AMD Javalabs + * @version Alpha, 21/09/2010 + * @deprecated It is no longer recommended that {@code EXECUTION_MODE}s are used, as a more sophisticated {@link com.aparapi.device.Device} + * preference mechanism is in place, see {@link com.aparapi.internal.kernel.KernelManager}. Though {@link #setExecutionMode(EXECUTION_MODE)} + * is still honored, the default EXECUTION_MODE is now {@link EXECUTION_MODE#AUTO}, which indicates that the KernelManager + * will determine execution behaviours. + *+ * @PrivateMemorySpace(32) int[] buffer = new int[32]; + *+ *
+ * The execution mode ENUM enumerates the possible modes of executing a kernel. + * One can request a mode of execution using the values below, and query a kernel after it first executes to + * determine how it executed. + *
+ *
+ * Aparapi supports 5 execution modes. Default is GPU. + *
+ *
+ *+ *
+ *+ * Enum value Execution + * GPUExecute using OpenCL on first available GPU device + * ACCExecute using OpenCL on first available Accelerator device + * CPUExecute using OpenCL on first available CPU device + * JTPExecute using a Java Thread Pool (one thread spawned per available core) + * SEQExecute using a single loop. This is useful for debugging but will be less + * performant than the other modes + * To request that a kernel is executed in a specific mode, call
Kernel.setExecutionMode(EXECUTION_MODE)before the + * kernel first executes. + *+ *
+ *+ * int[] values = new int[1024]; + * // fill values array + * SquareKernel kernel = new SquareKernel(values); + * kernel.setExecutionMode(Kernel.EXECUTION_MODE.JTP); + * kernel.execute(values.length); + *+ * <<<<<<< HEAD:src/main/java/com/codegen/Kernel.java + * Alternatively, the property
com.codegen.executionModecan be set to one ofJTP,GPU,ACC,CPU,SEQ+ * when an application is launched. + *+ * java -classpath ....;codegen.jar -Dcom.codegen.executionMode=GPU MyApplication + * ======= + * Alternatively, the propertycom.amd.codegen.executionModecan be set to one ofJTP,GPU,ACC,CPU,SEQ+ * when an application is launched. + *+ * java -classpath ....;codegen.jar -Dcom.amd.codegen.executionMode=GPU MyApplication + * >>>>>>> b118aad... added method to set execution mode without any fallback:com.amd.codegen/src/java/com/amd/codegen/Kernel.java + *+ * Generally setting the execution mode is not recommended (it is best to let Aparapi decide automatically) but the option + * provides a way to compare a kernel's performance under multiple execution modes. + */ + @Deprecated + public enum EXECUTION_MODE { + /** + * + */ + AUTO, + /** + * A dummy value to indicate an unknown state. + */ + NONE, + /** + * The value representing execution on a GPU device via OpenCL. + */ + GPU, + /** + * The value representing execution on a CPU device via OpenCL. + *
+ * Note not all OpenCL implementations support OpenCL compute on the CPU. + */ + CPU, + /** + * The value representing execution on a Java Thread Pool. + *
+ * By default one Java thread is started for each available core and each core will execute
globalSize/coreswork items. + * This creates a total ofglobalSize%coresthreads to complete the work. + * Choose suitable values forglobalSizeto minimize the number of threads that are spawned. + */ + JTP, + /** + * The value representing execution sequentially in a single loop. + *+ * This is meant to be used for debugging a kernel. + */ + SEQ, + /** + * The value representing execution on an accelerator device (Xeon Phi) via OpenCL. + */ + ACC; + + /** + * @deprecated See {@link EXECUTION_MODE}. + */ + @Deprecated + static LinkedHashSet
getDefaultExecutionModes() { + LinkedHashSet defaultExecutionModes = new LinkedHashSet<>(); + + if (OpenCLLoader.isOpenCLAvailable()) { + defaultExecutionModes.add(GPU); + defaultExecutionModes.add(JTP); + } else { + defaultExecutionModes.add(JTP); + } - /** - * @return the localIds - */ - public int[] getLocalIds() { - return localIds; - } + final String executionMode = Config.executionMode; + + if (executionMode != null) { + try { + LinkedHashSet requestedExecutionModes; + requestedExecutionModes = EXECUTION_MODE.getExecutionModeFromString(executionMode); + if (logger.isLoggable(Level.FINE)) { + logger.fine("requested execution mode ="); + for (final EXECUTION_MODE mode : requestedExecutionModes) { + logger.fine(" " + mode); + } + } + if ((OpenCLLoader.isOpenCLAvailable() && EXECUTION_MODE.anyOpenCL(requestedExecutionModes)) + || !EXECUTION_MODE.anyOpenCL(requestedExecutionModes)) { + defaultExecutionModes = requestedExecutionModes; + } + } catch (final Throwable t) { + // we will take the default + } + } - /** - * @param localIds the localIds to set - */ - public void setLocalIds(int[] localIds) { - this.localIds = localIds; - } + logger.info("default execution modes = " + defaultExecutionModes); - /** - * Set a specific index value - * - * @param _index - * @param value - */ - public void setLocalId(int _index, int value) { - localIds[_index] = value; - } + for (final EXECUTION_MODE e : defaultExecutionModes) { + logger.info("SETTING DEFAULT MODE: " + e.toString()); + } - /** - * @return the groupIds - */ - public int[] getGroupIds() { - return groupIds; - } + return (defaultExecutionModes); + } - /** - * @param groupIds the groupIds to set - */ - public void setGroupIds(int[] groupIds) { - this.groupIds = groupIds; - } + static LinkedHashSet getExecutionModeFromString(String executionMode) { + final LinkedHashSet executionModes = new LinkedHashSet<>(); + for (final String mode : executionMode.split(",")) { + executionModes.add(valueOf(mode.toUpperCase())); + } + return executionModes; + } + + static EXECUTION_MODE getFallbackExecutionMode() { + final EXECUTION_MODE defaultFallbackExecutionMode = JTP; + logger.info("fallback execution mode = " + defaultFallbackExecutionMode); + return (defaultFallbackExecutionMode); + } + + static boolean anyOpenCL(LinkedHashSet _executionModes) { + for (final EXECUTION_MODE mode : _executionModes) { + if ((mode == GPU) || (mode == ACC) || (mode == CPU)) { + return true; + } + } + return false; + } - /** - * Set a specific index value - * - * @param _index - * @param value - */ - public void setGroupId(int _index, int value) { - groupIds[_index] = value; - } + public boolean isOpenCL() { + return (this == GPU) || (this == ACC) || (this == CPU); + } + } - /** - * @return the range - */ - public Range getRange() { - return range; - } - /** - * @param range the range to set - */ - public void setRange(Range range) { - this.range = range; - } + private boolean autoCleanUpArrays = false; - /** - * @return the passId - */ - public int getPassId() { - return passId; - } + private KernelState kernelState = new KernelState(); - /** - * @param passId the passId to set - */ - public void setPassId(int passId) { - this.passId = passId; - } + /** + * This class is for internal Kernel state management + * NOT INTENDED FOR USE BY USERS + */ + public final class KernelState { + + public final AtomicIntegerArray globalIds = new AtomicIntegerArray(3); + + private int[] localIds = new int[]{0, 0, 0}; + + private int[] groupIds = new int[]{0, 0, 0}; + + private Range range; + + private int passId; + + private volatile CyclicBarrier localBarrier; + + private boolean localBarrierDisabled; + + /** + * Default constructor + */ + KernelState() { + + } + + /** + * Copy constructor + */ + KernelState(KernelState kernelState) { + for (int i = 0; i < globalIds.length(); i++) + globalIds.set(i, kernelState.globalIds.get(i)); + localIds = kernelState.getLocalIds(); + groupIds = kernelState.getGroupIds(); + range = kernelState.getRange(); + passId = kernelState.getPassId(); + localBarrier = kernelState.getLocalBarrier(); + } + + /** + * @param globalIds the globalIds to set + */ + void setGlobalIds(int[] globalIds) { + int i = 0; + for (int x : globalIds) + this.globalIds.set(i++, x); + } + + /** + * Set a specific index value + * + * @param _index + * @param value + */ + public void setGlobalId(int _index, int value) { + globalIds.set(_index, value); + } + + /** + * @return the localIds + */ + public int[] getLocalIds() { + return localIds; + } + + /** + * @param localIds the localIds to set + */ + void setLocalIds(int[] localIds) { + this.localIds = localIds; + } + + /** + * Set a specific index value + * + * @param _index + * @param value + */ + public void setLocalId(int _index, int value) { + localIds[_index] = value; + } + + /** + * @return the groupIds + */ + int[] getGroupIds() { + return groupIds; + } + + /** + * @param groupIds the groupIds to set + */ + void setGroupIds(int[] groupIds) { + this.groupIds = groupIds; + } + + /** + * Set a specific index value + * + * @param _index + * @param value + */ + public void setGroupId(int _index, int value) { + groupIds[_index] = value; + } + + /** + * @return the range + */ + Range getRange() { + return range; + } + + /** + * @param range the range to set + */ + public void setRange(Range range) { + this.range = range; + } + + /** + * @return the passId + */ + int getPassId() { + return passId; + } + + /** + * @param passId the passId to set + */ + public void setPassId(int passId) { + this.passId = passId; + } + + /** + * @return the localBarrier + */ + CyclicBarrier getLocalBarrier() { + return localBarrier; + } + + /** + * @param localBarrier the localBarrier to set + */ + public void setLocalBarrier(CyclicBarrier localBarrier) { + this.localBarrier = localBarrier; + } + + void awaitOnLocalBarrier() { + if (!localBarrierDisabled) { + try { + kernelState.getLocalBarrier().await(); + } catch (final InterruptedException | BrokenBarrierException e) { + // TODO Auto-generated catch block + e.printStackTrace(); + } + } + } - /** - * @return the localBarrier - */ - public CyclicBarrier getLocalBarrier() { - return localBarrier; - } + public void disableLocalBarrier() { + localBarrierDisabled = true; + } + } - /** - * @param localBarrier the localBarrier to set - */ - public void setLocalBarrier(CyclicBarrier localBarrier) { - this.localBarrier = localBarrier; - } + /** + * Determine the globalId of an executing kernel. + *
+ * The kernel implementation uses the globalId to determine which of the executing kernels (in the global domain space) this invocation is expected to deal with. + *
+ * For example in a
SquareKernelimplementation: + *+ *
+ *+ * class SquareKernel extends Kernel{ + * private int values[]; + * private int squares[]; + * public SquareKernel(int values[]){ + * this.values = values; + * squares = new int[values.length]; + * } + * public void run() { + * int gid = getGlobalID(); + * squares[gid] = values[gid]*values[gid]; + * } + * public int[] getSquares(){ + * return(squares); + * } + * } + *+ * Each invocation of
SquareKernel.run()retrieves it's globalId by callinggetGlobalId(), and then computes the value ofsquare[gid]for a given value ofvalue[gid]. + *+ * + * @return The globalId for the Kernel being executed + * @see #getLocalId() + * @see #getGroupId() + * @see #getGlobalSize() + * @see #getNumGroups() + * @see #getLocalSize() + */ - public void awaitOnLocalBarrier() { - if (!localBarrierDisabled) { - try { - kernelState.getLocalBarrier().await(); - } catch (final InterruptedException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final BrokenBarrierException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } - } - } + @OpenCLDelegate + protected final int getGlobalId() { + return getGlobalId(0); + } - public void disableLocalBarrier() { - localBarrierDisabled = true; - } - } - - /** - * Determine the globalId of an executing kernel. - *
- * The kernel implementation uses the globalId to determine which of the executing kernels (in the global domain space) this invocation is expected to deal with. - *
- * For example in a
SquareKernelimplementation: - *- *
- *- * class SquareKernel extends Kernel{ - * private int values[]; - * private int squares[]; - * public SquareKernel(int values[]){ - * this.values = values; - * squares = new int[values.length]; - * } - * public void run() { - * int gid = getGlobalID(); - * squares[gid] = values[gid]*values[gid]; - * } - * public int[] getSquares(){ - * return(squares); - * } - * } - *- * Each invocation of
SquareKernel.run()retrieves it's globalId by callinggetGlobalId(), and then computes the value ofsquare[gid]for a given value ofvalue[gid]. - *- * @return The globalId for the Kernel being executed - * - * @see #getLocalId() - * @see #getGroupId() - * @see #getGlobalSize() - * @see #getNumGroups() - * @see #getLocalSize() - */ - - @OpenCLDelegate - protected final int getGlobalId() { - return getGlobalId(0); - } - - @OpenCLDelegate - protected final int getGlobalId(int _dim) { - return kernelState.getGlobalIds()[_dim]; - } + @OpenCLDelegate + protected final int getGlobalId(int _dim) { + return kernelState.globalIds.get(_dim); + } /* @OpenCLDelegate protected final int getGlobalX() { @@ -737,46 +733,46 @@ protected final int getGlobalId(int _dim) { return (getGlobalId(2)); } */ - /** - * Determine the groupId of an executing kernel. - *
- * When a
Kernel.execute(int globalSize)is invoked for a particular kernel, the runtime will break the work into various 'groups'. - *- * A kernel can use
getGroupId()to determine which group a kernel is currently - * dispatched to - *- * The following code would capture the groupId for each kernel and map it against globalId. - *
- * - * @see #getLocalId() - * @see #getGlobalId() - * @see #getGlobalSize() - * @see #getNumGroups() - * @see #getLocalSize() - * - * @return The groupId for this Kernel being executed - */ - @OpenCLDelegate - protected final int getGroupId() { - return getGroupId(0); - } - - @OpenCLDelegate - protected final int getGroupId(int _dim) { - return kernelState.getGroupIds()[_dim]; - } + + /** + * Determine the groupId of an executing kernel. + *- * final int[] groupIds = new int[1024]; - * Kernel kernel = new Kernel(){ - * public void run() { - * int gid = getGlobalId(); - * groupIds[gid] = getGroupId(); - * } - * }; - * kernel.execute(groupIds.length); - * for (int i=0; i< values.length; i++){ - * System.out.printf("%4d %4d\n", i, groupIds[i]); - * } - *+ * When a
Kernel.execute(int globalSize)is invoked for a particular kernel, the runtime will break the work into various 'groups'. + *+ * A kernel can use
getGroupId()to determine which group a kernel is currently + * dispatched to + *+ * The following code would capture the groupId for each kernel and map it against globalId. + *
+ * + * @return The groupId for this Kernel being executed + * @see #getLocalId() + * @see #getGlobalId() + * @see #getGlobalSize() + * @see #getNumGroups() + * @see #getLocalSize() + */ + @OpenCLDelegate + protected final int getGroupId() { + return getGroupId(0); + } + + @OpenCLDelegate + private int getGroupId(int _dim) { + return kernelState.getGroupIds()[_dim]; + } /* @OpenCLDelegate protected final int getGroupX() { @@ -791,65 +787,64 @@ protected final int getGroupId(int _dim) { return (getGroupId(2)); } */ - /** - * Determine the passId of an executing kernel. - *+ * final int[] groupIds = new int[1024]; + * Kernel kernel = new Kernel(){ + * public void run() { + * int gid = getGlobalId(); + * groupIds[gid] = getGroupId(); + * } + * }; + * kernel.execute(groupIds.length); + * for (int i=0; i< values.length; i++){ + * System.out.printf("%4d %4d\n", i, groupIds[i]); + * } + *- * When a
Kernel.execute(int globalSize, int passes)is invoked for a particular kernel, the runtime will break the work into various 'groups'. - *- * A kernel can use
getPassId()to determine which pass we are in. This is ideal for 'reduce' type phases - * - * @see #getLocalId() - * @see #getGlobalId() - * @see #getGlobalSize() - * @see #getNumGroups() - * @see #getLocalSize() - * - * @return The groupId for this Kernel being executed - */ - @OpenCLDelegate - protected final int getPassId() { - return kernelState.getPassId(); - } - - /** - * Determine the local id of an executing kernel. - *- * When a
Kernel.execute(int globalSize)is invoked for a particular kernel, the runtime will break the work into - * various 'groups'. - *getLocalId()can be used to determine the relative id of the current kernel within a specific group. - *- * The following code would capture the groupId for each kernel and map it against globalId. - *
- * - * @see #getGroupId() - * @see #getGlobalId() - * @see #getGlobalSize() - * @see #getNumGroups() - * @see #getLocalSize() - * - * @return The local id for this Kernel being executed - */ - @OpenCLDelegate - protected final int getLocalId() { - return getLocalId(0); - } - - @OpenCLDelegate - protected final int getLocalId(int _dim) { - return kernelState.getLocalIds()[_dim]; - } + + /** + * Determine the passId of an executing kernel. + *- * final int[] localIds = new int[1024]; - * Kernel kernel = new Kernel(){ - * public void run() { - * int gid = getGlobalId(); - * localIds[gid] = getLocalId(); - * } - * }; - * kernel.execute(localIds.length); - * for (int i=0; i< values.length; i++){ - * System.out.printf("%4d %4d\n", i, localIds[i]); - * } - *+ * When a
Kernel.execute(int globalSize, int passes)is invoked for a particular kernel, the runtime will break the work into various 'groups'. + *+ * A kernel can use
getPassId()to determine which pass we are in. This is ideal for 'reduce' type phases + * + * @return The groupId for this Kernel being executed + * @see #getLocalId() + * @see #getGlobalId() + * @see #getGlobalSize() + * @see #getNumGroups() + * @see #getLocalSize() + */ + @OpenCLDelegate + protected final int getPassId() { + return kernelState.getPassId(); + } + + /** + * Determine the local id of an executing kernel. + *+ * When a
Kernel.execute(int globalSize)is invoked for a particular kernel, the runtime will break the work into + * various 'groups'. + *getLocalId()can be used to determine the relative id of the current kernel within a specific group. + *+ * The following code would capture the groupId for each kernel and map it against globalId. + *
+ * + * @return The local id for this Kernel being executed + * @see #getGroupId() + * @see #getGlobalId() + * @see #getGlobalSize() + * @see #getNumGroups() + * @see #getLocalSize() + */ + @OpenCLDelegate + protected final int getLocalId() { + return getLocalId(0); + } + + @OpenCLDelegate + protected final int getLocalId(int _dim) { + return kernelState.getLocalIds()[_dim]; + } /* @OpenCLDelegate protected final int getLocalX() { @@ -864,32 +859,32 @@ protected final int getLocalId(int _dim) { return (getLocalId(2)); } */ - /** - * Determine the size of the group that an executing kernel is a member of. - *+ * final int[] localIds = new int[1024]; + * Kernel kernel = new Kernel(){ + * public void run() { + * int gid = getGlobalId(); + * localIds[gid] = getLocalId(); + * } + * }; + * kernel.execute(localIds.length); + * for (int i=0; i< values.length; i++){ + * System.out.printf("%4d %4d\n", i, localIds[i]); + * } + *- * When a
Kernel.execute(int globalSize)is invoked for a particular kernel, the runtime will break the work into - * various 'groups'.getLocalSize()allows a kernel to determine the size of the current group. - *- * Note groups may not all be the same size. In particular, if
(global size)%(# of compute devices)!=0, the runtime can choose to dispatch kernels to - * groups with differing sizes. - * - * @see #getGroupId() - * @see #getGlobalId() - * @see #getGlobalSize() - * @see #getNumGroups() - * @see #getLocalSize() - * - * @return The size of the currently executing group. - */ - @OpenCLDelegate - protected final int getLocalSize() { - return kernelState.getRange().getLocalSize(0); - } - - @OpenCLDelegate - protected final int getLocalSize(int _dim) { - return kernelState.getRange().getLocalSize(_dim); - } + + /** + * Determine the size of the group that an executing kernel is a member of. + *+ * When a
Kernel.execute(int globalSize)is invoked for a particular kernel, the runtime will break the work into + * various 'groups'.getLocalSize()allows a kernel to determine the size of the current group. + *+ * Note groups may not all be the same size. In particular, if
(global size)%(# of compute devices)!=0, the runtime can choose to dispatch kernels to + * groups with differing sizes. + * + * @return The size of the currently executing group. + * @see #getGroupId() + * @see #getGlobalId() + * @see #getGlobalSize() + * @see #getNumGroups() + * @see #getLocalSize() + */ + @OpenCLDelegate + protected final int getLocalSize() { + return kernelState.getRange().getLocalSize(0); + } + + @OpenCLDelegate + protected final int getLocalSize(int _dim) { + return kernelState.getRange().getLocalSize(_dim); + } /* @OpenCLDelegate protected final int getLocalWidth() { @@ -904,25 +899,25 @@ protected final int getLocalSize(int _dim) { return (range.getLocalSize(2)); } */ - /** - * Determine the value that was passed toKernel.execute(int globalSize)method. - * - * @see #getGroupId() - * @see #getGlobalId() - * @see #getNumGroups() - * @see #getLocalSize() - * - * @return The value passed toKernel.execute(int globalSize)causing the current execution. - */ - @OpenCLDelegate - protected final int getGlobalSize() { - return kernelState.getRange().getGlobalSize(0); - } - - @OpenCLDelegate - protected final int getGlobalSize(int _dim) { - return kernelState.getRange().getGlobalSize(_dim); - } + + /** + * Determine the value that was passed toKernel.execute(int globalSize)method. + * + * @return The value passed toKernel.execute(int globalSize)causing the current execution. + * @see #getGroupId() + * @see #getGlobalId() + * @see #getNumGroups() + * @see #getLocalSize() + */ + @OpenCLDelegate + protected final int getGlobalSize() { + return kernelState.getRange().getGlobalSize(0); + } + + @OpenCLDelegate + protected final int getGlobalSize(int _dim) { + return kernelState.getRange().getGlobalSize(_dim); + } /* @OpenCLDelegate protected final int getGlobalWidth() { @@ -937,29 +932,29 @@ protected final int getGlobalSize(int _dim) { return (range.getGlobalSize(2)); } */ - /** - * Determine the number of groups that will be used to execute a kernel - *- * When
Kernel.execute(int globalSize)is invoked, the runtime will split the work into - * multiple 'groups'.getNumGroups()returns the total number of groups that will be used. - * - * @see #getGroupId() - * @see #getGlobalId() - * @see #getGlobalSize() - * @see #getNumGroups() - * @see #getLocalSize() - * - * @return The number of groups that kernels will be dispatched into. - */ - @OpenCLDelegate - protected final int getNumGroups() { - return kernelState.getRange().getNumGroups(0); - } - - @OpenCLDelegate - protected final int getNumGroups(int _dim) { - return kernelState.getRange().getNumGroups(_dim); - } + + /** + * Determine the number of groups that will be used to execute a kernel + *+ * When
Kernel.execute(int globalSize)is invoked, the runtime will split the work into + * multiple 'groups'.getNumGroups()returns the total number of groups that will be used. + * + * @return The number of groups that kernels will be dispatched into. + * @see #getGroupId() + * @see #getGlobalId() + * @see #getGlobalSize() + * @see #getNumGroups() + * @see #getLocalSize() + */ + @OpenCLDelegate + protected final int getNumGroups() { + return kernelState.getRange().getNumGroups(0); + } + + @OpenCLDelegate + protected final int getNumGroups(int _dim) { + return kernelState.getRange().getNumGroups(_dim); + } /* @OpenCLDelegate protected final int getNumGroupsWidth() { @@ -974,377 +969,365 @@ protected final int getNumGroups(int _dim) { return (range.getGroups(2)); } */ - /** - * The entry point of a kernel. - * - *- * Every kernel must override this method. - */ - public abstract void run(); - - /** False by default. In the event that all preferred devices fail to execute a kernel, it is possible to supply an alternate (possibly non-parallel) - * execution algorithm by overriding this method to return true, and overriding {@link #executeFallbackAlgorithm(Range, int)} with the alternate - * algorithm. - */ - public boolean hasFallbackAlgorithm() { - return false; - } - - /** If {@link #hasFallbackAlgorithm()} has been overriden to return true, this method should be overriden so as to - * apply a single pass of the kernel's logic to the entire _range. - * - *
- * This is not normally required, as fallback to {@link JavaDevice#THREAD_POOL} will implement the algorithm in parallel. However - * in the event that thread pool execution may be prohibitively slow, this method might implement a "quick and dirty" approximation - * to the desired result (for example, a simple box-blur as opposed to a gaussian blur in an image processing application). - */ - public void executeFallbackAlgorithm(Range _range, int _passId) { - // nothing - } - - /** - * Invoking this method flags that once the current pass is complete execution should be abandoned. Due to the complexity of intercommunication - * between java (or C) and executing OpenCL, this is the best we can do for general cancellation of execution at present. OpenCL 2.0 should introduce - * pipe mechanisms which will support mid-pass cancellation easily. - * - *
- * Note that in the case of thread-pool/pure java execution we could do better already, using Thread.interrupt() (and/or other means) to abandon - * execution mid-pass. However at present this is not attempted. - * - * @see #execute(int, int) - * @see #execute(Range, int) - * @see #execute(String, Range, int) - */ - public void cancelMultiPass() { - if (kernelRunner == null) { - return; - } - kernelRunner.cancelMultiPass(); - } - - public int getCancelState() { - return kernelRunner == null ? KernelRunner.CANCEL_STATUS_FALSE : kernelRunner.getCancelState(); - } - - /** - * @see KernelRunner#getCurrentPass() - */ - public int getCurrentPass() { - if (kernelRunner == null) { - return KernelRunner.PASS_ID_COMPLETED_EXECUTION; - } - return kernelRunner.getCurrentPass(); - } - - /** - * @see KernelRunner#isExecuting() - */ - public boolean isExecuting() { - if (kernelRunner == null) { - return false; - } - return kernelRunner.isExecuting(); - } - - /** - * When using a Java Thread Pool Aparapi uses clone to copy the initial instance to each thread. - * - *
- * If you choose to override
clone()you are responsible for delegating tosuper.clone();- */ - @Override - public Kernel clone() { - try { - final Kernel worker = (Kernel) super.clone(); - - // We need to be careful to also clone the KernelState - worker.kernelState = worker.new KernelState(kernelState); // Qualified copy constructor - - worker.kernelState.setGroupIds(new int[] {0, 0, 0}); - - worker.kernelState.setLocalIds(new int[] {0, 0, 0}); - - worker.kernelState.setGlobalIds(new int[] {0, 0, 0}); - - return worker; - } catch (final CloneNotSupportedException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - return (null); - } - } - /** - * Delegates to either {@link java.lang.Math#acos(double)} (Java) oracos(float)(OpenCL). - * + /** + * The entry point of a kernel. + *+ *
+ * Every kernel must override this method. + */ + public abstract void run(); + + /** + * False by default. In the event that all preferred devices fail to execute a kernel, it is possible to supply an alternate (possibly non-parallel) + * execution algorithm by overriding this method to return true, and overriding {@link #executeFallbackAlgorithm(Range, int)} with the alternate + * algorithm. + */ + public boolean hasFallbackAlgorithm() { + return false; + } + + /** + * If {@link #hasFallbackAlgorithm()} has been overriden to return true, this method should be overriden so as to + * apply a single pass of the kernel's logic to the entire _range. + *
+ *
+ * This is not normally required, as fallback to {@link JavaDevice#THREAD_POOL} will implement the algorithm in parallel. However + * in the event that thread pool execution may be prohibitively slow, this method might implement a "quick and dirty" approximation + * to the desired result (for example, a simple box-blur as opposed to a gaussian blur in an image processing application). + */ + public void executeFallbackAlgorithm(Range _range, int _passId) { + // nothing + } + +// /** +// * Invoking this method flags that once the current pass is complete execution should be abandoned. Due to the complexity of intercommunication +// * between java (or C) and executing OpenCL, this is the best we can do for general cancellation of execution at present. OpenCL 2.0 should introduce +// * pipe mechanisms which will support mid-pass cancellation easily. +// *
+// *
+// * Note that in the case of thread-pool/pure java execution we could do better already, using Thread.interrupt() (and/or other means) to abandon +// * execution mid-pass. However at present this is not attempted. +// * +// * @see #execute(int, int) +// * @see #execute(Range, int) +// * @see #execute(String, Range, int) +// */ +// public void cancelMultiPass() { +// if (kernelRunner == null) { +// return; +// } +// kernelRunner.cancelMultiPass(); +// } +// +// public int getCancelState() { +// return kernelRunner == null ? KernelRunner.CANCEL_STATUS_FALSE : kernelRunner.getCancelState(); +// } +// +// /** +// * @see KernelRunner#getCurrentPass() +// */ +// public int getCurrentPass() { +// if (kernelRunner == null) { +// return KernelRunner.PASS_ID_COMPLETED_EXECUTION; +// } +// return kernelRunner.getCurrentPass(); +// } +// +// /** +// * @see KernelRunner#isExecuting() +// */ +// public boolean isExecuting() { +// if (kernelRunner == null) { +// return false; +// } +// return kernelRunner.isExecuting(); +// } + + /** + * When using a Java Thread Pool Aparapi uses clone to copy the initial instance to each thread. + *
+ *
+ * If you choose to override
clone()you are responsible for delegating tosuper.clone();+ */ + @Override + public Kernel clone() { + try { + final Kernel worker = (Kernel) super.clone(); + + // We need to be careful to also clone the KernelState + worker.kernelState = worker.new KernelState(kernelState); // Qualified copy constructor + + worker.kernelState.setGroupIds(new int[]{0, 0, 0}); + + worker.kernelState.setLocalIds(new int[]{0, 0, 0}); + + worker.kernelState.setGlobalIds(new int[]{0, 0, 0}); + + return worker; + } catch (final CloneNotSupportedException e) { + // TODO Auto-generated catch block + e.printStackTrace(); + return (null); + } + } + + /** + * Delegates to either {@link java.lang.Math#acos(double)} (Java) oracos(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param a value to delegate to {@link java.lang.Math#acos(double)}/
acos(float)* @return {@link java.lang.Math#acos(double)} casted to float/acos(float)- * * @see java.lang.Math#acos(double) * @seeacos(float)*/ - @OpenCLMapping(mapTo = "acos") - protected float acos(float a) { - return (float) Math.acos(a); - } - - /** - * Delegates to either {@link java.lang.Math#acos(double)} (Java) oracos(double)(OpenCL). - * - * User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. - * - * @param a value to delegate to {@link java.lang.Math#acos(double)}/acos(double)- * @return {@link java.lang.Math#acos(double)}/acos(double)- * - * @see java.lang.Math#acos(double) - * @seeacos(double)- */ - @OpenCLMapping(mapTo = "acos") - protected double acos(double a) { - return Math.acos(a); - } - - /** - * Delegates to either {@link java.lang.Math#asin(double)} (Java) orasin(float)(OpenCL). + @OpenCLMapping(mapTo = "acos") + protected float acos(float a) { + return (float) Math.acos(a); + } + + /** + * Delegates to either {@link java.lang.Math#acos(double)} (Java) oracos(double)(OpenCL). + *+ * User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * + * @param a value to delegate to {@link java.lang.Math#acos(double)}/
acos(double)+ * @return {@link java.lang.Math#acos(double)}/acos(double)+ * @see java.lang.Math#acos(double) + * @seeacos(double)+ */ + @OpenCLMapping(mapTo = "acos") + protected double acos(double a) { + return Math.acos(a); + } + + /** + * Delegates to either {@link java.lang.Math#asin(double)} (Java) orasin(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#asin(double)}/
asin(float)* @return {@link java.lang.Math#asin(double)} casted to float/asin(float)- * * @see java.lang.Math#asin(double) * @seeasin(float)*/ - @OpenCLMapping(mapTo = "asin") - protected float asin(float _f) { - return (float) Math.asin(_f); - } + @OpenCLMapping(mapTo = "asin") + protected float asin(float _f) { + return (float) Math.asin(_f); + } - /** - * Delegates to either {@link java.lang.Math#asin(double)} (Java) orasin(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#asin(double)} (Java) orasin(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#asin(double)}/
asin(double)* @return {@link java.lang.Math#asin(double)}/asin(double)- * * @see java.lang.Math#asin(double) * @seeasin(double)*/ - @OpenCLMapping(mapTo = "asin") - protected double asin(double _d) { - return Math.asin(_d); - } + @OpenCLMapping(mapTo = "asin") + protected double asin(double _d) { + return Math.asin(_d); + } - /** - * Delegates to either {@link java.lang.Math#atan(double)} (Java) oratan(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#atan(double)} (Java) oratan(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#atan(double)}/
atan(float)* @return {@link java.lang.Math#atan(double)} casted to float/atan(float)- * * @see java.lang.Math#atan(double) * @seeatan(float)*/ - @OpenCLMapping(mapTo = "atan") - protected float atan(float _f) { - return (float) Math.atan(_f); - } + @OpenCLMapping(mapTo = "atan") + protected float atan(float _f) { + return (float) Math.atan(_f); + } - /** - * Delegates to either {@link java.lang.Math#atan(double)} (Java) oratan(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#atan(double)} (Java) oratan(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#atan(double)}/
atan(double)* @return {@link java.lang.Math#atan(double)}/atan(double)- * * @see java.lang.Math#atan(double) * @seeatan(double)*/ - @OpenCLMapping(mapTo = "atan") - protected double atan(double _d) { - return Math.atan(_d); - } + @OpenCLMapping(mapTo = "atan") + protected double atan(double _d) { + return Math.atan(_d); + } - /** - * Delegates to either {@link java.lang.Math#atan2(double, double)} (Java) oratan2(float, float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#atan2(double, double)} (Java) oratan2(float, float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f1 value to delegate to first argument of {@link java.lang.Math#atan2(double, double)}/
atan2(float, float)* @param _f2 value to delegate to second argument of {@link java.lang.Math#atan2(double, double)}/atan2(float, float)* @return {@link java.lang.Math#atan2(double, double)} casted to float/atan2(float, float)- * * @see java.lang.Math#atan2(double, double) * @seeatan2(float, float)*/ - @OpenCLMapping(mapTo = "atan2") - protected float atan2(float _f1, float _f2) { - return (float) Math.atan2(_f1, _f2); - } + @OpenCLMapping(mapTo = "atan2") + protected float atan2(float _f1, float _f2) { + return (float) Math.atan2(_f1, _f2); + } - /** - * Delegates to either {@link java.lang.Math#atan2(double, double)} (Java) oratan2(double, double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#atan2(double, double)} (Java) oratan2(double, double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d1 value to delegate to first argument of {@link java.lang.Math#atan2(double, double)}/
atan2(double, double)* @param _d2 value to delegate to second argument of {@link java.lang.Math#atan2(double, double)}/atan2(double, double)* @return {@link java.lang.Math#atan2(double, double)}/atan2(double, double)- * * @see java.lang.Math#atan2(double, double) * @seeatan2(double, double)*/ - @OpenCLMapping(mapTo = "atan2") - protected double atan2(double _d1, double _d2) { - return Math.atan2(_d1, _d2); - } + @OpenCLMapping(mapTo = "atan2") + protected double atan2(double _d1, double _d2) { + return Math.atan2(_d1, _d2); + } - /** - * Delegates to either {@link java.lang.Math#ceil(double)} (Java) orceil(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#ceil(double)} (Java) orceil(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#ceil(double)}/
ceil(float)* @return {@link java.lang.Math#ceil(double)} casted to float/ceil(float)- * * @see java.lang.Math#ceil(double) * @seeceil(float)*/ - @OpenCLMapping(mapTo = "ceil") - protected float ceil(float _f) { - return (float) Math.ceil(_f); - } + @OpenCLMapping(mapTo = "ceil") + protected float ceil(float _f) { + return (float) Math.ceil(_f); + } - /** - * Delegates to either {@link java.lang.Math#ceil(double)} (Java) orceil(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#ceil(double)} (Java) orceil(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#ceil(double)}/
ceil(double)* @return {@link java.lang.Math#ceil(double)}/ceil(double)- * * @see java.lang.Math#ceil(double) * @seeceil(double)*/ - @OpenCLMapping(mapTo = "ceil") - protected double ceil(double _d) { - return Math.ceil(_d); - } + @OpenCLMapping(mapTo = "ceil") + protected double ceil(double _d) { + return Math.ceil(_d); + } - /** - * Delegates to either {@link java.lang.Math#cos(double)} (Java) orcos(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#cos(double)} (Java) orcos(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#cos(double)}/
cos(float)* @return {@link java.lang.Math#cos(double)} casted to float/cos(float)- * * @see java.lang.Math#cos(double) * @seecos(float)*/ - @OpenCLMapping(mapTo = "cos") - protected float cos(float _f) { - return (float) Math.cos(_f); - } + @OpenCLMapping(mapTo = "cos") + protected float cos(float _f) { + return (float) Math.cos(_f); + } - /** - * Delegates to either {@link java.lang.Math#cos(double)} (Java) orcos(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#cos(double)} (Java) orcos(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#cos(double)}/
cos(double)* @return {@link java.lang.Math#cos(double)}/cos(double)- * * @see java.lang.Math#cos(double) * @seecos(double)*/ - @OpenCLMapping(mapTo = "cos") - protected double cos(double _d) { - return Math.cos(_d); - } + @OpenCLMapping(mapTo = "cos") + protected double cos(double _d) { + return Math.cos(_d); + } - /** - * Delegates to either {@link java.lang.Math#exp(double)} (Java) orexp(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#exp(double)} (Java) orexp(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#exp(double)}/
exp(float)* @return {@link java.lang.Math#exp(double)} casted to float/exp(float)- * * @see java.lang.Math#exp(double) * @seeexp(float)*/ - @OpenCLMapping(mapTo = "exp") - protected float exp(float _f) { - return (float) Math.exp(_f); - } + @OpenCLMapping(mapTo = "exp") + protected float exp(float _f) { + return (float) Math.exp(_f); + } - /** - * Delegates to either {@link java.lang.Math#exp(double)} (Java) orexp(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#exp(double)} (Java) orexp(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#exp(double)}/
exp(double)* @return {@link java.lang.Math#exp(double)}/exp(double)- * * @see java.lang.Math#exp(double) * @seeexp(double)*/ - @OpenCLMapping(mapTo = "exp") - protected double exp(double _d) { - return Math.exp(_d); - } + @OpenCLMapping(mapTo = "exp") + protected double exp(double _d) { + return Math.exp(_d); + } - /** - * Delegates to either {@link java.lang.Math#abs(float)} (Java) orfabs(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#abs(float)} (Java) orfabs(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#abs(float)}/
fabs(float)* @return {@link java.lang.Math#abs(float)}/fabs(float)- * * @see java.lang.Math#abs(float) * @seefabs(float)*/ - @OpenCLMapping(mapTo = "fabs") - protected float abs(float _f) { - return Math.abs(_f); - } + @OpenCLMapping(mapTo = "fabs") + protected float abs(float _f) { + return Math.abs(_f); + } - /** - * Delegates to either {@link java.lang.Integer#bitCount(int)} (Java) orpopcount(int)(OpenCL). + /** + * Delegates to either {@link java.lang.Integer#bitCount(int)} (Java) orpopcount(int)(OpenCL). * * @param _i value to delegate to {@link java.lang.Integer#bitCount(int)}/popcount(int)* @return {@link java.lang.Integer#bitCount(int)}/popcount(int)* @see java.lang.Integer#bitCount(int) * @seepopcount(int)*/ - @OpenCLMapping(mapTo = "popcount") - protected int popcount(int _i) { - return Integer.bitCount(_i); - } + @OpenCLMapping(mapTo = "popcount") + protected int popcount(int _i) { + return Integer.bitCount(_i); + } - /** - * Delegates to either {@link java.lang.Long#bitCount(long)} (Java) orpopcount(long)(OpenCL). + /** + * Delegates to either {@link java.lang.Long#bitCount(long)} (Java) orpopcount(long)(OpenCL). * * @param _i value to delegate to {@link java.lang.Long#bitCount(long)}/popcount(long)* @return {@link java.lang.Long#bitCount(long)}/popcount(long)* @see java.lang.Long#bitCount(long) * @seepopcount(long)*/ - @OpenCLMapping(mapTo = "popcount") - protected long popcount(long _i) { - return Long.bitCount(_i); - } + @OpenCLMapping(mapTo = "popcount") + protected long popcount(long _i) { + return Long.bitCount(_i); + } - /** - * Delegates to either {@link java.lang.Integer#numberOfLeadingZeros(int)} (Java) orclz(int)(OpenCL). + /** + * Delegates to either {@link java.lang.Integer#numberOfLeadingZeros(int)} (Java) orclz(int)(OpenCL). * * @param _i value to delegate to {@link java.lang.Integer#numberOfLeadingZeros(int)}/clz(int)* @return {@link java.lang.Integer#numberOfLeadingZeros(int)}/clz(int)@@ -1352,566 +1335,531 @@ protected long popcount(long _i) { * @seeclz(int)*/ - @OpenCLMapping(mapTo = "clz") - protected int clz(int _i) { - return Integer.numberOfLeadingZeros(_i); - } + @OpenCLMapping(mapTo = "clz") + protected int clz(int _i) { + return Integer.numberOfLeadingZeros(_i); + } - /** - * Delegates to either {@link java.lang.Long#numberOfLeadingZeros(long)} (Java) orclz(long)(OpenCL). + /** + * Delegates to either {@link java.lang.Long#numberOfLeadingZeros(long)} (Java) orclz(long)(OpenCL). * * @param _l value to delegate to {@link java.lang.Long#numberOfLeadingZeros(long)}/clz(long)* @return {@link java.lang.Long#numberOfLeadingZeros(long)}/clz(long)* @see java.lang.Long#numberOfLeadingZeros(long) * @seeclz(long)*/ - - - @OpenCLMapping(mapTo = "clz") - protected long clz(long _l) { - return Long.numberOfLeadingZeros(_l); - } + @OpenCLMapping(mapTo = "clz") + protected long clz(long _l) { + return Long.numberOfLeadingZeros(_l); + } /** - * Delegates to either {@link java.lang.Math#abs(double)} (Java) orfabs(double)(OpenCL). - * + * Delegates to either {@link java.lang.Math#abs(double)} (Java) orfabs(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#abs(double)}/
fabs(double)* @return {@link java.lang.Math#abs(double)}/fabs(double)- * * @see java.lang.Math#abs(double) * @seefabs(double)*/ - @OpenCLMapping(mapTo = "fabs") - protected double abs(double _d) { - return Math.abs(_d); - } + @OpenCLMapping(mapTo = "fabs") + protected double abs(double _d) { + return Math.abs(_d); + } - /** - * Delegates to either {@link java.lang.Math#abs(int)} (Java) orabs(int)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#abs(int)} (Java) orabs(int)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param n value to delegate to {@link java.lang.Math#abs(int)}/
abs(int)* @return {@link java.lang.Math#abs(int)}/abs(int)- * * @see java.lang.Math#abs(int) * @seeabs(int)*/ - @OpenCLMapping(mapTo = "abs") - protected int abs(int n) { - return Math.abs(n); - } + @OpenCLMapping(mapTo = "abs") + protected int abs(int n) { + return Math.abs(n); + } - /** - * Delegates to either {@link java.lang.Math#abs(long)} (Java) orabs(long)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#abs(long)} (Java) orabs(long)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param n value to delegate to {@link java.lang.Math#abs(long)}/
abs(long)* @return {@link java.lang.Math#abs(long)}/abs(long)- * * @see java.lang.Math#abs(long) * @seeabs(long)*/ - @OpenCLMapping(mapTo = "abs") - protected long abs(long n) { - return Math.abs(n); - } + @OpenCLMapping(mapTo = "abs") + protected long abs(long n) { + return Math.abs(n); + } - /** - * Delegates to either {@link java.lang.Math#floor(double)} (Java) orfloor(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#floor(double)} (Java) orfloor(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#floor(double)}/
floor(float)* @return {@link java.lang.Math#floor(double)} casted to float/floor(float)- * * @see java.lang.Math#floor(double) * @seefloor(float)*/ - @OpenCLMapping(mapTo = "floor") - protected float floor(float _f) { - return (float) Math.floor(_f); - } + @OpenCLMapping(mapTo = "floor") + protected float floor(float _f) { + return (float) Math.floor(_f); + } - /** - * Delegates to either {@link java.lang.Math#floor(double)} (Java) orfloor(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#floor(double)} (Java) orfloor(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#floor(double)}/
floor(double)* @return {@link java.lang.Math#floor(double)}/floor(double)- * * @see java.lang.Math#floor(double) * @seefloor(double)*/ - @OpenCLMapping(mapTo = "floor") - protected double floor(double _d) { - return Math.floor(_d); - } + @OpenCLMapping(mapTo = "floor") + protected double floor(double _d) { + return Math.floor(_d); + } - /** - * Delegates to either {@link java.lang.Math#max(float, float)} (Java) orfmax(float, float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#max(float, float)} (Java) orfmax(float, float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f1 value to delegate to first argument of {@link java.lang.Math#max(float, float)}/
fmax(float, float)* @param _f2 value to delegate to second argument of {@link java.lang.Math#max(float, float)}/fmax(float, float)* @return {@link java.lang.Math#max(float, float)}/fmax(float, float)- * * @see java.lang.Math#max(float, float) * @seefmax(float, float)*/ - @OpenCLMapping(mapTo = "fmax") - protected float max(float _f1, float _f2) { - return Math.max(_f1, _f2); - } + @OpenCLMapping(mapTo = "fmax") + protected float max(float _f1, float _f2) { + return Math.max(_f1, _f2); + } - /** - * Delegates to either {@link java.lang.Math#max(double, double)} (Java) orfmax(double, double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#max(double, double)} (Java) orfmax(double, double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d1 value to delegate to first argument of {@link java.lang.Math#max(double, double)}/
fmax(double, double)* @param _d2 value to delegate to second argument of {@link java.lang.Math#max(double, double)}/fmax(double, double)* @return {@link java.lang.Math#max(double, double)}/fmax(double, double)- * * @see java.lang.Math#max(double, double) * @seefmax(double, double)*/ - @OpenCLMapping(mapTo = "fmax") - protected double max(double _d1, double _d2) { - return Math.max(_d1, _d2); - } + @OpenCLMapping(mapTo = "fmax") + protected double max(double _d1, double _d2) { + return Math.max(_d1, _d2); + } - /** - * Delegates to either {@link java.lang.Math#max(int, int)} (Java) ormax(int, int)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#max(int, int)} (Java) ormax(int, int)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param n1 value to delegate to {@link java.lang.Math#max(int, int)}/
max(int, int)* @param n2 value to delegate to {@link java.lang.Math#max(int, int)}/max(int, int)* @return {@link java.lang.Math#max(int, int)}/max(int, int)- * * @see java.lang.Math#max(int, int) * @seemax(int, int)*/ - @OpenCLMapping(mapTo = "max") - protected int max(int n1, int n2) { - return Math.max(n1, n2); - } + @OpenCLMapping(mapTo = "max") + protected int max(int n1, int n2) { + return Math.max(n1, n2); + } - /** - * Delegates to either {@link java.lang.Math#max(long, long)} (Java) ormax(long, long)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#max(long, long)} (Java) ormax(long, long)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param n1 value to delegate to first argument of {@link java.lang.Math#max(long, long)}/
max(long, long)* @param n2 value to delegate to second argument of {@link java.lang.Math#max(long, long)}/max(long, long)* @return {@link java.lang.Math#max(long, long)}/max(long, long)- * * @see java.lang.Math#max(long, long) * @seemax(long, long)*/ - @OpenCLMapping(mapTo = "max") - protected long max(long n1, long n2) { - return Math.max(n1, n2); - } + @OpenCLMapping(mapTo = "max") + protected long max(long n1, long n2) { + return Math.max(n1, n2); + } - /** - * Delegates to either {@link java.lang.Math#min(float, float)} (Java) orfmin(float, float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#min(float, float)} (Java) orfmin(float, float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f1 value to delegate to first argument of {@link java.lang.Math#min(float, float)}/
fmin(float, float)* @param _f2 value to delegate to second argument of {@link java.lang.Math#min(float, float)}/fmin(float, float)* @return {@link java.lang.Math#min(float, float)}/fmin(float, float)- * * @see java.lang.Math#min(float, float) * @seefmin(float, float)*/ - @OpenCLMapping(mapTo = "fmin") - protected float min(float _f1, float _f2) { - return Math.min(_f1, _f2); - } + @OpenCLMapping(mapTo = "fmin") + protected float min(float _f1, float _f2) { + return Math.min(_f1, _f2); + } - /** - * Delegates to either {@link java.lang.Math#min(double, double)} (Java) orfmin(double, double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#min(double, double)} (Java) orfmin(double, double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d1 value to delegate to first argument of {@link java.lang.Math#min(double, double)}/
fmin(double, double)* @param _d2 value to delegate to second argument of {@link java.lang.Math#min(double, double)}/fmin(double, double)* @return {@link java.lang.Math#min(double, double)}/fmin(double, double)- * * @see java.lang.Math#min(double, double) * @seefmin(double, double)*/ - @OpenCLMapping(mapTo = "fmin") - protected double min(double _d1, double _d2) { - return Math.min(_d1, _d2); - } + @OpenCLMapping(mapTo = "fmin") + protected double min(double _d1, double _d2) { + return Math.min(_d1, _d2); + } - /** - * Delegates to either {@link java.lang.Math#min(int, int)} (Java) ormin(int, int)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#min(int, int)} (Java) ormin(int, int)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param n1 value to delegate to first argument of {@link java.lang.Math#min(int, int)}/
min(int, int)* @param n2 value to delegate to second argument of {@link java.lang.Math#min(int, int)}/min(int, int)* @return {@link java.lang.Math#min(int, int)}/min(int, int)- * * @see java.lang.Math#min(int, int) * @seemin(int, int)*/ - @OpenCLMapping(mapTo = "min") - protected int min(int n1, int n2) { - return Math.min(n1, n2); - } + @OpenCLMapping(mapTo = "min") + protected int min(int n1, int n2) { + return Math.min(n1, n2); + } - /** - * Delegates to either {@link java.lang.Math#min(long, long)} (Java) ormin(long, long)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#min(long, long)} (Java) ormin(long, long)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param n1 value to delegate to first argument of {@link java.lang.Math#min(long, long)}/
min(long, long)* @param n2 value to delegate to second argument of {@link java.lang.Math#min(long, long)}/min(long, long)* @return {@link java.lang.Math#min(long, long)}/min(long, long)- * * @see java.lang.Math#min(long, long) * @seemin(long, long)*/ - @OpenCLMapping(mapTo = "min") - protected long min(long n1, long n2) { - return Math.min(n1, n2); - } + @OpenCLMapping(mapTo = "min") + protected long min(long n1, long n2) { + return Math.min(n1, n2); + } - /** - * Delegates to either {@link java.lang.Math#log(double)} (Java) orlog(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#log(double)} (Java) orlog(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#log(double)}/
log(float)* @return {@link java.lang.Math#log(double)} casted to float/log(float)- * * @see java.lang.Math#log(double) * @seelog(float)*/ - @OpenCLMapping(mapTo = "log") - protected float log(float _f) { - return (float) Math.log(_f); - } + @OpenCLMapping(mapTo = "log") + protected float log(float _f) { + return (float) Math.log(_f); + } - /** - * Delegates to either {@link java.lang.Math#log(double)} (Java) orlog(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#log(double)} (Java) orlog(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#log(double)}/
log(double)* @return {@link java.lang.Math#log(double)}/log(double)- * * @see java.lang.Math#log(double) * @seelog(double)*/ - @OpenCLMapping(mapTo = "log") - protected double log(double _d) { - return Math.log(_d); - } + @OpenCLMapping(mapTo = "log") + protected double log(double _d) { + return Math.log(_d); + } - /** - * Delegates to either {@link java.lang.Math#pow(double, double)} (Java) orpow(float, float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#pow(double, double)} (Java) orpow(float, float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f1 value to delegate to first argument of {@link java.lang.Math#pow(double, double)}/
pow(float, float)* @param _f2 value to delegate to second argument of {@link java.lang.Math#pow(double, double)}/pow(float, float)* @return {@link java.lang.Math#pow(double, double)} casted to float/pow(float, float)- * * @see java.lang.Math#pow(double, double) * @seepow(float, float)*/ - @OpenCLMapping(mapTo = "pow") - protected float pow(float _f1, float _f2) { - return (float) Math.pow(_f1, _f2); - } + @OpenCLMapping(mapTo = "pow") + protected float pow(float _f1, float _f2) { + return (float) Math.pow(_f1, _f2); + } - /** - * Delegates to either {@link java.lang.Math#pow(double, double)} (Java) orpow(double, double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#pow(double, double)} (Java) orpow(double, double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d1 value to delegate to first argument of {@link java.lang.Math#pow(double, double)}/
pow(double, double)* @param _d2 value to delegate to second argument of {@link java.lang.Math#pow(double, double)}/pow(double, double)* @return {@link java.lang.Math#pow(double, double)}/pow(double, double)- * * @see java.lang.Math#pow(double, double) * @seepow(double, double)*/ - @OpenCLMapping(mapTo = "pow") - protected double pow(double _d1, double _d2) { - return Math.pow(_d1, _d2); - } + @OpenCLMapping(mapTo = "pow") + protected double pow(double _d1, double _d2) { + return Math.pow(_d1, _d2); + } - /** - * Delegates to either {@link java.lang.Math#IEEEremainder(double, double)} (Java) orremainder(float, float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#IEEEremainder(double, double)} (Java) orremainder(float, float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f1 value to delegate to first argument of {@link java.lang.Math#IEEEremainder(double, double)}/
remainder(float, float)* @param _f2 value to delegate to second argument of {@link java.lang.Math#IEEEremainder(double, double)}/remainder(float, float)* @return {@link java.lang.Math#IEEEremainder(double, double)} casted to float/remainder(float, float)- * * @see java.lang.Math#IEEEremainder(double, double) * @seeremainder(float, float)*/ - @OpenCLMapping(mapTo = "remainder") - protected float IEEEremainder(float _f1, float _f2) { - return (float) Math.IEEEremainder(_f1, _f2); - } + @OpenCLMapping(mapTo = "remainder") + protected float IEEEremainder(float _f1, float _f2) { + return (float) Math.IEEEremainder(_f1, _f2); + } - /** - * Delegates to either {@link java.lang.Math#IEEEremainder(double, double)} (Java) orremainder(double, double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#IEEEremainder(double, double)} (Java) orremainder(double, double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d1 value to delegate to first argument of {@link java.lang.Math#IEEEremainder(double, double)}/
remainder(double, double)* @param _d2 value to delegate to second argument of {@link java.lang.Math#IEEEremainder(double, double)}/remainder(double, double)* @return {@link java.lang.Math#IEEEremainder(double, double)}/remainder(double, double)- * * @see java.lang.Math#IEEEremainder(double, double) * @seeremainder(double, double)*/ - @OpenCLMapping(mapTo = "remainder") - protected double IEEEremainder(double _d1, double _d2) { - return Math.IEEEremainder(_d1, _d2); - } + @OpenCLMapping(mapTo = "remainder") + protected double IEEEremainder(double _d1, double _d2) { + return Math.IEEEremainder(_d1, _d2); + } - /** - * Delegates to either {@link java.lang.Math#toRadians(double)} (Java) orradians(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#toRadians(double)} (Java) orradians(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#toRadians(double)}/
radians(float)* @return {@link java.lang.Math#toRadians(double)} casted to float/radians(float)- * * @see java.lang.Math#toRadians(double) * @seeradians(float)*/ - @OpenCLMapping(mapTo = "radians") - protected float toRadians(float _f) { - return (float) Math.toRadians(_f); - } + @OpenCLMapping(mapTo = "radians") + protected float toRadians(float _f) { + return (float) Math.toRadians(_f); + } - /** - * Delegates to either {@link java.lang.Math#toRadians(double)} (Java) orradians(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#toRadians(double)} (Java) orradians(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#toRadians(double)}/
radians(double)* @return {@link java.lang.Math#toRadians(double)}/radians(double)- * * @see java.lang.Math#toRadians(double) * @seeradians(double)*/ - @OpenCLMapping(mapTo = "radians") - protected double toRadians(double _d) { - return Math.toRadians(_d); - } + @OpenCLMapping(mapTo = "radians") + protected double toRadians(double _d) { + return Math.toRadians(_d); + } - /** - * Delegates to either {@link java.lang.Math#toDegrees(double)} (Java) ordegrees(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#toDegrees(double)} (Java) ordegrees(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#toDegrees(double)}/
degrees(float)* @return {@link java.lang.Math#toDegrees(double)} casted to float/degrees(float)- * * @see java.lang.Math#toDegrees(double) * @seedegrees(float)*/ - @OpenCLMapping(mapTo = "degrees") - protected float toDegrees(float _f) { - return (float) Math.toDegrees(_f); - } + @OpenCLMapping(mapTo = "degrees") + protected float toDegrees(float _f) { + return (float) Math.toDegrees(_f); + } - /** - * Delegates to either {@link java.lang.Math#toDegrees(double)} (Java) ordegrees(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#toDegrees(double)} (Java) ordegrees(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#toDegrees(double)}/
degrees(double)* @return {@link java.lang.Math#toDegrees(double)}/degrees(double)- * * @see java.lang.Math#toDegrees(double) * @seedegrees(double)*/ - @OpenCLMapping(mapTo = "degrees") - protected double toDegrees(double _d) { - return Math.toDegrees(_d); - } + @OpenCLMapping(mapTo = "degrees") + protected double toDegrees(double _d) { + return Math.toDegrees(_d); + } - /** - * Delegates to either {@link java.lang.Math#rint(double)} (Java) orrint(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#rint(double)} (Java) orrint(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#rint(double)}/
rint(float)* @return {@link java.lang.Math#rint(double)} casted to float/rint(float)- * * @see java.lang.Math#rint(double) * @seerint(float)*/ - @OpenCLMapping(mapTo = "rint") - protected float rint(float _f) { - return (float) Math.rint(_f); - } + @OpenCLMapping(mapTo = "rint") + protected float rint(float _f) { + return (float) Math.rint(_f); + } - /** - * Delegates to either {@link java.lang.Math#rint(double)} (Java) orrint(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#rint(double)} (Java) orrint(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#rint(double)}/
rint(double)* @return {@link java.lang.Math#rint(double)}/rint(double)- * * @see java.lang.Math#rint(double) * @seerint(double)*/ - @OpenCLMapping(mapTo = "rint") - protected double rint(double _d) { - return Math.rint(_d); - } + @OpenCLMapping(mapTo = "rint") + protected double rint(double _d) { + return Math.rint(_d); + } - /** - * Delegates to either {@link java.lang.Math#round(float)} (Java) orround(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#round(float)} (Java) orround(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#round(float)}/
round(float)* @return {@link java.lang.Math#round(float)}/round(float)- * * @see java.lang.Math#round(float) * @seeround(float)*/ - @OpenCLMapping(mapTo = "round") - protected int round(float _f) { - return Math.round(_f); - } + @OpenCLMapping(mapTo = "round") + protected int round(float _f) { + return Math.round(_f); + } - /** - * Delegates to either {@link java.lang.Math#round(double)} (Java) orround(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#round(double)} (Java) orround(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#round(double)}/
round(double)* @return {@link java.lang.Math#round(double)}/round(double)- * * @see java.lang.Math#round(double) * @seeround(double)*/ - @OpenCLMapping(mapTo = "round") - protected long round(double _d) { - return Math.round(_d); - } + @OpenCLMapping(mapTo = "round") + protected long round(double _d) { + return Math.round(_d); + } - /** - * Delegates to either {@link java.lang.Math#sin(double)} (Java) orsin(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#sin(double)} (Java) orsin(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#sin(double)}/
sin(float)* @return {@link java.lang.Math#sin(double)} casted to float/sin(float)- * * @see java.lang.Math#sin(double) * @seesin(float)*/ - @OpenCLMapping(mapTo = "sin") - protected float sin(float _f) { - return (float) Math.sin(_f); - } + @OpenCLMapping(mapTo = "sin") + protected float sin(float _f) { + return (float) Math.sin(_f); + } - /** - * Delegates to either {@link java.lang.Math#sin(double)} (Java) orsin(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#sin(double)} (Java) orsin(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#sin(double)}/
sin(double)* @return {@link java.lang.Math#sin(double)}/sin(double)- * * @see java.lang.Math#sin(double) * @seesin(double)*/ - @OpenCLMapping(mapTo = "sin") - protected double sin(double _d) { - return Math.sin(_d); - } + @OpenCLMapping(mapTo = "sin") + protected double sin(double _d) { + return Math.sin(_d); + } - /** - * Delegates to either {@link java.lang.Math#sqrt(double)} (Java) orsqrt(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#sqrt(double)} (Java) orsqrt(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#sqrt(double)}/
sqrt(float)* @return {@link java.lang.Math#sqrt(double)} casted to float/sqrt(float)- * * @see java.lang.Math#sqrt(double) * @seesqrt(float)*/ - @OpenCLMapping(mapTo = "sqrt") - protected float sqrt(float _f) { - return (float) Math.sqrt(_f); - } + @OpenCLMapping(mapTo = "sqrt") + protected float sqrt(float _f) { + return (float) Math.sqrt(_f); + } - /** - * Delegates to either {@link java.lang.Math#sqrt(double)} (Java) orsqrt(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#sqrt(double)} (Java) orsqrt(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#sqrt(double)}/
sqrt(double)* @return {@link java.lang.Math#sqrt(double)}/sqrt(double)- * * @see java.lang.Math#sqrt(double) * @seesqrt(double)*/ - @OpenCLMapping(mapTo = "sqrt") - protected double sqrt(double _d) { - return Math.sqrt(_d); - } + @OpenCLMapping(mapTo = "sqrt") + protected double sqrt(double _d) { + return Math.sqrt(_d); + } - /** - * Delegates to either {@link java.lang.Math#tan(double)} (Java) ortan(float)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#tan(double)} (Java) ortan(float)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#tan(double)}/
tan(float)* @return {@link java.lang.Math#tan(double)} casted to float/tan(float)- * * @see java.lang.Math#tan(double) * @seetan(float)*/ - @OpenCLMapping(mapTo = "tan") - protected float tan(float _f) { - return (float) Math.tan(_f); - } + @OpenCLMapping(mapTo = "tan") + protected float tan(float _f) { + return (float) Math.tan(_f); + } - /** - * Delegates to either {@link java.lang.Math#tan(double)} (Java) ortan(double)(OpenCL). - * + /** + * Delegates to either {@link java.lang.Math#tan(double)} (Java) ortan(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#tan(double)}/
tan(double)* @return {@link java.lang.Math#tan(double)}/tan(double)- * * @see java.lang.Math#tan(double) * @seetan(double)*/ - @OpenCLMapping(mapTo = "tan") - protected double tan(double _d) { - return Math.tan(_d); - } + @OpenCLMapping(mapTo = "tan") + protected double tan(double _d) { + return Math.tan(_d); + } private static final double LOG_2_RECIPROCAL = 1.0D / Math.log(2.0D); private static final double PI_RECIPROCAL = 1.0D / Math.PI; @@ -1923,7 +1871,7 @@ protected final double acospi(final double a) { @OpenCLMapping(mapTo = "acospi") protected final float acospi(final float a) { - return (float)(Math.acos(a) * PI_RECIPROCAL); + return (float) (Math.acos(a) * PI_RECIPROCAL); } @OpenCLMapping(mapTo = "asinpi") @@ -1933,7 +1881,7 @@ protected final double asinpi(final double a) { @OpenCLMapping(mapTo = "asinpi") protected final float asinpi(final float a) { - return (float)(Math.asin(a) * PI_RECIPROCAL); + return (float) (Math.asin(a) * PI_RECIPROCAL); } @OpenCLMapping(mapTo = "atanpi") @@ -1943,7 +1891,7 @@ protected final double atanpi(final double a) { @OpenCLMapping(mapTo = "atanpi") protected final float atanpi(final float a) { - return (float)(Math.atan(a) * PI_RECIPROCAL); + return (float) (Math.atan(a) * PI_RECIPROCAL); } @OpenCLMapping(mapTo = "atan2pi") @@ -1953,7 +1901,7 @@ protected final double atan2pi(final double y, final double x) { @OpenCLMapping(mapTo = "atan2pi") protected final float atan2pi(final float y, final double x) { - return (float)(Math.atan2(y, x) * PI_RECIPROCAL); + return (float) (Math.atan2(y, x) * PI_RECIPROCAL); } @OpenCLMapping(mapTo = "cbrt") @@ -1963,7 +1911,7 @@ protected final double cbrt(final double a) { @OpenCLMapping(mapTo = "cbrt") protected final float cbrt(final float a) { - return (float)(Math.cbrt(a)); + return (float) (Math.cbrt(a)); } @OpenCLMapping(mapTo = "cosh") @@ -1973,7 +1921,7 @@ protected final double cosh(final double x) { @OpenCLMapping(mapTo = "cosh") protected final float cosh(final float x) { - return (float)(Math.cosh(x)); + return (float) (Math.cosh(x)); } @OpenCLMapping(mapTo = "cospi") @@ -1983,7 +1931,7 @@ protected final double cospi(final double a) { @OpenCLMapping(mapTo = "cospi") protected final float cospi(final float a) { - return (float)(Math.cos(a * Math.PI)); + return (float) (Math.cos(a * Math.PI)); } @OpenCLMapping(mapTo = "exp2") @@ -1993,7 +1941,7 @@ protected final double exp2(final double a) { @OpenCLMapping(mapTo = "exp2") protected final float exp2(final float a) { - return (float)(Math.pow(2.0D, a)); + return (float) (Math.pow(2.0D, a)); } @OpenCLMapping(mapTo = "exp10") @@ -2003,7 +1951,7 @@ protected final double exp10(final double a) { @OpenCLMapping(mapTo = "exp10") protected final float exp10(final float a) { - return (float)(Math.pow(10.0D, a)); + return (float) (Math.pow(10.0D, a)); } @OpenCLMapping(mapTo = "expm1") @@ -2013,7 +1961,7 @@ protected final double expm1(final double x) { @OpenCLMapping(mapTo = "expm1") protected final float expm1(final float x) { - return (float)(Math.expm1(x)); + return (float) (Math.expm1(x)); } @OpenCLMapping(mapTo = "log2") @@ -2023,7 +1971,7 @@ protected final double log2(final double a) { @OpenCLMapping(mapTo = "log2") protected final float log2(final float a) { - return (float)(log(a) * LOG_2_RECIPROCAL); + return (float) (log(a) * LOG_2_RECIPROCAL); } @OpenCLMapping(mapTo = "log10") @@ -2033,7 +1981,7 @@ protected final double log10(final double a) { @OpenCLMapping(mapTo = "log10") protected final float log10(final float a) { - return (float)(Math.log10(a)); + return (float) (Math.log10(a)); } @OpenCLMapping(mapTo = "log1p") @@ -2043,7 +1991,7 @@ protected final double log1p(final double x) { @OpenCLMapping(mapTo = "log1p") protected final float log1p(final float x) { - return (float)(Math.log1p(x)); + return (float) (Math.log1p(x)); } @OpenCLMapping(mapTo = "mad") @@ -2058,36 +2006,34 @@ protected final float mad(final float a, final float b, final float c) { /** * Delegates to either {code}a*b+c{code} (Java) orfma(float, float, float)(OpenCL). - * - * User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. - * - * @param a value to delegate to first argument offma(float, float, float)- * @param b value to delegate to second argument offma(float, float, float)- * @param c value to delegate to third argument offma(float, float, float)- * @return a * b + c /fma(float, float, float)- * - * @seefma(float, float, float)- */ + *+ * User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. + * + * @param a value to delegate to first argument of
fma(float, float, float)+ * @param b value to delegate to second argument offma(float, float, float)+ * @param c value to delegate to third argument offma(float, float, float)+ * @return a * b + c /fma(float, float, float)+ * @seefma(float, float, float)+ */ @OpenCLMapping(mapTo = "fma") protected float fma(final float a, final float b, final float c) { - return a * b + c; + return a * b + c; } /** * Delegates to either {code}a*b+c{code} (Java) orfma(double, double, double)(OpenCL). - * - * User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. - * - * @param a value to delegate to first argument offma(double, double, double)- * @param b value to delegate to second argument offma(double, double, double)- * @param c value to delegate to third argument offma(double, double, double)- * @return a * b + c /fma(double, double, double)- * - * @seefma(double, double, double)- */ + *+ * User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. + * + * @param a value to delegate to first argument of
fma(double, double, double)+ * @param b value to delegate to second argument offma(double, double, double)+ * @param c value to delegate to third argument offma(double, double, double)+ * @return a * b + c /fma(double, double, double)+ * @seefma(double, double, double)+ */ @OpenCLMapping(mapTo = "fma") protected double fma(final double a, final double b, final double c) { - return a * b + c; + return a * b + c; } @OpenCLMapping(mapTo = "nextafter") @@ -2097,17 +2043,16 @@ protected final double nextAfter(final double start, final double direction) { @OpenCLMapping(mapTo = "nextafter") protected final float nextAfter(final float start, final float direction) { - return (float)(Math.nextAfter(start, direction)); + return Math.nextAfter(start, direction); } /** * Delegates to either {@link java.lang.Math#sinh(double)} (Java) orsinh(double)(OpenCL). - * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param x value to delegate to {@link java.lang.Math#sinh(double)}/
sinh(double)* @return {@link java.lang.Math#sinh(double)}/sinh(double)- * * @see java.lang.Math#sinh(double) * @seesinh(double)*/ @@ -2118,30 +2063,28 @@ protected final double sinh(final double x) { /** * Delegates to either {@link java.lang.Math#sinh(double)} (Java) orsinh(float)(OpenCL). - * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param x value to delegate to {@link java.lang.Math#sinh(double)}/
sinh(float)* @return {@link java.lang.Math#sinh(double)}/sinh(float)- * * @see java.lang.Math#sinh(double) * @seesinh(float)*/ @OpenCLMapping(mapTo = "sinh") protected final float sinh(final float x) { - return (float)(Math.sinh(x)); + return (float) (Math.sinh(x)); } /** * Backed by either {@link java.lang.Math#sin(double)} (Java) orsinpi(double)(OpenCL). - * + ** This method is equivelant to
Math.sin(a * Math.PI)- * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param a value to delegate to
sinpi(double)or java equivelant * @returnsinpi(double)or java equivelant - * * @see java.lang.Math#sin(double) * @seesinpi(double)*/ @@ -2152,30 +2095,28 @@ protected final double sinpi(final double a) { /** * Backed by either {@link java.lang.Math#sin(double)} (Java) orsinpi(float)(OpenCL). - * + ** This method is equivelant to
Math.sin(a * Math.PI)- * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param a value to delegate to
sinpi(float)or java equivelant * @returnsinpi(float)or java equivelant - * * @see java.lang.Math#sin(double) * @seesinpi(float)*/ @OpenCLMapping(mapTo = "sinpi") protected final float sinpi(final float a) { - return (float)(Math.sin(a * Math.PI)); + return (float) (Math.sin(a * Math.PI)); } /** * Delegates to either {@link java.lang.Math#tanh(double)} (Java) ortanh(double)(OpenCL). - * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param x value to delegate to {@link java.lang.Math#tanh(double)}/
tanh(double)* @return {@link java.lang.Math#tanh(double)}/tanh(double)- * * @see java.lang.Math#tanh(double) * @seetanh(double)*/ @@ -2186,30 +2127,28 @@ protected final double tanh(final double x) { /** * Delegates to either {@link java.lang.Math#tanh(float)} (Java) ortanh(float)(OpenCL). - * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param x value to delegate to {@link java.lang.Math#tanh(float)}/
tanh(float)* @return {@link java.lang.Math#tanh(float)}/tanh(float)- * * @see java.lang.Math#tanh(float) * @seetanh(float)*/ @OpenCLMapping(mapTo = "tanh") protected final float tanh(final float x) { - return (float)(Math.tanh(x)); + return (float) (Math.tanh(x)); } /** * Backed by either {@link java.lang.Math#tan(double)} (Java) ortanpi(double)(OpenCL). - * + ** This method is equivelant to
Math.tan(a * Math.PI)- * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param a value to delegate to
tanpi(double)or java equivelant * @returntanpi(double)or java equivelant - * * @see java.lang.Math#tan(double) * @seetanpi(double)*/ @@ -2220,1188 +2159,1213 @@ protected final double tanpi(final double a) { /** * Backed by either {@link java.lang.Math#tan(double)} (Java) ortanpi(float)(OpenCL). - * + ** This method is equivelant to
Math.tan(a * Math.PI)- * + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param a value to delegate to
tanpi(float)or java equivelant * @returntanpi(float)or java equivelant - * * @see java.lang.Math#tan(double) * @seetanpi(float)*/ @OpenCLMapping(mapTo = "tanpi") protected final float tanpi(final float a) { - return (float)(Math.tan(a * Math.PI)); + return (float) (Math.tan(a * Math.PI)); } - // the following rsqrt and native_sqrt and native_rsqrt don't exist in java Math - // but added them here for nbody testing, not sure if we want to expose them - /** - * Computes inverse square root using {@link java.lang.Math#sqrt(double)} (Java) or delegates torsqrt(double)(OpenCL). - * + // the following rsqrt and native_sqrt and native_rsqrt don't exist in java Math + // but added them here for nbody testing, not sure if we want to expose them + + /** + * Computes inverse square root using {@link java.lang.Math#sqrt(double)} (Java) or delegates torsqrt(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _f value to delegate to {@link java.lang.Math#sqrt(double)}/
rsqrt(double)* @return( 1.0f / {@link java.lang.Math#sqrt(double)} casted to float )/rsqrt(double)- * * @see java.lang.Math#sqrt(double) * @seersqrt(double)*/ - @OpenCLMapping(mapTo = "rsqrt") - protected float rsqrt(float _f) { - return (1.0f / (float) Math.sqrt(_f)); - } + @OpenCLMapping(mapTo = "rsqrt") + protected float rsqrt(float _f) { + return (1.0f / (float) Math.sqrt(_f)); + } - /** - * Computes inverse square root using {@link java.lang.Math#sqrt(double)} (Java) or delegates torsqrt(double)(OpenCL). - * + /** + * Computes inverse square root using {@link java.lang.Math#sqrt(double)} (Java) or delegates torsqrt(double)(OpenCL). + ** User should note the differences in precision between Java and OpenCL's implementation of arithmetic functions to determine whether the difference in precision is acceptable. * * @param _d value to delegate to {@link java.lang.Math#sqrt(double)}/
rsqrt(double)* @return( 1.0f / {@link java.lang.Math#sqrt(double)} )/rsqrt(double)- * * @see java.lang.Math#sqrt(double) * @seersqrt(double)*/ - @OpenCLMapping(mapTo = "rsqrt") - protected double rsqrt(double _d) { - return (1.0 / Math.sqrt(_d)); - } - - @OpenCLMapping(mapTo = "native_sqrt") - private float native_sqrt(float _f) { - int j = Float.floatToIntBits(_f); - j = ((1 << 29) + (j >> 1)) - (1 << 22) - 0x4c00; - return (Float.intBitsToFloat(j)); - // could add more precision using one iteration of newton's method, use the following - } - - @OpenCLMapping(mapTo = "native_rsqrt") - private float native_rsqrt(float _f) { - int j = Float.floatToIntBits(_f); - j = 0x5f3759df - (j >> 1); - final float x = (Float.intBitsToFloat(j)); - return x; - // if want more precision via one iteration of newton's method, use the following - // float fhalf = 0.5f*_f; - // return (x *(1.5f - fhalf * x * x)); - } - - // Hacked from AtomicIntegerArray.getAndAdd(i, delta) - /** - * Atomically adds_deltavalue to_indexelement of array_arr(Java) or delegates toatomic_add(volatile int*, int)(OpenCL). - * - * - * @param _arr array for which an element value needs to be atomically incremented by_delta+ @OpenCLMapping(mapTo = "rsqrt") + protected double rsqrt(double _d) { + return (1.0 / Math.sqrt(_d)); + } + + @OpenCLMapping(mapTo = "native_sqrt") + private float native_sqrt(float _f) { + int j = Float.floatToIntBits(_f); + j = ((1 << 29) + (j >> 1)) - (1 << 22) - 0x4c00; + return (Float.intBitsToFloat(j)); + // could add more precision using one iteration of newton's method, use the following + } + + @OpenCLMapping(mapTo = "native_rsqrt") + private float native_rsqrt(float _f) { + int j = Float.floatToIntBits(_f); + j = 0x5f3759df - (j >> 1); + final float x = (Float.intBitsToFloat(j)); + return x; + // if want more precision via one iteration of newton's method, use the following + // float fhalf = 0.5f*_f; + // return (x *(1.5f - fhalf * x * x)); + } + + // Hacked from AtomicIntegerArray.getAndAdd(i, delta) + + /** + * Atomically adds_deltavalue to_indexelement of array_arr(Java) or delegates toatomic_add(volatile int*, int)(OpenCL). + * + * @param _arr array for which an element value needs to be atomically incremented by_delta* @param _index index of the_arrarray that needs to be atomically incremented by_delta* @param _delta value by which_indexelement of_arrarray needs to be atomically incremented * @return previous value of_indexelement of_arrarray - * * @seeatomic_add(volatile int*, int)*/ - @OpenCLMapping(atomic32 = true) - protected int atomicAdd(int[] _arr, int _index, int _delta) { - if (!Config.disableUnsafe) { - return UnsafeWrapper.atomicAdd(_arr, _index, _delta); - } else { - synchronized (_arr) { - final int previous = _arr[_index]; - _arr[_index] += _delta; - return previous; - } - } - } - - /** - * Wait for all kernels in the current group to rendezvous at this call before continuing execution. - * - * @annotion Experimental - */ - @OpenCLDelegate - @Experimental - protected final void localBarrier() { - kernelState.awaitOnLocalBarrier(); - } - - /** - * Wait for all kernels in the current group to rendezvous at this call before continuing execution. - * - * - * Java version is identical to localBarrier() - * - * @annotion Experimental - * @deprecated - */ - @OpenCLDelegate - @Experimental - @Deprecated - protected final void globalBarrier() throws DeprecatedException { - throw new DeprecatedException( + @OpenCLMapping(atomic32 = true) + protected int atomicAdd(int[] _arr, int _index, int _delta) { + if (!Config.disableUnsafe) { + return UnsafeWrapper.atomicAdd(_arr, _index, _delta); + } else { + //synchronized (_arr) { + final int previous = _arr[_index]; + _arr[_index] += _delta; + return previous; + //} + } + } + + /** + * Wait for all kernels in the current group to rendezvous at this call before continuing execution. + * + * @annotion Experimental + */ + @OpenCLDelegate + @Experimental + protected final void localBarrier() { + kernelState.awaitOnLocalBarrier(); + } + + /** + * Wait for all kernels in the current group to rendezvous at this call before continuing execution. + *+ *
+ * Java version is identical to localBarrier() + * + * @annotion Experimental + * @deprecated + */ + @OpenCLDelegate + @Experimental + @Deprecated + protected final void globalBarrier() throws DeprecatedException { + throw new DeprecatedException( "Kernel.globalBarrier() has been deprecated. It was based an incorrect understanding of OpenCL functionality."); - } + } + + @OpenCLMapping(mapTo = "hypot") + protected float hypot(final float a, final float b) { + return (float) Math.hypot(a, b); + } - @OpenCLMapping(mapTo = "hypot") - protected float hypot(final float a, final float b) { - return (float) Math.hypot(a, b); - } + @OpenCLMapping(mapTo = "hypot") + protected double hypot(final double a, final double b) { + return Math.hypot(a, b); + } - @OpenCLMapping(mapTo = "hypot") - protected double hypot(final double a, final double b) { - return Math.hypot(a, b); - } + public KernelState getKernelState() { + return kernelState; + } - public KernelState getKernelState() { - return kernelState; - } - private KernelRunner prepareKernelRunner() { - if (kernelRunner == null) { - kernelRunner = new KernelRunner(this); - } - return kernelRunner; - } - - /** - * Determine the execution time of the previous Kernel.execute(range) call. - * - * Note that for the first call this will include the conversion time. - * - * @return The time spent executing the kernel (ms) - * - * @see #getConversionTime(); - * @see #getAccumulatedExecutionTime(); - * - */ - public double getExecutionTime() { - KernelProfile profile = KernelManager.instance().getProfile(getClass()); - synchronized (profile) { - return profile.getLastExecutionTime(); - } - } - - /** - * Determine the total execution time of all previous Kernel.execute(range) calls. - * - * Note that this will include the initial conversion time. - * - * @return The total time spent executing the kernel (ms) - * - * @see #getExecutionTime(); - * @see #getConversionTime(); - * - */ - public double getAccumulatedExecutionTime() { - KernelProfile profile = KernelManager.instance().getProfile(getClass()); - synchronized (profile) { - return profile.getAccumulatedTotalTime(); - } - } - - /** - * Determine the time taken to convert bytecode to OpenCL for first Kernel.execute(range) call. - * @return The time spent preparing the kernel for execution using GPU - * - * @see #getExecutionTime(); - * @see #getAccumulatedExecutionTime(); - */ - public double getConversionTime() { - KernelProfile profile = KernelManager.instance().getProfile(getClass()); - synchronized (profile) { - return profile.getLastConversionTime(); - } - } - - /** - * Start execution of
_rangekernels. - *- * When
kernel.execute(globalSize)is invoked, Aparapi will schedule the execution ofglobalSizekernels. If the execution mode is GPU then - * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. - *- * @param _range The number of Kernels that we would like to initiate. - * @returnThe Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) - * - */ - public synchronized Kernel execute(Range _range) { - return (execute(_range, 1)); - } - - @Override - @SuppressWarnings("deprecation") - public String toString() { - if (executionMode == EXECUTION_MODE.AUTO) { - List
preferredDevices = KernelManager.instance().getPreferences(this).getPreferredDevices(this); - StringBuilder preferredDevicesSummary = new StringBuilder("{"); - for (int i = 0; i < preferredDevices.size(); ++i) { - Device device = preferredDevices.get(i); - preferredDevicesSummary.append(device.getShortDescription()); - if (i < preferredDevices.size() - 1) { - preferredDevicesSummary.append("|"); - } - } - preferredDevicesSummary.append("}"); - return Reflection.getSimpleName(getClass()) + ", devices=" + preferredDevicesSummary.toString(); - } else { - return Reflection.getSimpleName(getClass()) + ", modes=" + executionModes + ", current = " + executionMode; - } - } - - /** - * Start execution of _rangekernels. - *- * When
kernel.execute(_range)is 1invoked, Aparapi will schedule the execution of_rangekernels. If the execution mode is GPU then - * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. - *- * Since adding the new
Range classthis method offers backward compatibility and merely defers toreturn (execute(Range.create(_range), 1));. - * @param _range The number of Kernels that we would like to initiate. - * @returnThe Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) - * - */ - public synchronized Kernel execute(int _range) { - return (execute(createRange(_range), 1)); - } - - @SuppressWarnings("deprecation") - protected Range createRange(int _range) { - if (executionMode.equals(EXECUTION_MODE.AUTO)) { - Device device = getTargetDevice(); - Range range = Range.create(device, _range); - return range; - } else { - return Range.create(null, _range); - } - } - - /** - * Start execution of_passesiterations of_rangekernels. - *- * When
kernel.execute(_range, _passes)is invoked, Aparapi will schedule the execution of_reangekernels. If the execution mode is GPU then - * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. - *- * @param _passes The number of passes to make - * @return The Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) - * - */ - public synchronized Kernel execute(Range _range, int _passes) { - return (execute("run", _range, _passes)); - } - - /** - * Start execution of
_passesiterations over the_rangeof kernels. - *- * When
kernel.execute(_range)is invoked, Aparapi will schedule the execution of_rangekernels. If the execution mode is GPU then - * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. - *- * Since adding the new
Range classthis method offers backward compatibility and merely defers toreturn (execute(Range.create(_range), 1));. - * @param _range The number of Kernels that we would like to initiate. - * @returnThe Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) - * - */ - public synchronized Kernel execute(int _range, int _passes) { - return (execute(createRange(_range), _passes)); - } - - /** - * Start execution ofglobalSizekernels for the given entrypoint. - *- * When
kernel.execute("entrypoint", globalSize)is invoked, Aparapi will schedule the execution ofglobalSizekernels. If the execution mode is GPU then - * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. - *- * @param _entrypoint is the name of the method we wish to use as the entrypoint to the kernel - * @return The Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) - * - */ - public synchronized Kernel execute(String _entrypoint, Range _range) { - return (execute(_entrypoint, _range, 1)); - } - - /** - * Start execution of
globalSizekernels for the given entrypoint. - *- * When
kernel.execute("entrypoint", globalSize)is invoked, Aparapi will schedule the execution ofglobalSizekernels. If the execution mode is GPU then - * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. - *- * @param _entrypoint is the name of the method we wish to use as the entrypoint to the kernel - * @return The Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) - * - */ - public synchronized Kernel execute(String _entrypoint, Range _range, int _passes) { - return prepareKernelRunner().execute(_entrypoint, _range, _passes); - } - - public boolean isAutoCleanUpArrays() { - return autoCleanUpArrays; - } - - /** - * Property which if true enables automatic calling of {@link #cleanUpArrays()} following each execution. - */ - public void setAutoCleanUpArrays(boolean autoCleanUpArrays) { - this.autoCleanUpArrays = autoCleanUpArrays; - } - - /** - * Frees the bulk of the resources used by this kernel, by setting array sizes in non-primitive {@link KernelArg}s to 1 (0 size is prohibited) and invoking kernel - * execution on a zero size range. Unlike {@link #dispose()}, this does not prohibit further invocations of this kernel, as sundry resources such as OpenCL queues are - * not freed by this method. - * - *
This allows a "dormant" Kernel to remain in existence without undue strain on GPU resources, which may be strongly preferable to disposing a Kernel and - * recreating another one later, as creation/use of a new Kernel (specifically creation of its associated OpenCL context) is expensive.
- * - *Note that where the underlying array field is declared final, for obvious reasons it is not resized to zero.
- */ - public synchronized void cleanUpArrays() { - if (kernelRunner != null) { - kernelRunner.cleanUpArrays(); - } - } - - /** - * Release any resources associated with this Kernel. - *- * When the execution mode is
CPUorGPU, Aparapi stores some OpenCL resources in a data structure associated with the kernel instance. The - *dispose()method must be called to release these resources. - *- * If
execute(int _globalSize)is called afterdispose()is called the results are undefined. - */ - public synchronized void dispose() { - if (kernelRunner != null) { - kernelRunner.dispose(); - kernelRunner = null; - } - } - - public boolean isRunningCL() { - return getTargetDevice() instanceof OpenCLDevice; - } - - public final Device getTargetDevice() { - return KernelManager.instance().getPreferences(this).getPreferredDevice(this); - } - - /** @return true by default, may be overriden to allow vetoing of a device or devices by a given Kernel instance. */ - public boolean isAllowDevice(Device _device) { - return true; - } - - /** - * @deprecated See {@link EXECUTION_MODE} - *- * Return the current execution mode. - * - * Before a Kernel executes, this return value will be the execution mode as determined by the setting of - * the EXECUTION_MODE enumeration. By default, this setting is either GPU - * if OpenCL is available on the target system, or JTP otherwise. This default setting can be - * changed by calling setExecutionMode(). - * - *
- * After a Kernel executes, the return value will be the mode in which the Kernel actually executed. - * - * @return The current execution mode. - * - * @see #setExecutionMode(EXECUTION_MODE) - */ - @Deprecated - public EXECUTION_MODE getExecutionMode() { - return (executionMode); - } - - /** - * @deprecated See {@link EXECUTION_MODE} - *
- * Set the execution mode. - *
- * This should be regarded as a request. The real mode will be determined at runtime based on the availability of OpenCL and the characteristics of the workload. - * - * @param _executionMode the requested execution mode. - * - * @see #getExecutionMode() - */ - @Deprecated - public void setExecutionMode(EXECUTION_MODE _executionMode) { - executionMode = _executionMode; - } - - public void setExecutionModeWithoutFallback(EXECUTION_MODE _executionMode) { - executionModes.clear(); - executionModes.add(_executionMode); - currentMode = executionModes.iterator(); - executionMode = currentMode.next(); } - - /** - * @deprecated See {@link EXECUTION_MODE} - */ - @Deprecated - public void setFallbackExecutionMode() { - executionMode = EXECUTION_MODE.getFallbackExecutionMode(); - } - - final static Map
typeToLetterMap = new HashMap (); - - static { - // only primitive types for now - typeToLetterMap.put("double", "D"); - typeToLetterMap.put("float", "F"); - typeToLetterMap.put("int", "I"); - typeToLetterMap.put("long", "J"); - typeToLetterMap.put("boolean", "Z"); - typeToLetterMap.put("byte", "B"); - typeToLetterMap.put("char", "C"); - typeToLetterMap.put("short", "S"); - typeToLetterMap.put("void", "V"); - } - - private static String descriptorToReturnTypeLetter(String desc) { - // find the letter after the closed parenthesis - return desc.substring(desc.lastIndexOf(')') + 1); - } - - private static String getReturnTypeLetter(Method meth) { - return toClassShortNameIfAny(meth.getReturnType()); - } - - private static String toClassShortNameIfAny(final Class> retClass) { - if (retClass.isArray()) { - return "[" + toClassShortNameIfAny(retClass.getComponentType()); - } - final String strRetClass = retClass.toString(); - final String mapping = typeToLetterMap.get(strRetClass); - // System.out.println("strRetClass = <" + strRetClass + ">, mapping = " + mapping); - if (mapping == null) - return "[" + retClass.getName() + ";"; - return mapping; - } - - public static String getMappedMethodName(MethodReferenceEntry _methodReferenceEntry) { - if (CacheEnabler.areCachesEnabled()) - return getProperty(mappedMethodNamesCache, _methodReferenceEntry, null); - String mappedName = null; - final String name = _methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8(); - Class> currentClass = _methodReferenceEntry.getOwnerClassModel().getClassWeAreModelling(); - while (currentClass != Object.class) { - for (final Method kernelMethod : currentClass.getDeclaredMethods()) { - if (kernelMethod.isAnnotationPresent(OpenCLMapping.class)) { - // ultimately, need a way to constrain this based upon signature (to disambiguate abs(float) from abs(int); - // for Alpha, we will just disambiguate based on the return type - if (false) { - System.out.println("kernelMethod is ... " + kernelMethod.toGenericString()); - System.out.println("returnType = " + kernelMethod.getReturnType()); - System.out.println("returnTypeLetter = " + getReturnTypeLetter(kernelMethod)); - System.out.println("kernelMethod getName = " + kernelMethod.getName()); - System.out.println("methRefName = " + name + " descriptor = " - + _methodReferenceEntry.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8()); - System.out.println("descToReturnTypeLetter = " - + descriptorToReturnTypeLetter(_methodReferenceEntry.getNameAndTypeEntry().getDescriptorUTF8Entry() - .getUTF8())); - } - if (toSignature(_methodReferenceEntry).equals(toSignature(kernelMethod))) { - final OpenCLMapping annotation = kernelMethod.getAnnotation(OpenCLMapping.class); - final String mapTo = annotation.mapTo(); - if (!mapTo.equals("")) { - mappedName = mapTo; - // System.out.println("mapTo = " + mapTo); - } - } - } - } - if (mappedName != null) - break; - currentClass = currentClass.getSuperclass(); - } - // System.out.println("... in getMappedMethodName, returning = " + mappedName); - return (mappedName); - } - - public static boolean isMappedMethod(MethodReferenceEntry methodReferenceEntry) { - if (CacheEnabler.areCachesEnabled()) - return getBoolean(mappedMethodFlags, methodReferenceEntry); - boolean isMapped = false; - for (final Method kernelMethod : Kernel.class.getDeclaredMethods()) { - if (kernelMethod.isAnnotationPresent(OpenCLMapping.class)) { - if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { - - // well they have the same name ;) - isMapped = true; - } - } - } - return (isMapped); - } - - public static boolean isOpenCLDelegateMethod(MethodReferenceEntry methodReferenceEntry) { - if (CacheEnabler.areCachesEnabled()) - return getBoolean(openCLDelegateMethodFlags, methodReferenceEntry); - boolean isMapped = false; - for (final Method kernelMethod : Kernel.class.getDeclaredMethods()) { - if (kernelMethod.isAnnotationPresent(OpenCLDelegate.class)) { - if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { - - // well they have the same name ;) - isMapped = true; - } - } - } - return (isMapped); - } - - public static boolean usesAtomic32(MethodReferenceEntry methodReferenceEntry) { - if (CacheEnabler.areCachesEnabled()) - return getProperty(atomic32Cache, methodReferenceEntry, false); - for (final Method kernelMethod : Kernel.class.getDeclaredMethods()) { - if (kernelMethod.isAnnotationPresent(OpenCLMapping.class)) { - if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { - final OpenCLMapping annotation = kernelMethod.getAnnotation(OpenCLMapping.class); - return annotation.atomic32(); - } - } - } - return (false); - } - - // For alpha release atomic64 is not supported - public static boolean usesAtomic64(MethodReferenceEntry methodReferenceEntry) { - // if (CacheEnabler.areCachesEnabled()) - // return getProperty(atomic64Cache, methodReferenceEntry, false); - //for (java.lang.reflect.Method kernelMethod : Kernel.class.getDeclaredMethods()) { - // if (kernelMethod.isAnnotationPresent(Kernel.OpenCLMapping.class)) { - // if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { - // OpenCLMapping annotation = kernelMethod.getAnnotation(Kernel.OpenCLMapping.class); - // return annotation.atomic64(); - // } - // } - //} - return (false); - } - - // the flag useNullForLocalSize is useful for testing that what we compute for localSize is what OpenCL - // would also compute if we passed in null. In non-testing mode, we just call execute with the - // same localSize that we computed in getLocalSizeJNI. We don't want do publicize these of course. - // GRF we can't access this from test classes without exposing in in javadoc so I left the flag but made the test/set of the flag reflectively - boolean useNullForLocalSize = false; - - // Explicit memory management API's follow - - /** - * For dev purposes (we should remove this for production) allow us to define that this Kernel uses explicit memory management - * @param _explicit (true if we want explicit memory management) - */ - public void setExplicit(boolean _explicit) { - prepareKernelRunner().setExplicit(_explicit); - } - - /** - * For dev purposes (we should remove this for production) determine whether this Kernel uses explicit memory management - * @return (true if we kernel is using explicit memory management) - */ - public boolean isExplicit() { - return prepareKernelRunner().isExplicit(); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(long[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(long[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(long[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(double[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(double[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(double[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(float[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(float[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(float[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(int[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(int[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(int[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(byte[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(byte[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(byte[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API + + //TODO use WeakConcurrent HashMap + private final Map runners = Collections.synchronizedMap( + new WeakHashMap (Runtime.getRuntime().availableProcessors() * 2, 0.95f) + ); + + + private KernelRunner runner() { + return runners.computeIfAbsent(Thread.currentThread(), (tt)-> new KernelRunner(this)); + } + + /** + * Determine the execution time of the previous Kernel.execute(range) call. + * + * Note that for the first call this will include the conversion time. + * + * @return The time spent executing the kernel (ms) + * @see #getConversionTime(); + * @see #getAccumulatedExecutionTime(); */ - public Kernel put(char[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(char[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(char[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(boolean[] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(boolean[][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Tag this array so that it is explicitly enqueued before the kernel is executed - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel put(boolean[][][] array) { - prepareKernelRunner().put(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(long[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(long[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(long[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(double[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(double[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(double[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(float[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(float[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(float[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(int[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(int[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(int[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(byte[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(byte[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(byte[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(char[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(char[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(char[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(boolean[] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(boolean[][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. - * @param array - * @return This kernel so that we can use the 'fluent' style API - */ - public Kernel get(boolean[][][] array) { - prepareKernelRunner().get(array); - return (this); - } - - /** - * Get the profiling information from the last successful call to Kernel.execute(). - * @return A list of ProfileInfo records - */ - public List
getProfileInfo() { - return prepareKernelRunner().getProfileInfo(); - } - - /** - * @deprecated See {@link EXECUTION_MODE}. - */ - @Deprecated - private final LinkedHashSet executionModes = (Config.executionMode != null) ? EXECUTION_MODE.getDefaultExecutionModes() : new LinkedHashSet<>(Collections.singleton(EXECUTION_MODE.AUTO)); - - /** - * @deprecated See {@link EXECUTION_MODE}. - */ - @Deprecated - private Iterator currentMode = executionModes.iterator(); - - /** - * @deprecated See {@link EXECUTION_MODE}. - */ - @Deprecated - private EXECUTION_MODE executionMode = currentMode.next(); - - /** - * @deprecated See {@link EXECUTION_MODE}. - * - * set possible fallback path for execution modes. - * for example setExecutionFallbackPath(GPU,CPU,JTP) will try to use the GPU - * if it fails it will fall back to OpenCL CPU and finally it will try JTP. - */ - @Deprecated - public void addExecutionModes(EXECUTION_MODE... platforms) { - executionModes.addAll(Arrays.asList(platforms)); - currentMode = executionModes.iterator(); - executionMode = currentMode.next(); - } - - /** - * @deprecated See {@link EXECUTION_MODE}. - * @return is there another execution path we can try - */ - @Deprecated - public boolean hasNextExecutionMode() { - return currentMode.hasNext(); - } - - /** - * @deprecated See {@link EXECUTION_MODE}. - * try the next execution path in the list if there aren't any more than give up - */ - @Deprecated - public void tryNextExecutionMode() { - if (currentMode.hasNext()) { - executionMode = currentMode.next(); - } - } + public double getExecutionTime() { + KernelProfile profile = KernelManager.instance().getProfile(getClass()); - private static final ValueCache
, Map , RuntimeException> mappedMethodFlags = markedWith(OpenCLMapping.class); + return profile.getTotalExecutionTime(); - private static final ValueCache , Map , RuntimeException> openCLDelegateMethodFlags = markedWith(OpenCLDelegate.class); + } - private static final ValueCache , Map , RuntimeException> atomic32Cache = cacheProperty(new ValueComputer , Map >() { - @Override - public Map compute(Class> key) { - Map properties = new HashMap<>(); - for (final Method method : key.getDeclaredMethods()) { - if (isRelevant(method) && method.isAnnotationPresent(OpenCLMapping.class)) { - properties.put(toSignature(method), method.getAnnotation(OpenCLMapping.class).atomic32()); - } - } - return properties; - } - }); + /** + * Determine the total execution time of all previous Kernel.execute(range) calls. + * + * Note that this will include the initial conversion time. + * + * @return The total time spent executing the kernel (ms) + * @see #getExecutionTime(); + * @see #getConversionTime(); + */ + public double getAccumulatedExecutionTime() { + KernelProfile profile = KernelManager.instance().getProfile(getClass()); - private static final ValueCache
, Map , RuntimeException> atomic64Cache = cacheProperty(new ValueComputer , Map >() { - @Override - public Map compute(Class> key) { - Map properties = new HashMap<>(); - for (final Method method : key.getDeclaredMethods()) { - if (isRelevant(method) && method.isAnnotationPresent(OpenCLMapping.class)) { - properties.put(toSignature(method), method.getAnnotation(OpenCLMapping.class).atomic64()); - } - } - return properties; - } - }); - - private static boolean getBoolean(ValueCache , Map , RuntimeException> methodNamesCache, - MethodReferenceEntry methodReferenceEntry) { - return getProperty(methodNamesCache, methodReferenceEntry, false); - } - - private static ValueCache , Map , RuntimeException> markedWith( - final Class annotationClass) { - return cacheProperty(new ValueComputer , Map >() { - @Override - public Map compute(Class> key) { - Map markedMethodNames = new HashMap<>(); - for (final Method method : key.getDeclaredMethods()) { - markedMethodNames.put(toSignature(method), method.isAnnotationPresent(annotationClass)); - } - return markedMethodNames; - } - }); - } - - static String toSignature(Method method) { - return method.getName() + getArgumentsLetters(method) + getReturnTypeLetter(method); - } - - private static String getArgumentsLetters(Method method) { - StringBuilder sb = new StringBuilder("("); - for (Class> parameterClass : method.getParameterTypes()) { - sb.append(toClassShortNameIfAny(parameterClass)); - } - sb.append(")"); - return sb.toString(); - } - - private static boolean isRelevant(Method method) { - return !method.isSynthetic() && !method.isBridge(); - } - - private static V getProperty(ValueCache , Map , T> cache, - MethodReferenceEntry methodReferenceEntry, V defaultValue) throws T { - Map map = cache.computeIfAbsent(methodReferenceEntry.getOwnerClassModel().getClassWeAreModelling()); - String key = toSignature(methodReferenceEntry); - if (map.containsKey(key)) - return map.get(key); - return defaultValue; - } - - private static String toSignature(MethodReferenceEntry methodReferenceEntry) { - NameAndTypeEntry nameAndTypeEntry = methodReferenceEntry.getNameAndTypeEntry(); - return nameAndTypeEntry.getNameUTF8Entry().getUTF8() + nameAndTypeEntry.getDescriptorUTF8Entry().getUTF8(); - } - - private static final ValueCache , Map , RuntimeException> mappedMethodNamesCache = cacheProperty(new ValueComputer , Map >() { - @Override - public Map compute(Class> key) { - Map properties = new HashMap<>(); - for (final Method method : key.getDeclaredMethods()) { - if (isRelevant(method) && method.isAnnotationPresent(OpenCLMapping.class)) { - // ultimately, need a way to constrain this based upon signature (to disambiguate abs(float) from abs(int); - final OpenCLMapping annotation = method.getAnnotation(OpenCLMapping.class); - final String mapTo = annotation.mapTo(); - if (mapTo != null && !mapTo.equals("")) { - properties.put(toSignature(method), mapTo); - } - } - } - return properties; - } - }); + return profile.getTotalTime(); - private static ValueCache , Map , T> cacheProperty( - final ThrowingValueComputer , Map , T> throwingValueComputer) { - return ValueCache.on(new ThrowingValueComputer , Map , T>() { - @Override - public Map compute(Class> key) throws T { - Map properties = new HashMap<>(); - Deque > superclasses = new ArrayDeque<>(); - Class> currentSuperClass = key; - do { - superclasses.push(currentSuperClass); - currentSuperClass = currentSuperClass.getSuperclass(); - } while (currentSuperClass != Object.class); - for (Class> clazz : superclasses) { - // Overwrite property values for shadowed/overriden methods - properties.putAll(throwingValueComputer.compute(clazz)); + } + + /** + * Determine the time taken to convert bytecode to OpenCL for first Kernel.execute(range) call. + * + * @return The time spent preparing the kernel for execution using GPU + * @see #getExecutionTime(); + * @see #getAccumulatedExecutionTime(); + */ + public double getTotalConversionTime() { + KernelProfile profile = KernelManager.instance().getProfile(getClass()); + + return profile.getTotalConversionTime(); + + } + + /** + * Start execution of _rangekernels. + *+ * When
kernel.execute(globalSize)is invoked, Aparapi will schedule the execution ofglobalSizekernels. If the execution mode is GPU then + * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. + *+ * + * @param _range The number of Kernels that we would like to initiate. + * @returnThe Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) + */ + public Kernel execute(Range _range) { + return execute(_range, 1); + } + + @Override + @SuppressWarnings("deprecation") + public String toString() { + if (executionMode == EXECUTION_MODE.AUTO) { + List
preferredDevices = KernelManager.instance().getPreferences(this).getPreferredDevices(this); + StringBuilder preferredDevicesSummary = new StringBuilder("{"); + for (int i = 0; i < preferredDevices.size(); ++i) { + Device device = preferredDevices.get(i); + preferredDevicesSummary.append(device.getShortDescription()); + if (i < preferredDevices.size() - 1) { + preferredDevicesSummary.append('|'); + } } - return properties; - } - }); - } - - public static void invalidateCaches() { - atomic32Cache.invalidate(); - atomic64Cache.invalidate(); - mappedMethodFlags.invalidate(); - mappedMethodNamesCache.invalidate(); - openCLDelegateMethodFlags.invalidate(); - } + preferredDevicesSummary.append('}'); + return Reflection.getSimpleName(getClass()) + ", devices=" + preferredDevicesSummary; + } else { + return Reflection.getSimpleName(getClass()) + ", modes=" + executionModes + ", current = " + executionMode; + } + } + + /** + * Start execution of _rangekernels. + *+ * When
kernel.execute(_range)is 1invoked, Aparapi will schedule the execution of_rangekernels. If the execution mode is GPU then + * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. + *+ * Since adding the new
Range classthis method offers backward compatibility and merely defers toreturn (execute(Range.create(_range), 1));. + * + * @param _range The number of Kernels that we would like to initiate. + * @returnThe Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) + */ + public Kernel execute(int _range) { + return execute(createRange(_range), 1); + } + + @SuppressWarnings("deprecation") + private Range createRange(int _range) { + if (executionMode.equals(EXECUTION_MODE.AUTO)) { + Device device = getTargetDevice(); + Range range = Range.create(device, _range); + return range; + } else { + return Range.create(null, _range); + } + } + + /** + * Start execution of_passesiterations of_rangekernels. + *+ * When
kernel.execute(_range, _passes)is invoked, Aparapi will schedule the execution of_reangekernels. If the execution mode is GPU then + * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. + *+ * + * @param _passes The number of passes to make + * @return The Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) + */ + public Kernel execute(Range _range, int _passes) { + return execute("run", _range, _passes); + } + + /** + * Start execution of
_passesiterations over the_rangeof kernels. + *+ * When
kernel.execute(_range)is invoked, Aparapi will schedule the execution of_rangekernels. If the execution mode is GPU then + * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. + *+ * Since adding the new
Range classthis method offers backward compatibility and merely defers toreturn (execute(Range.create(_range), 1));. + * + * @param _range The number of Kernels that we would like to initiate. + * @returnThe Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) + */ + public Kernel execute(int _range, int _passes) { + return execute(createRange(_range), _passes); + } + + /** + * Start execution ofglobalSizekernels for the given entrypoint. + *+ * When
kernel.execute("entrypoint", globalSize)is invoked, Aparapi will schedule the execution ofglobalSizekernels. If the execution mode is GPU then + * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. + *+ * + * @param _entrypoint is the name of the method we wish to use as the entrypoint to the kernel + * @return The Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) + */ + public Kernel execute(String _entrypoint, Range _range) { + return execute(_entrypoint, _range, 1); + } + + /** + * Start execution of
globalSizekernels for the given entrypoint. + *+ * When
kernel.execute("entrypoint", globalSize)is invoked, Aparapi will schedule the execution ofglobalSizekernels. If the execution mode is GPU then + * the kernels will execute as OpenCL code on the GPU device. Otherwise, if the mode is JTP, the kernels will execute as a pool of Java threads on the CPU. + *+ * + * @param _entrypoint is the name of the method we wish to use as the entrypoint to the kernel + * @return The Kernel instance (this) so we can chain calls to put(arr).execute(range).get(arr) + */ + private Kernel execute(String _entrypoint, Range _range, int _passes) { + return runner().execute(_entrypoint, _range, _passes); + } + + public boolean isAutoCleanUpArrays() { + return autoCleanUpArrays; + } + + /** + * Property which if true enables automatic calling of {@link #cleanUpArrays()} following each execution. + */ + public void setAutoCleanUpArrays(boolean autoCleanUpArrays) { + this.autoCleanUpArrays = autoCleanUpArrays; + } + + /** + * Frees the bulk of the resources used by this kernel, by setting array sizes in non-primitive {@link KernelArg}s to 1 (0 size is prohibited) and invoking kernel + * execution on a zero size range. Unlike {@link #dispose()}, this does not prohibit further invocations of this kernel, as sundry resources such as OpenCL queues are + * not freed by this method. + *
+ *
This allows a "dormant" Kernel to remain in existence without undue strain on GPU resources, which may be strongly preferable to disposing a Kernel and + * recreating another one later, as creation/use of a new Kernel (specifically creation of its associated OpenCL context) is expensive.
+ *+ *
Note that where the underlying array field is declared final, for obvious reasons it is not resized to zero.
+ */ + public void cleanUpArrays() { + runners.values().forEach(KernelRunner::cleanUpArrays); + } + + /** + * Release any resources associated with this Kernel. + *+ * When the execution mode is
CPUorGPU, Aparapi stores some OpenCL resources in a data structure associated with the kernel instance. The + *dispose()method must be called to release these resources. + *+ * If
execute(int _globalSize)is called afterdispose()is called the results are undefined. + */ + public void dispose() { + runners.entrySet().removeIf(e -> { + e.getValue().dispose(); + return true; + }); + } + + public boolean isRunningCL() { + return getTargetDevice() instanceof OpenCLDevice; + } + + public final Device getTargetDevice() { + return KernelManager.instance().getPreferences(this).getPreferredDevice(this); + } + + /** + * @return true by default, may be overriden to allow vetoing of a device or devices by a given Kernel instance. + */ + public boolean isAllowDevice(Device _device) { + return true; + } + + /** + * @return The current execution mode. + * @see #setExecutionMode(EXECUTION_MODE) + * @deprecated See {@link EXECUTION_MODE} + *+ * Return the current execution mode. + *
+ * Before a Kernel executes, this return value will be the execution mode as determined by the setting of + * the EXECUTION_MODE enumeration. By default, this setting is either GPU + * if OpenCL is available on the target system, or JTP otherwise. This default setting can be + * changed by calling setExecutionMode(). + *
+ *
+ * After a Kernel executes, the return value will be the mode in which the Kernel actually executed. + */ + @Deprecated + public EXECUTION_MODE getExecutionMode() { + return (executionMode); + } + + /** + * @param _executionMode the requested execution mode. + * @see #getExecutionMode() + * @deprecated See {@link EXECUTION_MODE} + *
+ * Set the execution mode. + *
+ * This should be regarded as a request. The real mode will be determined at runtime based on the availability of OpenCL and the characteristics of the workload. + */ + @Deprecated + public void setExecutionMode(EXECUTION_MODE _executionMode) { + executionMode = _executionMode; + } + + public void setExecutionModeWithoutFallback(EXECUTION_MODE _executionMode) { + executionModes.clear(); + executionModes.add(_executionMode); + currentMode = executionModes.iterator(); + executionMode = currentMode.next(); + } + + /** + * @deprecated See {@link EXECUTION_MODE} + */ + @Deprecated + public void setFallbackExecutionMode() { + executionMode = EXECUTION_MODE.getFallbackExecutionMode(); + } + + private final static Map
typeToLetterMap = new HashMap<>(); + + static { + // only primitive types for now + typeToLetterMap.put("double", "D"); + typeToLetterMap.put("float", "F"); + typeToLetterMap.put("int", "I"); + typeToLetterMap.put("long", "J"); + typeToLetterMap.put("boolean", "Z"); + typeToLetterMap.put("byte", "B"); + typeToLetterMap.put("char", "C"); + typeToLetterMap.put("short", "S"); + typeToLetterMap.put("void", "V"); + } + + private static String descriptorToReturnTypeLetter(String desc) { + // find the letter after the closed parenthesis + return desc.substring(desc.lastIndexOf(')') + 1); + } + + private static String getReturnTypeLetter(Method meth) { + return toClassShortNameIfAny(meth.getReturnType()); + } + + private static String toClassShortNameIfAny(final Class> retClass) { + if (retClass.isArray()) { + return '[' + toClassShortNameIfAny(retClass.getComponentType()); + } + final String strRetClass = retClass.toString(); + final String mapping = typeToLetterMap.get(strRetClass); + // System.out.println("strRetClass = <" + strRetClass + ">, mapping = " + mapping); + if (mapping == null) + return '[' + retClass.getName() + ';'; + return mapping; + } + + public static String getMappedMethodName(MethodReferenceEntry _methodReferenceEntry) { + if (CacheEnabler.areCachesEnabled()) + return getProperty(mappedMethodNamesCache, _methodReferenceEntry, null); + String mappedName = null; + final String name = _methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8(); + Class> currentClass = _methodReferenceEntry.getOwnerClassModel().clazz; + while (currentClass != Object.class) { + for (final Method kernelMethod : currentClass.getDeclaredMethods()) { + if (kernelMethod.isAnnotationPresent(OpenCLMapping.class)) { + // ultimately, need a way to constrain this based upon signature (to disambiguate abs(float) from abs(int); + // for Alpha, we will just disambiguate based on the return type + if (false) { + System.out.println("kernelMethod is ... " + kernelMethod.toGenericString()); + System.out.println("returnType = " + kernelMethod.getReturnType()); + System.out.println("returnTypeLetter = " + getReturnTypeLetter(kernelMethod)); + System.out.println("kernelMethod getName = " + kernelMethod.getName()); + System.out.println("methRefName = " + name + " descriptor = " + + _methodReferenceEntry.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8()); + System.out.println("descToReturnTypeLetter = " + + descriptorToReturnTypeLetter(_methodReferenceEntry.getNameAndTypeEntry().getDescriptorUTF8Entry() + .getUTF8())); + } + if (toSignature(_methodReferenceEntry).equals(toSignature(kernelMethod))) { + final OpenCLMapping annotation = kernelMethod.getAnnotation(OpenCLMapping.class); + final String mapTo = annotation.mapTo(); + if (!mapTo.isEmpty()) { + mappedName = mapTo; + // System.out.println("mapTo = " + mapTo); + } + } + } + } + if (mappedName != null) + break; + currentClass = currentClass.getSuperclass(); + } + // System.out.println("... in getMappedMethodName, returning = " + mappedName); + return (mappedName); + } + + public static boolean isMappedMethod(MethodReferenceEntry methodReferenceEntry) { + if (CacheEnabler.areCachesEnabled()) + return getBoolean(mappedMethodFlags, methodReferenceEntry); + boolean isMapped = false; + for (final Method kernelMethod : Kernel.class.getDeclaredMethods()) { + if (kernelMethod.isAnnotationPresent(OpenCLMapping.class)) { + if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { + + // well they have the same name ;) + isMapped = true; + } + } + } + return (isMapped); + } + + public static boolean isOpenCLDelegateMethod(MethodReferenceEntry methodReferenceEntry) { + if (CacheEnabler.areCachesEnabled()) + return getBoolean(openCLDelegateMethodFlags, methodReferenceEntry); + boolean isMapped = false; + for (final Method kernelMethod : Kernel.class.getDeclaredMethods()) { + if (kernelMethod.isAnnotationPresent(OpenCLDelegate.class)) { + if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { + + // well they have the same name ;) + isMapped = true; + } + } + } + return (isMapped); + } + + public static boolean usesAtomic32(MethodReferenceEntry methodReferenceEntry) { + if (CacheEnabler.areCachesEnabled()) + return getProperty(atomic32Cache, methodReferenceEntry, false); + for (final Method kernelMethod : Kernel.class.getDeclaredMethods()) { + if (kernelMethod.isAnnotationPresent(OpenCLMapping.class)) { + if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { + final OpenCLMapping annotation = kernelMethod.getAnnotation(OpenCLMapping.class); + return annotation.atomic32(); + } + } + } + return (false); + } + + // For alpha release atomic64 is not supported + public static boolean usesAtomic64(MethodReferenceEntry methodReferenceEntry) { + // if (CacheEnabler.areCachesEnabled()) + // return getProperty(atomic64Cache, methodReferenceEntry, false); + //for (java.lang.reflect.Method kernelMethod : Kernel.class.getDeclaredMethods()) { + // if (kernelMethod.isAnnotationPresent(Kernel.OpenCLMapping.class)) { + // if (methodReferenceEntry.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(kernelMethod.getName())) { + // OpenCLMapping annotation = kernelMethod.getAnnotation(Kernel.OpenCLMapping.class); + // return annotation.atomic64(); + // } + // } + //} + return (false); + } + + // the flag useNullForLocalSize is useful for testing that what we compute for localSize is what OpenCL + // would also compute if we passed in null. In non-testing mode, we just call execute with the + // same localSize that we computed in getLocalSizeJNI. We don't want do publicize these of course. + // GRF we can't access this from test classes without exposing in in javadoc so I left the flag but made the test/set of the flag reflectively + boolean useNullForLocalSize = false; + + // Explicit memory management API's follow + + /** + * For dev purposes (we should remove this for production) allow us to define that this Kernel uses explicit memory management + * + * @param _explicit (true if we want explicit memory management) + */ + public void setExplicit(boolean _explicit) { + runner().setExplicit(_explicit); + } + + /** + * For dev purposes (we should remove this for production) determine whether this Kernel uses explicit memory management + * + * @return (true if we kernel is using explicit memory management) + */ + public boolean isExplicit() { + return runner().isExplicit(); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(long[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(long[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(long[][][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(double[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(double[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(double[][][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(float[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(float[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(float[][][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(int[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(int[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(int[][][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(byte[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(byte[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(byte[][][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(char[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(char[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(char[][][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(boolean[] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(boolean[][] array) { + runner().put(array); + return (this); + } + + /** + * Tag this array so that it is explicitly enqueued before the kernel is executed + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel put(boolean[][][] array) { + runner().put(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(long[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(long[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(long[][][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(double[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(double[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(double[][][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(float[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(float[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(float[][][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(int[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(int[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(int[][][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(byte[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(byte[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(byte[][][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(char[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(char[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(char[][][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(boolean[] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(boolean[][] array) { + runner().get(array); + return (this); + } + + /** + * Enqueue a request to return this buffer from the GPU. This method blocks until the array is available. + * + * @param array + * @return This kernel so that we can use the 'fluent' style API + */ + public Kernel get(boolean[][][] array) { + runner().get(array); + return (this); + } + + /** + * Get the profiling information from the last successful call to Kernel.execute(). + * + * @return A list of ProfileInfo records + */ + public List getProfileInfo() { + return runner().getProfileInfo(); + } + + /** + * @deprecated See {@link EXECUTION_MODE}. + */ + @Deprecated + private final LinkedHashSet executionModes = (Config.executionMode != null) ? EXECUTION_MODE.getDefaultExecutionModes() : new LinkedHashSet<>(Collections.singleton(EXECUTION_MODE.AUTO)); + + /** + * @deprecated See {@link EXECUTION_MODE}. + */ + @Deprecated + private Iterator currentMode = executionModes.iterator(); + + /** + * @deprecated See {@link EXECUTION_MODE}. + */ + @Deprecated + private EXECUTION_MODE executionMode = currentMode.next(); + + /** + * @deprecated See {@link EXECUTION_MODE}. + * + * set possible fallback path for execution modes. + * for example setExecutionFallbackPath(GPU,CPU,JTP) will try to use the GPU + * if it fails it will fall back to OpenCL CPU and finally it will try JTP. + */ + @Deprecated + public void addExecutionModes(EXECUTION_MODE... platforms) { + executionModes.addAll(Arrays.asList(platforms)); + currentMode = executionModes.iterator(); + executionMode = currentMode.next(); + } + + /** + * @return is there another execution path we can try + * @deprecated See {@link EXECUTION_MODE}. + */ + @Deprecated + public boolean hasNextExecutionMode() { + return currentMode.hasNext(); + } + + /** + * @deprecated See {@link EXECUTION_MODE}. + * try the next execution path in the list if there aren't any more than give up + */ + @Deprecated + public void tryNextExecutionMode() { + if (currentMode.hasNext()) { + executionMode = currentMode.next(); + } + } + + private static final ValueCache
, Map , RuntimeException> mappedMethodFlags = markedWith(OpenCLMapping.class); + + private static final ValueCache , Map , RuntimeException> openCLDelegateMethodFlags = markedWith(OpenCLDelegate.class); + + private static final ValueCache , Map , RuntimeException> atomic32Cache = cacheProperty((ValueComputer , Map >) key -> { + Map properties = new HashMap<>(); + for (final Method method : key.getDeclaredMethods()) { + if (isRelevant(method) && method.isAnnotationPresent(OpenCLMapping.class)) { + properties.put(toSignature(method), method.getAnnotation(OpenCLMapping.class).atomic32()); + } + } + return properties; + }); + + private static final ValueCache , Map , RuntimeException> atomic64Cache = cacheProperty((ValueComputer , Map >) key -> { + Map properties = new HashMap<>(); + for (final Method method : key.getDeclaredMethods()) { + if (isRelevant(method) && method.isAnnotationPresent(OpenCLMapping.class)) { + properties.put(toSignature(method), method.getAnnotation(OpenCLMapping.class).atomic64()); + } + } + return properties; + }); + + private static boolean getBoolean(ValueCache , Map , RuntimeException> methodNamesCache, + MethodReferenceEntry methodReferenceEntry) { + return getProperty(methodNamesCache, methodReferenceEntry, false); + } + + private static ValueCache , Map , RuntimeException> markedWith( + final Class annotationClass) { + return cacheProperty((ValueComputer , Map >) key -> { + Map markedMethodNames = new HashMap<>(); + for (final Method method : key.getDeclaredMethods()) { + markedMethodNames.put(toSignature(method), method.isAnnotationPresent(annotationClass)); + } + return markedMethodNames; + }); + } + + private static String toSignature(Method method) { + return method.getName() + getArgumentsLetters(method) + getReturnTypeLetter(method); + } + + private static String getArgumentsLetters(Method method) { + StringBuilder sb = new StringBuilder("("); + for (Class> parameterClass : method.getParameterTypes()) { + sb.append(toClassShortNameIfAny(parameterClass)); + } + sb.append(')'); + return sb.toString(); + } + + private static boolean isRelevant(Method method) { + return !method.isSynthetic() && !method.isBridge(); + } + + private static V getProperty(ValueCache , Map , T> cache, + MethodReferenceEntry methodReferenceEntry, V defaultValue) throws T { + Map map = cache.computeIfAbsent(methodReferenceEntry.getOwnerClassModel().clazz); + String key = toSignature(methodReferenceEntry); + if (map.containsKey(key)) + return map.get(key); + return defaultValue; + } + + private static String toSignature(MethodReferenceEntry methodReferenceEntry) { + NameAndTypeEntry nameAndTypeEntry = methodReferenceEntry.getNameAndTypeEntry(); + return nameAndTypeEntry.getNameUTF8Entry().getUTF8() + nameAndTypeEntry.getDescriptorUTF8Entry().getUTF8(); + } + + private static final ValueCache , Map , RuntimeException> mappedMethodNamesCache = cacheProperty((ValueComputer , Map >) key -> { + Map properties = new HashMap<>(); + for (final Method method : key.getDeclaredMethods()) { + if (isRelevant(method) && method.isAnnotationPresent(OpenCLMapping.class)) { + // ultimately, need a way to constrain this based upon signature (to disambiguate abs(float) from abs(int); + final OpenCLMapping annotation = method.getAnnotation(OpenCLMapping.class); + final String mapTo = annotation.mapTo(); + if (mapTo != null && !mapTo.isEmpty()) { + properties.put(toSignature(method), mapTo); + } + } + } + return properties; + }); + + private static ValueCache , Map , T> cacheProperty( + final ThrowingValueComputer , Map , T> throwingValueComputer) { + return ValueCache.on(key -> { + Map properties = new HashMap<>(); + Deque > superclasses = new ArrayDeque<>(); + Class> currentSuperClass = key; + do { + superclasses.push(currentSuperClass); + currentSuperClass = currentSuperClass.getSuperclass(); + } while (currentSuperClass != Object.class); + for (Class> clazz : superclasses) { + // Overwrite property values for shadowed/overriden methods + properties.putAll(throwingValueComputer.compute(clazz)); + } + return properties; + }); + } + + public static void invalidateCaches() { + atomic32Cache.invalidate(); + atomic64Cache.invalidate(); + mappedMethodFlags.invalidate(); + mappedMethodNamesCache.invalidate(); + openCLDelegateMethodFlags.invalidate(); + } } diff --git a/src/main/java/com/aparapi/ProfileInfo.java b/src/main/java/com/aparapi/ProfileInfo.java index 47b1e5da..0391bab1 100644 --- a/src/main/java/com/aparapi/ProfileInfo.java +++ b/src/main/java/com/aparapi/ProfileInfo.java @@ -58,7 +58,7 @@ private enum TYPE { R, X, W - }; // 0 = write, 1 = execute, 2 = read + } // 0 = write, 1 = execute, 2 = read private final TYPE type; @@ -116,7 +116,7 @@ public TYPE getType() { } @Override public String toString() { - final StringBuilder sb = new StringBuilder(); + final StringBuilder sb = new StringBuilder(512); sb.append("ProfileInfo["); sb.append(type); sb.append(" '"); @@ -131,7 +131,7 @@ public TYPE getType() { sb.append(queued); sb.append(", duration="); sb.append((end - start)); - sb.append("]"); + sb.append(']'); return sb.toString(); } diff --git a/src/main/java/com/aparapi/Range.java b/src/main/java/com/aparapi/Range.java index c289d66a..5f27222e 100644 --- a/src/main/java/com/aparapi/Range.java +++ b/src/main/java/com/aparapi/Range.java @@ -64,14 +64,16 @@ */ public class Range extends RangeJNI{ - public static final int THREADS_PER_CORE = 16; + private static final int THREADS_PER_CORE = + 1; + //16; - public static final int MAX_OPENCL_GROUP_SIZE = 256; + private static final int MAX_OPENCL_GROUP_SIZE = 256; - public static final int MAX_GROUP_SIZE = Math.max(Runtime.getRuntime().availableProcessors() * THREADS_PER_CORE, + private static final int MAX_GROUP_SIZE = Math.max(Runtime.getRuntime().availableProcessors() * THREADS_PER_CORE, MAX_OPENCL_GROUP_SIZE); - private OpenCLDevice device = null; + public final Device device; private int maxWorkGroupSize; @@ -87,8 +89,8 @@ public class Range extends RangeJNI{ * @param _device * @param _dims */ - public Range(Device _device, int _dims) { - device = !(_device instanceof OpenCLDevice) ? null : (OpenCLDevice) _device; + private Range(Device _device, int _dims) { + device = _device; //!(_device instanceof OpenCLDevice) ? null : (OpenCLDevice) _device; dims = _dims; if (device != null) { @@ -99,6 +101,9 @@ public Range(Device _device, int _dims) { } } + public int volume() { + return globalSize_0 * globalSize_1 * globalSize_2; + } /** * Create a one dimensional range 0.._globalWidthwhich is processed in groups of size _localWidth. *
@@ -191,9 +196,7 @@ public static Range create(int _globalWidth, int _localWidth) { } public static Range create(int _globalWidth) { - final Range range = create(null, _globalWidth); - - return (range); + return create(null, _globalWidth); } /** @@ -429,15 +432,13 @@ public static Range create3D(int _globalWidth, int _globalHeight, int _globalDep switch (dims) { case 1: - sb.append("global:" + globalSize_0 + " local:" + (localIsDerived ? "(derived)" : "") + localSize_0); + sb.append("global:").append(globalSize_0).append(" local:").append(localIsDerived ? "(derived)" : "").append(localSize_0); break; case 2: - sb.append("2D(global:" + globalSize_0 + "x" + globalSize_1 + " local:" + (localIsDerived ? "(derived)" : "") - + localSize_0 + "x" + localSize_1 + ")"); + sb.append("2D(global:").append(globalSize_0).append('x').append(globalSize_1).append(" local:").append(localIsDerived ? "(derived)" : "").append(localSize_0).append('x').append(localSize_1).append(')'); break; case 3: - sb.append("3D(global:" + globalSize_0 + "x" + globalSize_1 + "x" + globalSize_2 + " local:" - + (localIsDerived ? "(derived)" : "") + localSize_0 + "x" + localSize_1 + "x" + localSize_2 + ")"); + sb.append("3D(global:").append(globalSize_0).append('x').append(globalSize_1).append('x').append(globalSize_2).append(" local:").append(localIsDerived ? "(derived)" : "").append(localSize_0).append('x').append(localSize_1).append('x').append(localSize_2).append(')'); break; } @@ -473,7 +474,19 @@ public int getGlobalSize(int _dim) { * @return the number of groups for the given dimension. */ public int getNumGroups(int _dim) { - return (_dim == 0 ? (globalSize_0 / localSize_0) : (_dim == 1 ? (globalSize_1 / localSize_1) : (globalSize_2 / localSize_2))); + int g, l; + switch (_dim) { + case 0: g = globalSize_0; l = localSize_0; break; + case 1: g = globalSize_1; l = localSize_1; break; + case 2: g = globalSize_2; l = localSize_2; break; + default: + throw new UnsupportedOperationException(); + } + return divideWork(g, l); + } + + static int divideWork(int total, int div) { + return (int) Math.ceil(((float)total)/div); } /** @@ -499,14 +512,14 @@ public int getGlobalSize_0() { * @param globalSize_0 * the globalSize_0 to set */ - public void setGlobalSize_0(int globalSize_0) { + private void setGlobalSize_0(int globalSize_0) { this.globalSize_0 = globalSize_0; } /** * @return the localSize_0 */ - public int getLocalSize_0() { + private int getLocalSize_0() { return localSize_0; } @@ -514,7 +527,7 @@ public int getLocalSize_0() { * @param localSize_0 * the localSize_0 to set */ - public void setLocalSize_0(int localSize_0) { + private void setLocalSize_0(int localSize_0) { this.localSize_0 = localSize_0; } @@ -529,14 +542,14 @@ public int getGlobalSize_1() { * @param globalSize_1 * the globalSize_1 to set */ - public void setGlobalSize_1(int globalSize_1) { + private void setGlobalSize_1(int globalSize_1) { this.globalSize_1 = globalSize_1; } /** * @return the localSize_1 */ - public int getLocalSize_1() { + private int getLocalSize_1() { return localSize_1; } @@ -544,7 +557,7 @@ public int getLocalSize_1() { * @param localSize_1 * the localSize_1 to set */ - public void setLocalSize_1(int localSize_1) { + private void setLocalSize_1(int localSize_1) { this.localSize_1 = localSize_1; } @@ -559,14 +572,14 @@ public int getGlobalSize_2() { * @param globalSize_2 * the globalSize_2 to set */ - public void setGlobalSize_2(int globalSize_2) { + private void setGlobalSize_2(int globalSize_2) { this.globalSize_2 = globalSize_2; } /** * @return the localSize_2 */ - public int getLocalSize_2() { + private int getLocalSize_2() { return localSize_2; } @@ -574,7 +587,7 @@ public int getLocalSize_2() { * @param localSize_2 * the localSize_2 to set */ - public void setLocalSize_2(int localSize_2) { + private void setLocalSize_2(int localSize_2) { this.localSize_2 = localSize_2; } @@ -598,7 +611,7 @@ public void setDims(int dims) { /** * @return the valid */ - public boolean isValid() { + private boolean isValid() { return valid; } @@ -606,7 +619,7 @@ public boolean isValid() { * @param valid * the valid to set */ - public void setValid(boolean valid) { + private void setValid(boolean valid) { this.valid = valid; } @@ -621,14 +634,14 @@ public boolean isLocalIsDerived() { * @param localIsDerived * the localIsDerived to set */ - public void setLocalIsDerived(boolean localIsDerived) { + private void setLocalIsDerived(boolean localIsDerived) { this.localIsDerived = localIsDerived; } /** * @return the maxWorkGroupSize */ - public int getMaxWorkGroupSize() { + private int getMaxWorkGroupSize() { return maxWorkGroupSize; } @@ -643,7 +656,7 @@ public void setMaxWorkGroupSize(int maxWorkGroupSize) { /** * @return the maxWorkItemSize */ - public int[] getMaxWorkItemSize() { + private int[] getMaxWorkItemSize() { return maxWorkItemSize; } diff --git a/src/main/java/com/aparapi/device/Device.java b/src/main/java/com/aparapi/device/Device.java index 670cf6ba..4408e8cf 100644 --- a/src/main/java/com/aparapi/device/Device.java +++ b/src/main/java/com/aparapi/device/Device.java @@ -15,168 +15,182 @@ */ package com.aparapi.device; -import com.aparapi.*; -import com.aparapi.internal.kernel.*; - -public abstract class Device{ - - public static enum TYPE { - UNKNOWN(Integer.MAX_VALUE), - GPU(2), - CPU(3), - JTP(5), - SEQ(6), - ACC(1), - ALT(4); - - /** Heuristic ranking of device types, lower is better. */ - public final int rank; - - TYPE(int rank) { - this.rank = rank; - } - }; - - /** @deprecated use {@link KernelManager#bestDevice()} - * @see com.aparapi.device - */ - @Deprecated - public static Device best() { - return KernelManager.instance().bestDevice(); - } - - /** - * @see com.aparapi.device - */ - @SuppressWarnings("deprecation") - @Deprecated - public static Device bestGPU() { - return firstGPU(); - } - - /** - * @see com.aparapi.device - */ - @Deprecated - public static Device first(final Device.TYPE _type) { - return KernelManager.DeprecatedMethods.firstDevice(_type); - } - - /** - * @see com.aparapi.device - */ - @SuppressWarnings("deprecation") - @Deprecated - public static Device firstGPU() { - return KernelManager.DeprecatedMethods.firstDevice(TYPE.GPU); - } - - /** - * @see com.aparapi.device - */ - @SuppressWarnings("deprecation") - @Deprecated - public static Device firstCPU() { - return KernelManager.DeprecatedMethods.firstDevice(TYPE.CPU); - } - - /** - * @see com.aparapi.device - */ - @Deprecated - public static Device bestACC() { - throw new UnsupportedOperationException(); - } - - protected TYPE type = TYPE.UNKNOWN; - - protected int maxWorkGroupSize; - - protected int maxWorkItemDimensions; - - protected int[] maxWorkItemSize = new int[] { - 0, - 0, - 0 - }; - - public abstract String getShortDescription(); - - public TYPE getType() { - return type; - } - - public void setType(TYPE type) { - this.type = type; - } - - public int getMaxWorkItemDimensions() { - return maxWorkItemDimensions; - } - - public void setMaxWorkItemDimensions(int _maxWorkItemDimensions) { - maxWorkItemDimensions = _maxWorkItemDimensions; - } - - public int getMaxWorkGroupSize() { - return maxWorkGroupSize; - } - - public void setMaxWorkGroupSize(int _maxWorkGroupSize) { - maxWorkGroupSize = _maxWorkGroupSize; - } - - public int[] getMaxWorkItemSize() { - return maxWorkItemSize; - } - - public void setMaxWorkItemSize(int[] maxWorkItemSize) { - this.maxWorkItemSize = maxWorkItemSize; - } - - public Range createRange(int _globalWidth) { - return (Range.create(this, _globalWidth)); - } - - public Range createRange(int _globalWidth, int _localWidth) { - return (Range.create(this, _globalWidth, _localWidth)); - } - - public Range createRange2D(int _globalWidth, int _globalHeight) { - return (Range.create2D(this, _globalWidth, _globalHeight)); - } - - public Range createRange2D(int _globalWidth, int _globalHeight, int _localWidth, int _localHeight) { - return (Range.create2D(this, _globalWidth, _globalHeight, _localWidth, _localHeight)); - } - - public Range createRange3D(int _globalWidth, int _globalHeight, int _globalDepth) { - return (Range.create3D(this, _globalWidth, _globalHeight, _globalDepth)); - } - - public Range createRange3D(int _globalWidth, int _globalHeight, int _globalDepth, int _localWidth, int _localHeight, - int _localDepth) { - return (Range.create3D(this, _globalWidth, _globalHeight, _globalDepth, _localWidth, _localHeight, _localDepth)); - } - - public abstract long getDeviceId(); - - @Override - public boolean equals(Object o) { - if (this == o) { - return true; - } - if (!(o instanceof Device)) { - return false; - } - - Device device = (Device) o; - - return getDeviceId() == device.getDeviceId(); - } - - @Override - public int hashCode() { - return Long.valueOf(getDeviceId()).hashCode(); - } +import com.aparapi.Range; +import com.aparapi.internal.kernel.KernelManager; + +public abstract class Device { + + public final long deviceId; + + protected Device(long deviceId) { + this.deviceId = deviceId; + } + + public enum TYPE { + UNKNOWN(Integer.MAX_VALUE), + GPU(2), + CPU(3), + JTP(5), + SEQ(6), + ACC(1), + ALT(4); + + /** + * Heuristic ranking of device types, lower is better. + */ + public final int rank; + + TYPE(int rank) { + this.rank = rank; + } + } + + /** + * @see com.aparapi.device + * @deprecated use {@link KernelManager#bestDevice()} + */ + @Deprecated + public static Device best() { + return KernelManager.instance().bestDevice(); + } + + /** + * @see com.aparapi.device + */ + @SuppressWarnings("deprecation") + @Deprecated + public static Device bestGPU() { + return firstGPU(); + } + + /** + * @see com.aparapi.device + */ + @Deprecated + public static Device first(final Device.TYPE _type) { + return KernelManager.DeprecatedMethods.firstDevice(_type); + } + + /** + * @see com.aparapi.device + */ + @SuppressWarnings("deprecation") + @Deprecated + public static Device firstGPU() { + return KernelManager.DeprecatedMethods.firstDevice(TYPE.GPU); + } + + /** + * @see com.aparapi.device + */ + @SuppressWarnings("deprecation") + @Deprecated + public static Device firstCPU() { + return KernelManager.DeprecatedMethods.firstDevice(TYPE.CPU); + } + + /** + * @see com.aparapi.device + */ + @Deprecated + public static Device bestACC() { + throw new UnsupportedOperationException(); + } + + protected TYPE type = TYPE.UNKNOWN; + + protected int maxWorkGroupSize; + + protected int maxWorkItemDimensions; + + protected int[] maxWorkItemSize = new int[]{ + 0, + 0, + 0 + }; + + public abstract String getShortDescription(); + + public TYPE getType() { + return type; + } + + protected final void setType(TYPE type) { + this.type = type; + } + + public int getMaxWorkItemDimensions() { + return maxWorkItemDimensions; + } + + public void setMaxWorkItemDimensions(int _maxWorkItemDimensions) { + maxWorkItemDimensions = _maxWorkItemDimensions; + } + + public int getMaxWorkGroupSize() { + return maxWorkGroupSize; + } + + public void setMaxWorkGroupSize(int _maxWorkGroupSize) { + maxWorkGroupSize = _maxWorkGroupSize; + } + + public int[] getMaxWorkItemSize() { + return maxWorkItemSize; + } + + public void setMaxWorkItemSize(int[] maxWorkItemSize) { + this.maxWorkItemSize = maxWorkItemSize; + } + + public Range createRange(int _globalWidth) { + return (Range.create(this, _globalWidth)); + } + + public Range createRange(int _globalWidth, int _localWidth) { + return (Range.create(this, _globalWidth, _localWidth)); + } + + public Range createRange2D(int _globalWidth, int _globalHeight) { + return (Range.create2D(this, _globalWidth, _globalHeight)); + } + + public Range createRange2D(int _globalWidth, int _globalHeight, int _localWidth, int _localHeight) { + return (Range.create2D(this, _globalWidth, _globalHeight, _localWidth, _localHeight)); + } + + public Range createRange3D(int _globalWidth, int _globalHeight, int _globalDepth) { + return (Range.create3D(this, _globalWidth, _globalHeight, _globalDepth)); + } + + public Range createRange3D(int _globalWidth, int _globalHeight, int _globalDepth, int _localWidth, int _localHeight, + int _localDepth) { + return (Range.create3D(this, _globalWidth, _globalHeight, _globalDepth, _localWidth, _localHeight, _localDepth)); + } + + public final long getDeviceId() { + return deviceId; + } + + @Override + public boolean equals(Object o) { + if (this == o) { + return true; + } + if (!(o instanceof Device)) { + return false; + } + + Device device = (Device) o; + + return deviceId == device.deviceId; + } + + @Override + public int hashCode() { + long i = deviceId; + return (int) (i ^ (i >>> 32)); + //return Long.valueOf(getDeviceId()).hashCode(); + } + } diff --git a/src/main/java/com/aparapi/device/JavaDevice.java b/src/main/java/com/aparapi/device/JavaDevice.java index b7c9ab47..ee379872 100644 --- a/src/main/java/com/aparapi/device/JavaDevice.java +++ b/src/main/java/com/aparapi/device/JavaDevice.java @@ -22,11 +22,10 @@ public class JavaDevice extends Device { public static final JavaDevice SEQUENTIAL = new JavaDevice(TYPE.SEQ, "Java Sequential", -1); private final String name; - private final long deviceId; private JavaDevice(TYPE _type, String _name, long deviceId) { - this.deviceId = deviceId; - this.type = _type; + super(deviceId); + setType(_type); this.name = _name; } @@ -35,11 +34,6 @@ public String getShortDescription() { return name; } - @Override - public long getDeviceId() { - return deviceId; - } - @Override public String toString() { return getShortDescription(); diff --git a/src/main/java/com/aparapi/device/OpenCLDevice.java b/src/main/java/com/aparapi/device/OpenCLDevice.java index ac27440f..1dac4ab0 100644 --- a/src/main/java/com/aparapi/device/OpenCLDevice.java +++ b/src/main/java/com/aparapi/device/OpenCLDevice.java @@ -25,10 +25,7 @@ import java.lang.reflect.InvocationHandler; import java.lang.reflect.Method; import java.lang.reflect.Proxy; -import java.util.ArrayList; -import java.util.HashMap; -import java.util.List; -import java.util.Map; +import java.util.*; import com.aparapi.Range; import com.aparapi.internal.opencl.OpenCLArgDescriptor; @@ -45,125 +42,126 @@ import com.aparapi.opencl.OpenCL.Resource; import com.aparapi.opencl.OpenCL.Source; -public class OpenCLDevice extends Device{ +public class OpenCLDevice extends Device { - private final OpenCLPlatform platform; + private final OpenCLPlatform platform; - private final long deviceId; + private int maxComputeUnits; - private int maxComputeUnits; + private long localMemSize; - private long localMemSize; + private long globalMemSize; - private long globalMemSize; - - private long maxMemAllocSize; - - private String shortDescription = null; - - private String name = null; - - /** - * Minimal constructor - * - * @param _platform - * @param _deviceId - * @param _type - */ - public OpenCLDevice(OpenCLPlatform _platform, long _deviceId, TYPE _type) { - platform = _platform; - deviceId = _deviceId; - type = _type; - } - - public OpenCLPlatform getOpenCLPlatform() { - return platform; - } - - public int getMaxComputeUnits() { - return maxComputeUnits; - } - - public void setMaxComputeUnits(int _maxComputeUnits) { - maxComputeUnits = _maxComputeUnits; - } - - public long getLocalMemSize() { - return localMemSize; - } - - public void setLocalMemSize(long _localMemSize) { - localMemSize = _localMemSize; - } - - public long getMaxMemAllocSize() { - return maxMemAllocSize; - } - - public void setMaxMemAllocSize(long _maxMemAllocSize) { - maxMemAllocSize = _maxMemAllocSize; - } - - public long getGlobalMemSize() { - return globalMemSize; - } - - public void setGlobalMemSize(long _globalMemSize) { - globalMemSize = _globalMemSize; - } - - void setMaxWorkItemSize(int _dim, int _value) { - maxWorkItemSize[_dim] = _value; - } - - public String getName() { - return name; - } - - public void setName(String name) { - this.name = name; - } - - @Override - public long getDeviceId() { - return (deviceId); - } - - @Override - public String getShortDescription() { - if (shortDescription == null) { - String vendor = platform.getName(); - // Hopefully(!) this equates to the recognisable name of the vendor, e.g. "Intel", "NVIDIA", "AMD" - // Note, it is not necessarily the hardware vendor, e.g. if the AMD CPU driver (i.e. platform) is used for an Intel CPU, this will be "AMD" - String[] split = vendor.split("[\\s\\(\\)]"); // split on whitespace or on '(' or ')' since Intel use "Intel(R)" here - shortDescription = split[0] + "<" + getType() + ">"; - } - return shortDescription; - } - - public static class OpenCLInvocationHandler> implements InvocationHandler{ - private final Map map; - - private final OpenCLProgram program; - private boolean disposed = false; - public OpenCLInvocationHandler(OpenCLProgram _program, Map _map) { - program = _program; - map = _map; - disposed = false; - } - - @Override public Object invoke(Object proxy, Method method, Object[] args) throws Throwable { - if (disposed){ - throw new IllegalStateException("bound interface already disposed"); - } - if (!isReservedInterfaceMethod(method)) { - final OpenCLKernel kernel = map.get(method.getName()); - if (kernel != null) { - kernel.invoke(args); + private long maxMemAllocSize; + + private String shortDescription = null; + + private String name = null; + + /** + * Minimal constructor + * + * @param _platform + * @param _deviceId + * @param _type + */ + public OpenCLDevice(OpenCLPlatform _platform, long _deviceId, TYPE _type) { + super(_deviceId); + platform = _platform; + setType(_type); + } + + public OpenCLPlatform getOpenCLPlatform() { + return platform; + } + + public int getMaxComputeUnits() { + return maxComputeUnits; + } + + public void setMaxComputeUnits(int _maxComputeUnits) { + maxComputeUnits = _maxComputeUnits; + } + + public long getLocalMemSize() { + return localMemSize; + } + + public void setLocalMemSize(long _localMemSize) { + localMemSize = _localMemSize; + } + + public long getMaxMemAllocSize() { + return maxMemAllocSize; + } + + public void setMaxMemAllocSize(long _maxMemAllocSize) { + maxMemAllocSize = _maxMemAllocSize; + } + + public long getGlobalMemSize() { + return globalMemSize; + } + + public void setGlobalMemSize(long _globalMemSize) { + globalMemSize = _globalMemSize; + } + + void setMaxWorkItemSize(int _dim, int _value) { + maxWorkItemSize[_dim] = _value; + } + + public String getName() { + return name; + } + + public void setName(String name) { + this.name = name; + } + + @Override + public String getShortDescription() { + if (shortDescription == null) { + String vendor = getName(); + // Hopefully(!) this equates to the recognisable name of the vendor, e.g. "Intel", "NVIDIA", "AMD" + // Note, it is not necessarily the hardware vendor, e.g. if the AMD CPU driver (i.e. platform) is used for an Intel CPU, this will be "AMD" + String[] split = vendor.split("[\\s\\(\\)]"); // split on whitespace or on '(' or ')' since Intel use "Intel(R)" here + if (split.length == 0 || split[0] == null + || !(split[0].equals("Intel") || split[0].equals("NVIDIA") || split[0].equals("AMD"))) { + vendor = platform.getName(); + split = vendor.split("[\\s\\(\\)]"); } - } else { - if (method.getName().equals("put")) { - System.out.println("put not implemented"); + shortDescription = split[0] + '<' + getType() + '>'; + } + return shortDescription; + } + + static class OpenCLInvocationHandler > implements InvocationHandler { + private final Map map; + + private final OpenCLProgram program; + private boolean disposed = false; + + OpenCLInvocationHandler(OpenCLProgram _program, Map _map) { + program = _program; + map = _map; + disposed = false; + } + + @Override + public Object invoke(Object proxy, Method method, Object[] args) { + if (disposed) { + throw new IllegalStateException("bound interface already disposed"); + } + if (!isReservedInterfaceMethod(method)) { + final OpenCLKernel kernel = map.get(method.getName()); + if (kernel != null) { + kernel.invoke(args); + } + } else { + switch (method.getName()) { + case "put": + throw new UnsupportedOperationException("put not implemented"); /* for (Object arg : args) { @@ -186,8 +184,9 @@ public OpenCLInvocationHandler(OpenCLProgram _program, Map } } */ - } else if (method.getName().equals("get")) { - System.out.println("get not implemented"); + + case "get": + throw new UnsupportedOperationException("get not implemented"); /* for (Object arg : args) { Class> argClass = arg.getClass(); @@ -209,337 +208,344 @@ public OpenCLInvocationHandler(OpenCLProgram _program, Map } } */ - } else if (method.getName().equals("begin")) { - System.out.println("begin not implemented"); - } else if (method.getName().equals("dispose")) { - // System.out.println("dispose"); - for (OpenCLKernel k:map.values()){ - k.dispose(); - } - program.dispose(); - map.clear(); - disposed=true; - } else if (method.getName().equals("end")) { - System.out.println("end not implemented"); - } else if (method.getName().equals("getProfileInfo")){ - proxy = program.getProfileInfo(); - } - } - return proxy; - } - } - - public List getArgs(Method m) { - final List args = new ArrayList (); - final Annotation[][] parameterAnnotations = m.getParameterAnnotations(); - final Class>[] parameterTypes = m.getParameterTypes(); - - for (int arg = 0; arg < parameterTypes.length; arg++) { - if (parameterTypes[arg].isAssignableFrom(Range.class)) { - - } else { - - long bits = 0L; - String name = null; - for (final Annotation pa : parameterAnnotations[arg]) { - if (pa instanceof GlobalReadOnly) { - name = ((GlobalReadOnly) pa).value(); - bits |= OpenCLArgDescriptor.ARG_GLOBAL_BIT | OpenCLArgDescriptor.ARG_READONLY_BIT; - } else if (pa instanceof GlobalWriteOnly) { - name = ((GlobalWriteOnly) pa).value(); - bits |= OpenCLArgDescriptor.ARG_GLOBAL_BIT | OpenCLArgDescriptor.ARG_WRITEONLY_BIT; - } else if (pa instanceof GlobalReadWrite) { - name = ((GlobalReadWrite) pa).value(); - bits |= OpenCLArgDescriptor.ARG_GLOBAL_BIT | OpenCLArgDescriptor.ARG_READWRITE_BIT; - } else if (pa instanceof Local) { - name = ((Local) pa).value(); - bits |= OpenCLArgDescriptor.ARG_LOCAL_BIT; - } else if (pa instanceof Constant) { - name = ((Constant) pa).value(); - bits |= OpenCLArgDescriptor.ARG_CONST_BIT | OpenCLArgDescriptor.ARG_READONLY_BIT; - } else if (pa instanceof Arg) { - name = ((Arg) pa).value(); - bits |= OpenCLArgDescriptor.ARG_ISARG_BIT; - } + case "begin": + throw new UnsupportedOperationException("begin not implemented"); + case "dispose": + // System.out.println("dispose"); + for (OpenCLKernel k : map.values()) + k.dispose(); + map.clear(); + program.dispose(); + disposed = true; + break; + case "end": + throw new UnsupportedOperationException("end not implemented"); + + case "getProfileInfo": + proxy = program.getProfileInfo(); + break; + } } - if (parameterTypes[arg].isArray()) { - if (parameterTypes[arg].isAssignableFrom(float[].class)) { - bits |= OpenCLArgDescriptor.ARG_FLOAT_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; - } else if (parameterTypes[arg].isAssignableFrom(int[].class)) { - bits |= OpenCLArgDescriptor.ARG_INT_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; - } else if (parameterTypes[arg].isAssignableFrom(double[].class)) { - bits |= OpenCLArgDescriptor.ARG_DOUBLE_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; - } else if (parameterTypes[arg].isAssignableFrom(byte[].class)) { - bits |= OpenCLArgDescriptor.ARG_BYTE_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; - } else if (parameterTypes[arg].isAssignableFrom(short[].class)) { - bits |= OpenCLArgDescriptor.ARG_SHORT_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; - } else if (parameterTypes[arg].isAssignableFrom(long[].class)) { - bits |= OpenCLArgDescriptor.ARG_LONG_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; - } - } else if (parameterTypes[arg].isPrimitive()) { - if (parameterTypes[arg].isAssignableFrom(float.class)) { - bits |= OpenCLArgDescriptor.ARG_FLOAT_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; - } else if (parameterTypes[arg].isAssignableFrom(int.class)) { - bits |= OpenCLArgDescriptor.ARG_INT_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; - } else if (parameterTypes[arg].isAssignableFrom(double.class)) { - bits |= OpenCLArgDescriptor.ARG_DOUBLE_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; - } else if (parameterTypes[arg].isAssignableFrom(byte.class)) { - bits |= OpenCLArgDescriptor.ARG_BYTE_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; - } else if (parameterTypes[arg].isAssignableFrom(short.class)) { - bits |= OpenCLArgDescriptor.ARG_SHORT_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; - } else if (parameterTypes[arg].isAssignableFrom(long.class)) { - bits |= OpenCLArgDescriptor.ARG_LONG_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; - } - } else { - System.out.println("OUch!"); - } - if (name == null) { - throw new IllegalStateException("no name!"); - } - final OpenCLArgDescriptor kernelArg = new OpenCLArgDescriptor(name, bits); - args.add(kernelArg); + return proxy; + } + } - } - } + private static List getArgs(Method m) { + final Annotation[][] parameterAnnotations = m.getParameterAnnotations(); + final Class>[] parameterTypes = m.getParameterTypes(); - return (args); - } + final List args = new ArrayList<>(parameterTypes.length); + for (int arg = 0; arg < parameterTypes.length; arg++) { + Class> t = parameterTypes[arg]; + if (t.isAssignableFrom(Range.class)) { - private static boolean isReservedInterfaceMethod(Method _methods) { - return ( _methods.getName().equals("put") - || _methods.getName().equals("get") - || _methods.getName().equals("dispose") - || _methods.getName().equals("begin") - || _methods.getName().equals("end") - || _methods.getName().equals("getProfileInfo")); - } - - private String streamToString(InputStream _inputStream) { - final StringBuilder sourceBuilder = new StringBuilder(); - - if (_inputStream != null) { - - final BufferedReader reader = new BufferedReader(new InputStreamReader(_inputStream)); + } else { - try { - for (String line = reader.readLine(); line != null; line = reader.readLine()) { - sourceBuilder.append(line).append("\n"); + long bits = 0L; + String name = null; + for (final Annotation pa : parameterAnnotations[arg]) { + if (pa instanceof GlobalReadOnly) { + name = ((GlobalReadOnly) pa).value(); + bits |= OpenCLArgDescriptor.ARG_GLOBAL_BIT | OpenCLArgDescriptor.ARG_READONLY_BIT; + } else if (pa instanceof GlobalWriteOnly) { + name = ((GlobalWriteOnly) pa).value(); + bits |= OpenCLArgDescriptor.ARG_GLOBAL_BIT | OpenCLArgDescriptor.ARG_WRITEONLY_BIT; + } else if (pa instanceof GlobalReadWrite) { + name = ((GlobalReadWrite) pa).value(); + bits |= OpenCLArgDescriptor.ARG_GLOBAL_BIT | OpenCLArgDescriptor.ARG_READWRITE_BIT; + } else if (pa instanceof Local) { + name = ((Local) pa).value(); + bits |= OpenCLArgDescriptor.ARG_LOCAL_BIT; + } else if (pa instanceof Constant) { + name = ((Constant) pa).value(); + bits |= OpenCLArgDescriptor.ARG_CONST_BIT | OpenCLArgDescriptor.ARG_READONLY_BIT; + } else if (pa instanceof Arg) { + name = ((Arg) pa).value(); + bits |= OpenCLArgDescriptor.ARG_ISARG_BIT; + } + + } + if (t.isArray()) { + if (t.isAssignableFrom(float[].class)) { + bits |= OpenCLArgDescriptor.ARG_FLOAT_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; + } else if (t.isAssignableFrom(int[].class)) { + bits |= OpenCLArgDescriptor.ARG_INT_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; + } else if (t.isAssignableFrom(double[].class)) { + bits |= OpenCLArgDescriptor.ARG_DOUBLE_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; + } else if (t.isAssignableFrom(byte[].class)) { + bits |= OpenCLArgDescriptor.ARG_BYTE_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; + } else if (t.isAssignableFrom(short[].class)) { + bits |= OpenCLArgDescriptor.ARG_SHORT_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; + } else if (t.isAssignableFrom(long[].class)) { + bits |= OpenCLArgDescriptor.ARG_LONG_BIT | OpenCLArgDescriptor.ARG_ARRAY_BIT; + } + } else if (t.isPrimitive()) { + if (t.isAssignableFrom(float.class)) { + bits |= OpenCLArgDescriptor.ARG_FLOAT_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; + } else if (t.isAssignableFrom(int.class)) { + bits |= OpenCLArgDescriptor.ARG_INT_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; + } else if (t.isAssignableFrom(double.class)) { + bits |= OpenCLArgDescriptor.ARG_DOUBLE_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; + } else if (t.isAssignableFrom(byte.class)) { + bits |= OpenCLArgDescriptor.ARG_BYTE_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; + } else if (t.isAssignableFrom(short.class)) { + bits |= OpenCLArgDescriptor.ARG_SHORT_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; + } else if (t.isAssignableFrom(long.class)) { + bits |= OpenCLArgDescriptor.ARG_LONG_BIT | OpenCLArgDescriptor.ARG_PRIMITIVE_BIT; + } + } else { + //System.out.println("OUch!"); + throw new IllegalStateException(t + " unsupported type"); + } + if (name == null) { + throw new IllegalStateException("no name!"); + } + + args.add(new OpenCLArgDescriptor(name, bits)); } - } catch (final IOException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } - - try { - _inputStream.close(); - } catch (final IOException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } - } - - - return (sourceBuilder.toString()); - } - - public > T bind(Class _interface, InputStream _inputStream) { - return (bind(_interface, streamToString(_inputStream))); - } - - public > T bind(Class _interface) { - return (bind(_interface, (String) null)); - } - - public > T bind(Class _interface, String _source) { - final Map > kernelNameToArgsMap = new HashMap >(); - - if (_source == null) { - final StringBuilder sourceBuilder = new StringBuilder(); - boolean interfaceIsAnnotated = false; - for (final Annotation a : _interface.getAnnotations()) { - if (a instanceof Source) { - final Source source = (Source) a; - sourceBuilder.append(source.value()).append("\n"); - interfaceIsAnnotated = true; - } else if (a instanceof Resource) { - final Resource sourceResource = (Resource) a; - final InputStream stream = _interface.getClassLoader().getResourceAsStream(sourceResource.value()); - sourceBuilder.append(streamToString(stream)); - interfaceIsAnnotated = true; + } + + return (args); + } + + final static Set reservedInterfaceMethods = new HashSet(6) {{ + add("put"); + add("get"); + add("dispose"); + add("begin"); + add("end"); + add("getProfileInfo"); + }}; + + private static boolean isReservedInterfaceMethod(Method _methods) { + return reservedInterfaceMethods.contains(_methods.getName()); + } + + private static String streamToString(InputStream _inputStream) { + final StringBuilder sourceBuilder = new StringBuilder(); + + if (_inputStream != null) { + + final BufferedReader reader = new BufferedReader(new InputStreamReader(_inputStream)); + + try { + for (String line = reader.readLine(); line != null; line = reader.readLine()) { + sourceBuilder.append(line).append('\n'); + } + } catch (final IOException e) { + // TODO Auto-generated catch block + e.printStackTrace(); } - } - if (interfaceIsAnnotated) { - // just crawl the methods (non put or get) and create kernels - for (final Method m : _interface.getDeclaredMethods()) { - if (!isReservedInterfaceMethod(m)) { - final List args = getArgs(m); - kernelNameToArgsMap.put(m.getName(), args); - } + try { + _inputStream.close(); + } catch (final IOException e) { + // TODO Auto-generated catch block + e.printStackTrace(); + } + } + + + return (sourceBuilder.toString()); + } + + public > T bind(Class _interface, InputStream _inputStream) { + return (bind(_interface, streamToString(_inputStream))); + } + + public > T bind(Class _interface) { + return (bind(_interface, (String) null)); + } + + public > T bind(Class _interface, String _source) { + final Map > kernelNameToArgsMap = new HashMap<>(); + + if (_source == null) { + final StringBuilder sourceBuilder = new StringBuilder(); + boolean interfaceIsAnnotated = false; + for (final Annotation a : _interface.getAnnotations()) { + if (a instanceof Source) { + final Source source = (Source) a; + sourceBuilder.append(source.value()).append('\n'); + interfaceIsAnnotated = true; + } else if (a instanceof Resource) { + final Resource sourceResource = (Resource) a; + final InputStream stream = _interface.getClassLoader().getResourceAsStream(sourceResource.value()); + sourceBuilder.append(streamToString(stream)); + interfaceIsAnnotated = true; + } } - } else { - for (final Method m : _interface.getDeclaredMethods()) { - if (!isReservedInterfaceMethod(m)) { - for (final Annotation a : m.getAnnotations()) { - // System.out.println(" annotation "+a); - // System.out.println(" annotation type " + a.annotationType()); - if (a instanceof Kernel) { - sourceBuilder.append("__kernel void " + m.getName() + "("); + if (interfaceIsAnnotated) { + // just crawl the methods (non put or get) and create kernels + for (final Method m : _interface.getDeclaredMethods()) { + if (!isReservedInterfaceMethod(m)) { final List args = getArgs(m); + kernelNameToArgsMap.put(m.getName(), args); + } + } + } else { - boolean first = true; - for (final OpenCLArgDescriptor arg : args) { - if (first) { - first = false; - } else { - sourceBuilder.append(","); - } - sourceBuilder.append("\n " + arg); + for (final Method m : _interface.getDeclaredMethods()) { + if (!isReservedInterfaceMethod(m)) { + for (final Annotation a : m.getAnnotations()) { + // System.out.println(" annotation "+a); + // System.out.println(" annotation type " + a.annotationType()); + if (a instanceof Kernel) { + sourceBuilder.append("__kernel void ").append(m.getName()).append('('); + final List args = getArgs(m); + + boolean first = true; + for (final OpenCLArgDescriptor arg : args) { + if (first) { + first = false; + } else { + sourceBuilder.append(','); + } + sourceBuilder.append("\n ").append(arg); + } + + sourceBuilder.append(')'); + final Kernel kernel = (Kernel) a; + sourceBuilder.append(kernel.value()); + kernelNameToArgsMap.put(m.getName(), args); + + } } + } + } - sourceBuilder.append(")"); - final Kernel kernel = (Kernel) a; - sourceBuilder.append(kernel.value()); - kernelNameToArgsMap.put(m.getName(), args); - - } - } - } } - - } - _source = sourceBuilder.toString(); - } else { - for (final Method m : _interface.getDeclaredMethods()) { - if (!isReservedInterfaceMethod(m)) { - final List args = getArgs(m); - kernelNameToArgsMap.put(m.getName(), args); + _source = sourceBuilder.toString(); + } else { + for (final Method m : _interface.getDeclaredMethods()) { + if (!isReservedInterfaceMethod(m)) { + kernelNameToArgsMap.put(m.getName(), getArgs(m)); + } } - } - } + } - final OpenCLProgram program = new OpenCLProgram(this, _source).createProgram(this); + final OpenCLProgram program = new OpenCLProgram(this, _source).createProgram(this); - final Map map = new HashMap (); - for (final String name : kernelNameToArgsMap.keySet()) { - final OpenCLKernel kernel = OpenCLKernel.createKernel(program, name, kernelNameToArgsMap.get(name)); - //final OpenCLKernel kernel = new OpenCLKernel(program, name, kernelNameToArgsMap.get(name)); - if (kernel == null) { - throw new IllegalStateException("kernel is null"); - } + final Map map = new HashMap<>(); + for (final Map.Entry > stringListEntry : kernelNameToArgsMap.entrySet()) { + final OpenCLKernel kernel = OpenCLKernel.createKernel(program, stringListEntry.getKey(), stringListEntry.getValue()); + //final OpenCLKernel kernel = new OpenCLKernel(program, name, kernelNameToArgsMap.get(name)); + if (kernel == null) { + throw new IllegalStateException("kernel is null"); + } - map.put(name, kernel); - } + map.put(stringListEntry.getKey(), kernel); + } - final OpenCLInvocationHandler invocationHandler = new OpenCLInvocationHandler (program, map); - final T instance = (T) Proxy.newProxyInstance(OpenCLDevice.class.getClassLoader(), new Class[] { + final OpenCLInvocationHandler invocationHandler = new OpenCLInvocationHandler<>(program, map); + final T instance = (T) Proxy.newProxyInstance(OpenCLDevice.class.getClassLoader(), new Class[]{ _interface, OpenCL.class - }, invocationHandler); + }, invocationHandler); - return instance; - } + return instance; + } - public interface DeviceSelector{ - OpenCLDevice select(OpenCLDevice _device); - } + public interface DeviceSelector { + OpenCLDevice select(OpenCLDevice _device); + } - public interface DeviceComparitor{ - OpenCLDevice select(OpenCLDevice _deviceLhs, OpenCLDevice _deviceRhs); - } + public interface DeviceComparitor { + OpenCLDevice select(OpenCLDevice _deviceLhs, OpenCLDevice _deviceRhs); + } - /** List OpenCLDevices of a given TYPE, or all OpenCLDevices if type == null. */ - public static List listDevices(TYPE type) { - final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); - final ArrayList results = new ArrayList<>(); + /** List OpenCLDevices of a given TYPE, or all OpenCLDevices if type == null. */ + public static List listDevices(TYPE type, Iterable platforms) { + final ArrayList results = new ArrayList<>(); - for (final OpenCLPlatform p : platform.getOpenCLPlatforms()) { - for (final OpenCLDevice device : p.getOpenCLDevices()) { - if (type == null || device.getType() == type) { - results.add(device); + for (final OpenCLPlatform p : platforms) { + for (final OpenCLDevice device : p.getOpenCLDevices()) { + if (type == null || device.getType() == type) { + results.add(device); + } } - } - } + } - return results; - } + return results; + } - public static OpenCLDevice select(DeviceSelector _deviceSelector) { - OpenCLDevice device = null; - final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); + public static OpenCLDevice select(DeviceSelector _deviceSelector) { + OpenCLDevice device = null; + final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); - for (final OpenCLPlatform p : platform.getOpenCLPlatforms()) { - for (final OpenCLDevice d : p.getOpenCLDevices()) { - device = _deviceSelector.select(d); + for (final OpenCLPlatform p : platform.getOpenCLPlatforms()) { + for (final OpenCLDevice d : p.getOpenCLDevices()) { + device = _deviceSelector.select(d); + if (device != null) { + break; + } + } if (device != null) { - break; + break; } - } - if (device != null) { - break; - } - } - - return (device); - } - - public static OpenCLDevice select(DeviceComparitor _deviceComparitor) { - OpenCLDevice device = null; - final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); - - List openCLPlatforms = platform.getOpenCLPlatforms(); - for (final OpenCLPlatform p : openCLPlatforms) { - List openCLDevices = p.getOpenCLDevices(); - for (final OpenCLDevice d : openCLDevices) { - if (device == null) { - device = d; - } else { - device = _deviceComparitor.select(device, d); + } + + return (device); + } + + public static OpenCLDevice select(DeviceComparitor _deviceComparitor) { + OpenCLDevice device = null; + final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); + + List openCLPlatforms = platform.getOpenCLPlatforms(); + for (final OpenCLPlatform p : openCLPlatforms) { + List openCLDevices = p.getOpenCLDevices(); + for (final OpenCLDevice d : openCLDevices) { + if (device == null) { + device = d; + } else { + device = _deviceComparitor.select(device, d); + } } - } - } - - return (device); - } - - public static OpenCLDevice select(DeviceComparitor _deviceComparitor, Device.TYPE _type) { - OpenCLDevice device = null; - final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); - - for (final OpenCLPlatform p : platform.getOpenCLPlatforms()) { - for (final OpenCLDevice d : p.getOpenCLDevices()) { - if (d.getType() != _type) continue; - if (device == null) { - device = d; + } + + return (device); + } + + public static OpenCLDevice select(DeviceComparitor _deviceComparitor, Device.TYPE _type) { + OpenCLDevice device = null; + final OpenCLPlatform platform = new OpenCLPlatform(0, null, null, null); + + for (final OpenCLPlatform p : platform.getOpenCLPlatforms()) { + for (final OpenCLDevice d : p.getOpenCLDevices()) { + if (d.getType() != _type) continue; + if (device == null) { + device = d; + } else { + device = _deviceComparitor.select(device, d); + } + } + } + + return (device); + } + + @Override + public String toString() { + final StringBuilder s = new StringBuilder("{"); + boolean first = true; + for (final int workItemSize : maxWorkItemSize) { + if (first) { + first = false; } else { - device = _deviceComparitor.select(device, d); + s.append(", "); } - } - } - - return (device); - } - - @Override public String toString() { - final StringBuilder s = new StringBuilder("{"); - boolean first = true; - for (final int workItemSize : maxWorkItemSize) { - if (first) { - first = false; - } else { - s.append(", "); - } - s.append(workItemSize); - } + s.append(workItemSize); + } - s.append("}"); + s.append('}'); - return ("Device " + deviceId + "\n vendor = " + getOpenCLPlatform().getVendor() + return ("Device " + deviceId + "\n vendor = " + getOpenCLPlatform().getVendor() + "\n type:" + type + "\n maxComputeUnits=" + maxComputeUnits + "\n maxWorkItemDimensions=" + maxWorkItemDimensions + "\n maxWorkItemSizes=" + s + "\n maxWorkWorkGroupSize=" + maxWorkGroupSize + "\n globalMemSize=" + globalMemSize + "\n localMemSize=" + localMemSize); - } + } } diff --git a/src/main/java/com/aparapi/internal/exception/ClassParseException.java b/src/main/java/com/aparapi/internal/exception/ClassParseException.java index 39aac1fa..a38e8261 100644 --- a/src/main/java/com/aparapi/internal/exception/ClassParseException.java +++ b/src/main/java/com/aparapi/internal/exception/ClassParseException.java @@ -65,7 +65,7 @@ to national security controls as identified on the Commerce Control List (curren */ @SuppressWarnings("serial") public class ClassParseException extends AparapiException{ - public static enum TYPE { + public enum TYPE { NONE("none"), // ARRAY_RETURN("We don't support areturn instructions"), // PUTFIELD("We don't support putstatic instructions"), // @@ -103,18 +103,18 @@ public static enum TYPE { MISSINGLOCALVARIABLETABLE("Method does not contain a local variable table (recompile with -g?)"), // IMPROPERPRIVATENAMEMANGLING("Could not parse private array size from field name"); - private String description; + private final String description; TYPE(final String _description) { description = _description; } - public String getDescription() { + String getDescription() { return (description); } - }; + } - private Instruction instruction; + private Instruction instruction; private TYPE type; @@ -125,13 +125,13 @@ public ClassParseException(final TYPE _type) { } public ClassParseException(final Instruction _instruction, final TYPE _type) { - super("@" + _instruction.getThisPC() + " " + _instruction.getByteCode() + " " + _type.getDescription()); + super("@" + _instruction.getThisPC() + ' ' + _instruction.getByteCode() + ' ' + _type.getDescription()); type = _type; instruction = _instruction; } public ClassParseException(final TYPE _type, final String _methodName) { - super("@" + _methodName + " " + _type.getDescription()); + super('@' + _methodName + ' ' + _type.getDescription()); type = _type; instruction = null; } diff --git a/src/main/java/com/aparapi/internal/exception/RangeException.java b/src/main/java/com/aparapi/internal/exception/RangeException.java index f8b8c757..ce0afeab 100644 --- a/src/main/java/com/aparapi/internal/exception/RangeException.java +++ b/src/main/java/com/aparapi/internal/exception/RangeException.java @@ -52,7 +52,8 @@ to national security controls as identified on the Commerce Control List (curren */ package com.aparapi.internal.exception; -@SuppressWarnings("serial") public class RangeException extends AparapiException{ +@SuppressWarnings("serial") +class RangeException extends AparapiException{ public RangeException(String msg) { super(msg); diff --git a/src/main/java/com/aparapi/internal/instruction/BranchSet.java b/src/main/java/com/aparapi/internal/instruction/BranchSet.java index a0437b55..f9044aa8 100644 --- a/src/main/java/com/aparapi/internal/instruction/BranchSet.java +++ b/src/main/java/com/aparapi/internal/instruction/BranchSet.java @@ -127,19 +127,19 @@ public static abstract class LogicalExpressionNode { private LogicalExpressionNode parent = null; - public void setParent(LogicalExpressionNode _parent) { + void setParent(LogicalExpressionNode _parent) { parent = _parent; } - public abstract int getTarget(); + protected abstract int getTarget(); - public abstract int getFallThrough(); + protected abstract int getFallThrough(); - public abstract void invert(); + protected abstract void invert(); public abstract LogicalExpressionNode cloneInverted(); - public LogicalExpressionNode getRoot() { + LogicalExpressionNode getRoot() { if (parent != null) { return (parent); } else { @@ -147,11 +147,11 @@ public LogicalExpressionNode getRoot() { } } - public LogicalExpressionNode getNext() { + LogicalExpressionNode getNext() { return (next == null ? next : next.getRoot()); } - public void setNext(LogicalExpressionNode _next) { + void setNext(LogicalExpressionNode _next) { next = _next == null ? _next : _next.getRoot(); } @@ -176,7 +176,7 @@ public static class SimpleLogicalExpressionNode extends LogicalExpressionNode { private boolean invert; - public SimpleLogicalExpressionNode(ConditionalBranch _branch) { + SimpleLogicalExpressionNode(ConditionalBranch _branch) { this(_branch, false); } @@ -215,7 +215,7 @@ public ConditionalBranch getBranch() { @Override public String toString() { - return invert ? ("!(" + getBranch() + ")") : getBranch().toString(); + return invert ? ("!(" + getBranch() + ')') : getBranch().toString(); } } @@ -253,7 +253,7 @@ private CompoundLogicalExpressionNode(boolean _and, LogicalExpressionNode _lhs, lhs.setParent(this); } - public CompoundLogicalExpressionNode(boolean _and, LogicalExpressionNode _lhs, LogicalExpressionNode _rhs) { + CompoundLogicalExpressionNode(boolean _and, LogicalExpressionNode _lhs, LogicalExpressionNode _rhs) { this(_and, _lhs, _rhs, true); } @@ -295,12 +295,12 @@ public LogicalExpressionNode getRhs() { @Override public String toString() { - return getLhs().toString() + " " + (isAnd() ? "&&" : "||") + " " + getRhs().toString(); + return getLhs().toString() + ' ' + (isAnd() ? "&&" : "||") + ' ' + getRhs().toString(); } } - private final List set = new ArrayList (); + private final List set = new ArrayList<>(); private final Instruction fallThrough; @@ -323,7 +323,7 @@ public BranchSet(Branch _branch) { target = _branch.getTarget(); last = _branch; - final Set expandedSet = new LinkedHashSet (); + final Set expandedSet = new LinkedHashSet<>(); final Instruction fallThroughRoot = last.getNextExpr(); fallThrough = fallThroughRoot == null ? last.getNextPC() : fallThroughRoot.getStartInstruction(); first = last; diff --git a/src/main/java/com/aparapi/internal/instruction/ExpressionList.java b/src/main/java/com/aparapi/internal/instruction/ExpressionList.java index 2aac1c00..e3134337 100644 --- a/src/main/java/com/aparapi/internal/instruction/ExpressionList.java +++ b/src/main/java/com/aparapi/internal/instruction/ExpressionList.java @@ -83,7 +83,7 @@ to national security controls as identified on the Commerce Control List (curren */ public class ExpressionList{ - private static Logger logger = Logger.getLogger(Config.getLoggerName()); + private static final Logger logger = Logger.getLogger(Config.getLoggerName()); private final MethodModel methodModel; @@ -124,12 +124,13 @@ public ExpressionList(MethodModel _methodModel) { * @param _extent * @return */ - public boolean doesNotContainContinueOrBreak(Instruction _start, Instruction _extent) { + private static boolean doesNotContainContinueOrBreak(Instruction _start, Instruction _extent) { boolean ok = true; boolean breakOrContinue = false; for (Instruction i = _start; i != null; i = i.getNextExpr()) { if (i.isBranch()) { - if (i.asBranch().isForwardUnconditional() && i.asBranch().getTarget().isAfter(_extent)) { + Branch br = i.asBranch(); + if (br.isForwardUnconditional() && br.getTarget().isAfter(_extent)) { breakOrContinue = true; } else { ok = false; @@ -149,7 +150,7 @@ public boolean doesNotContainContinueOrBreak(Instruction _start, Instruction _ex return (ok); } - public boolean doesNotContainCompositeOrBranch(Instruction _start, Instruction _exclusiveEnd) { + private static boolean doesNotContainCompositeOrBranch(Instruction _start, Instruction _exclusiveEnd) { boolean ok = true; for (Instruction i = _start; (i != null) && (i != _exclusiveEnd); i = i.getNextExpr()) { if (!(i instanceof CompositeInstruction) && (i.isBranch())) { @@ -160,7 +161,7 @@ public boolean doesNotContainCompositeOrBranch(Instruction _start, Instruction _ return (ok); } - public void unwind() { + private void unwind() { if (parent != null) { if (instruction != null) { tail.setNextExpr(instruction); @@ -481,7 +482,7 @@ public boolean foldComposite(final Instruction _instruction) throws ClassParseEx loopTop = loopTop.getPrevExpr(); if (loopTop instanceof AssignToLocalVariable) { - final LocalVariableInfo localVariableInfo = ((AssignToLocalVariable) loopTop).getLocalVariableInfo(); + final LocalVariableInfo localVariableInfo = ((InstructionSet.LocalVariableTableIndexAccessor) loopTop).getLocalVariableInfo(); if ((localVariableInfo.getStart() == loopTop.getNextExpr().getStartPC()) && (localVariableInfo.getEnd() == _instruction.getThisPC())) { loopTop = loopTop.getPrevExpr(); // back up over the initialization @@ -500,7 +501,7 @@ public boolean foldComposite(final Instruction _instruction) throws ClassParseEx // looptop == the unconditional? loopTop = loopTop.getPrevExpr(); if (loopTop instanceof AssignToLocalVariable) { - final LocalVariableInfo localVariableInfo = ((AssignToLocalVariable) loopTop).getLocalVariableInfo(); + final LocalVariableInfo localVariableInfo = ((InstructionSet.LocalVariableTableIndexAccessor) loopTop).getLocalVariableInfo(); if ((localVariableInfo.getStart() == loopTop.getNextExpr().getStartPC()) && (localVariableInfo.getEnd() == _instruction.getThisPC())) { loopTop = loopTop.getPrevExpr(); // back up over the initialization @@ -618,7 +619,7 @@ public boolean foldComposite(final Instruction _instruction) throws ClassParseEx Instruction loopTop = topOfRealLoop.getPrevExpr(); if (loopTop instanceof AssignToLocalVariable) { - final LocalVariableInfo localVariableInfo = ((AssignToLocalVariable) loopTop).getLocalVariableInfo(); + final LocalVariableInfo localVariableInfo = ((InstructionSet.LocalVariableTableIndexAccessor) loopTop).getLocalVariableInfo(); if ((localVariableInfo.getStart() == loopTop.getNextExpr().getStartPC()) && (localVariableInfo.getEnd() == _instruction.getThisPC())) { loopTop = loopTop.getPrevExpr(); // back up over the initialization @@ -652,7 +653,7 @@ public boolean foldComposite(final Instruction _instruction) throws ClassParseEx } if (loopTop instanceof AssignToLocalVariable) { - final LocalVariableInfo localVariableInfo = ((AssignToLocalVariable) loopTop).getLocalVariableInfo(); + final LocalVariableInfo localVariableInfo = ((InstructionSet.LocalVariableTableIndexAccessor) loopTop).getLocalVariableInfo(); if ((localVariableInfo.getStart() == loopTop.getNextExpr().getStartPC()) && (localVariableInfo.getEnd() == _instruction.getThisPC())) { loopTop = loopTop.getPrevExpr(); // back up over the initialization @@ -802,19 +803,26 @@ public boolean foldComposite(final Instruction _instruction) throws ClassParseEx for (final LocalVariableInfo localVariableInfo : localVariableTable) { if (localVariableInfo.getEnd() == _instruction.getThisPC()) { - logger.fine(localVariableInfo.getVariableName() + " scope " + localVariableInfo.getStart() + " ," - + localVariableInfo.getEnd()); + if (logger.isLoggable(Level.FINE)) { + logger.fine(localVariableInfo.getVariableName() + " scope " + localVariableInfo.getStart() + " ," + + localVariableInfo.getEnd()); + } if (localVariableInfo.getStart() < startPc) { startPc = localVariableInfo.getStart(); } } } if (startPc < Short.MAX_VALUE) { - logger.fine("Scope block from " + startPc + " to " + (tail.getThisPC() + tail.getLength())); + + if (logger.isLoggable(Level.FINE)) + logger.fine("Scope block from " + startPc + " to " + (tail.getThisPC() + tail.getLength())); + for (Instruction i = head; i != null; i = i.getNextPC()) { if (i.getThisPC() == startPc) { final Instruction startInstruction = i.getRootExpr().getPrevExpr(); - logger.fine("Start = " + startInstruction); + + if (logger.isLoggable(Level.FINE)) + logger.fine("Start = " + startInstruction); addAsComposites(ByteCode.COMPOSITE_ARBITRARY_SCOPE, startInstruction.getPrevExpr(), null); handled = true; @@ -830,8 +838,6 @@ public boolean foldComposite(final Instruction _instruction) throws ClassParseEx Config.instructionListener.showAndTell("after folding", head, _instruction); } - } catch (final ClassParseException _classParseException) { - throw new ClassParseException(_classParseException); } catch (final Throwable t) { throw new ClassParseException(t); @@ -861,7 +867,7 @@ private void addAsComposites(ByteCode _byteCode, Instruction _start, BranchSet _ */ public String dumpDiagram(Instruction _instruction) { final StringBuilder sb = new StringBuilder(); - final List list = new ArrayList (); + final List list = new ArrayList<>(); for (Instruction i = head; i != null; i = i.getNextExpr()) { list.add(i); @@ -871,20 +877,20 @@ public String dumpDiagram(Instruction _instruction) { list.add(i); } - final Instruction[] array = list.toArray(new Instruction[0]); + final Instruction[] array = list.toArray(new Instruction[list.size()]); boolean lastWasCursor = false; - final List branches = new ArrayList (); + final List branches = new ArrayList<>(); for (final Instruction i : list) { sb.append(String.format(" %3d", i.getStartPC())); } - sb.append("\n"); + sb.append('\n'); for (final Instruction i : list) { sb.append(String.format(" %3d", i.getThisPC())); } - sb.append("\n"); + sb.append('\n'); for (final Instruction i : list) { if (i == _instruction) { diff --git a/src/main/java/com/aparapi/internal/instruction/Instruction.java b/src/main/java/com/aparapi/internal/instruction/Instruction.java index 055f2424..d5a7c413 100644 --- a/src/main/java/com/aparapi/internal/instruction/Instruction.java +++ b/src/main/java/com/aparapi/internal/instruction/Instruction.java @@ -74,13 +74,13 @@ to national security controls as identified on the Commerce Control List (curren */ public abstract class Instruction{ - protected MethodModel method; + final MethodModel method; private final ByteCode byteCode; private int length; - protected int pc; + private final int pc; abstract String getDescription(); @@ -172,13 +172,13 @@ public int getStartPC() { return (getFirstChild() == null ? pc : getFirstChild().getStartPC()); } - protected Instruction(MethodModel _method, ByteCode _byteCode, int _pc) { + Instruction(MethodModel _method, ByteCode _byteCode, int _pc) { method = _method; pc = _pc; byteCode = _byteCode; } - protected Instruction(MethodModel _method, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + Instruction(MethodModel _method, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { this(_method, _byteCode, _wide ? _byteReader.getOffset() - 2 : _byteReader.getOffset() - 1); } @@ -187,7 +187,7 @@ public int getStackConsumeCount() { return (byteCode.getPop().getStackAdjust()); } - public int getStackProduceCount() { + int getStackProduceCount() { return (byteCode.getPush().getStackAdjust()); } @@ -247,7 +247,7 @@ public Instruction getPrevPC() { return (prevPC); } - public void setParentExpr(Instruction _parentExpr) { + private void setParentExpr(Instruction _parentExpr) { parentExpr = _parentExpr; } @@ -259,7 +259,7 @@ public Instruction getRootExpr() { return (parentExpr == null ? this : parentExpr.getRootExpr()); } - public boolean isReverseConditionalBranchTarget() { + private boolean isReverseConditionalBranchTarget() { return ((reverseConditionalBranchTargets != null) && (reverseConditionalBranchTargets.size() > 0)); } @@ -267,7 +267,7 @@ public boolean isForwardConditionalBranchTarget() { return ((forwardConditionalBranchTargets != null) && (forwardConditionalBranchTargets.size() > 0)); } - public boolean isReverseUnconditionalBranchTarget() { + private boolean isReverseUnconditionalBranchTarget() { return ((reverseUnconditionalBranchTargets != null) && (reverseUnconditionalBranchTargets.size() > 0)); } @@ -275,7 +275,7 @@ public boolean isForwardUnconditionalBranchTarget() { return ((forwardUnconditionalBranchTargets != null) && (forwardUnconditionalBranchTargets.size() > 0)); } - public boolean isReverseBranchTarget() { + private boolean isReverseBranchTarget() { return (isReverseConditionalBranchTarget() || isReverseUnconditionalBranchTarget()); } @@ -316,31 +316,31 @@ public void addBranchTarget(Branch _branch) { if (_branch.isReverse()) { if (_branch.isConditional()) { if (reverseConditionalBranchTargets == null) { - reverseConditionalBranchTargets = new LinkedList (); + reverseConditionalBranchTargets = new LinkedList<>(); } reverseConditionalBranchTargets.add((ConditionalBranch) _branch); } else { if (reverseUnconditionalBranchTargets == null) { - reverseUnconditionalBranchTargets = new LinkedList (); + reverseUnconditionalBranchTargets = new LinkedList<>(); } reverseUnconditionalBranchTargets.add(_branch); } } else { if (_branch.isConditional()) { if (forwardConditionalBranchTargets == null) { - forwardConditionalBranchTargets = new LinkedList (); + forwardConditionalBranchTargets = new LinkedList<>(); } forwardConditionalBranchTargets.add((ConditionalBranch) _branch); } else { if (forwardUnconditionalBranchTargets == null) { - forwardUnconditionalBranchTargets = new LinkedList (); + forwardUnconditionalBranchTargets = new LinkedList<>(); } forwardUnconditionalBranchTargets.add(_branch); } } } - public void removeBranchTarget(Branch _branch) { + void removeBranchTarget(Branch _branch) { if (_branch.isReverse()) { if (_branch.isConditional()) { if (reverseConditionalBranchTargets != null) { diff --git a/src/main/java/com/aparapi/internal/instruction/InstructionPattern.java b/src/main/java/com/aparapi/internal/instruction/InstructionPattern.java index c0ad9ac9..7b2ed535 100644 --- a/src/main/java/com/aparapi/internal/instruction/InstructionPattern.java +++ b/src/main/java/com/aparapi/internal/instruction/InstructionPattern.java @@ -68,7 +68,9 @@ to national security controls as identified on the Commerce Control List (curren public class InstructionPattern{ - @SuppressWarnings("unused") private boolean compareSubTrees(Instruction _lhs, Instruction _rhs) { + @SuppressWarnings("unused") private static boolean compareSubTrees(Instruction _lhs, Instruction _rhs) { + if (_lhs == _rhs) + return true; _lhs = _lhs.getReal(); _rhs = _rhs.getReal(); boolean same = _lhs.sameAs(_rhs); @@ -91,15 +93,15 @@ public class InstructionPattern{ public static class InstructionMatch{ public final boolean ok; - public static final InstructionMatch TRUE = new InstructionMatch(true); + static final InstructionMatch TRUE = new InstructionMatch(true); - public static final InstructionMatch FALSE = new InstructionMatch(false); + static final InstructionMatch FALSE = new InstructionMatch(false); - public InstructionMatch(boolean _ok) { + InstructionMatch(boolean _ok) { ok = _ok; } - public static InstructionMatch test(boolean _condition) { + static InstructionMatch test(boolean _condition) { return (_condition ? TRUE : FALSE); } } @@ -120,7 +122,7 @@ public InstructionMatch matches(Instruction _instruction, InstructionMatcher _in return (InstructionMatch.FALSE); } - public InstructionMatcher(String _description) { + InstructionMatcher(String _description) { description = _description; } @@ -129,7 +131,7 @@ public String getDescription() { } } - public class AssignableInstructionMatcher extends InstructionMatcher{ + static class AssignableInstructionMatcher extends InstructionMatcher{ private final Class>[] classes; public AssignableInstructionMatcher(Class>... _classes) { diff --git a/src/main/java/com/aparapi/internal/instruction/InstructionSet.java b/src/main/java/com/aparapi/internal/instruction/InstructionSet.java index b16536ab..d9b86562 100644 --- a/src/main/java/com/aparapi/internal/instruction/InstructionSet.java +++ b/src/main/java/com/aparapi/internal/instruction/InstructionSet.java @@ -66,7 +66,7 @@ to national security controls as identified on the Commerce Control List (curren public class InstructionSet{ - public static enum LoadSpec { + public enum LoadSpec { NONE, // F, // Float D, // Double @@ -76,7 +76,7 @@ public static enum LoadSpec { O, // Object } - public static enum StoreSpec { + public enum StoreSpec { NONE, // F, // Float D, // Double @@ -86,7 +86,7 @@ public static enum StoreSpec { O, // Object } - public static enum TypeSpec { + public enum TypeSpec { NONE("none", "none", 0, 0), // Z("Z", "boolean", 4, 1), // Note 'Z' is the java code for 'boolean' type C("C", "char", 2, 1), // @@ -114,7 +114,7 @@ public static enum TypeSpec { private final int slots; - private TypeSpec(String _shortName, String _longName, int _size, int _slots) { + TypeSpec(String _shortName, String _longName, int _size, int _slots) { shortName = _shortName; longName = _longName; size = _size; @@ -145,7 +145,7 @@ public String getShortName() { * */ - public static enum Operator { + public enum Operator { NONE, LogicalOr(true, "||"), // LogicalAnd(true, "&&", LogicalOr), // @@ -198,19 +198,19 @@ public static enum Operator { private Operator compliment; - private Operator(boolean _binary, String _text) { + Operator(boolean _binary, String _text) { text = _text; binary = _binary; } - private Operator(boolean _binary, String _text, Operator _c) { + Operator(boolean _binary, String _text, Operator _c) { this(_binary, _text); compliment = _c; compliment.compliment = this; } - private Operator() { + Operator() { this(false, null); } @@ -226,7 +226,7 @@ public String getText(boolean _invert) { return (_invert ? compliment.getText() : getText()); } - public boolean isBinary() { + boolean isBinary() { return (binary); } @@ -237,7 +237,7 @@ public boolean isUnary() { } } - public static enum PushSpec { + public enum PushSpec { NONE, // UNKNOWN, // I(TypeSpec.I), // @@ -256,7 +256,7 @@ public static enum PushSpec { LorD(TypeSpec.LorD), // RA(TypeSpec.RA); - private PushSpec(TypeSpec... _types) { + PushSpec(TypeSpec... _types) { types = _types; } @@ -267,7 +267,7 @@ public int getStackAdjust() { } } - public static enum PopSpec { + public enum PopSpec { NONE, // UNKNOWN(TypeSpec.UNKNOWN), // I(TypeSpec.I), // @@ -299,7 +299,7 @@ public static enum PopSpec { OARGS(TypeSpec.O, TypeSpec.ARGS), // ; - private PopSpec(TypeSpec... _types) { + PopSpec(TypeSpec... _types) { types = _types; } @@ -310,7 +310,7 @@ public int getStackAdjust() { } } - public static enum ImmediateSpec { + public enum ImmediateSpec { NONE("NONE"), // UNKNOWN("UNKNOWN"), // Bconst("byte constant value", TypeSpec.B), // @@ -329,7 +329,7 @@ public static enum ImmediateSpec { private final String name; - private ImmediateSpec(String _name, TypeSpec... _types) { + ImmediateSpec(String _name, TypeSpec... _types) { name = _name; types = _types; @@ -346,7 +346,7 @@ public TypeSpec[] getTypes() { } } - public static enum ByteCode { + public enum ByteCode { // name, operation type, immediateOperands, pop operands, push operands NOP(null, LoadSpec.NONE, StoreSpec.NONE, ImmediateSpec.NONE, PopSpec.NONE, PushSpec.NONE, Operator.NONE), // ACONST_NULL(I_ACONST_NULL.class, PushSpec.N), // @@ -640,14 +640,14 @@ public static enum ByteCode { private final Operator operator; - private LoadSpec loadSpec; + private final LoadSpec loadSpec; - private StoreSpec storeSpec; + private final StoreSpec storeSpec; private Constructor> constructor; - private ByteCode(Class> _class, LoadSpec _loadSpec, StoreSpec _storeSpec, ImmediateSpec _immediate, PopSpec _pop, - PushSpec _push, Operator _operator) { + ByteCode(Class> _class, LoadSpec _loadSpec, StoreSpec _storeSpec, ImmediateSpec _immediate, PopSpec _pop, + PushSpec _push, Operator _operator) { clazz = _class; immediate = _immediate; push = _push; @@ -660,68 +660,62 @@ private ByteCode(Class> _class, LoadSpec _loadSpec, StoreSpec _storeSpec, Imme try { constructor = clazz.getDeclaredConstructor(MethodModel.class, ByteReader.class, boolean.class); - } catch (final SecurityException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final NoSuchMethodException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final IllegalArgumentException e) { + } catch (final SecurityException | IllegalArgumentException | NoSuchMethodException e) { // TODO Auto-generated catch block e.printStackTrace(); } } } - private ByteCode(Class> _class, ImmediateSpec _immediate) { + ByteCode(Class> _class, ImmediateSpec _immediate) { this(_class, LoadSpec.NONE, StoreSpec.NONE, _immediate, PopSpec.NONE, PushSpec.NONE, Operator.NONE); } - private ByteCode(Class> _class, PushSpec _push) { + ByteCode(Class> _class, PushSpec _push) { this(_class, LoadSpec.NONE, StoreSpec.NONE, ImmediateSpec.NONE, PopSpec.NONE, _push, Operator.NONE); } - private ByteCode(Class> _class, StoreSpec _store, ImmediateSpec _immediate, PopSpec _pop) { + ByteCode(Class> _class, StoreSpec _store, ImmediateSpec _immediate, PopSpec _pop) { this(_class, LoadSpec.NONE, _store, _immediate, _pop, PushSpec.NONE, Operator.NONE); } - private ByteCode(Class> _class, StoreSpec _store, PopSpec _pop) { + ByteCode(Class> _class, StoreSpec _store, PopSpec _pop) { this(_class, LoadSpec.NONE, _store, ImmediateSpec.NONE, _pop, PushSpec.NONE, Operator.NONE); } - private ByteCode(Class> _class, ImmediateSpec _immediate, PopSpec _pop) { + ByteCode(Class> _class, ImmediateSpec _immediate, PopSpec _pop) { this(_class, LoadSpec.NONE, StoreSpec.NONE, _immediate, _pop, PushSpec.NONE, Operator.NONE); } - private ByteCode(Class> _class, ImmediateSpec _immediate, PopSpec _pop, Operator _operator) { + ByteCode(Class> _class, ImmediateSpec _immediate, PopSpec _pop, Operator _operator) { this(_class, LoadSpec.NONE, StoreSpec.NONE, _immediate, _pop, PushSpec.NONE, _operator); } - private ByteCode(Class> _class, LoadSpec _load, ImmediateSpec _immediate, PushSpec _push) { + ByteCode(Class> _class, LoadSpec _load, ImmediateSpec _immediate, PushSpec _push) { this(_class, _load, StoreSpec.NONE, _immediate, PopSpec.NONE, _push, Operator.NONE); } - private ByteCode(Class> _class, LoadSpec _load, PushSpec _push) { + ByteCode(Class> _class, LoadSpec _load, PushSpec _push) { this(_class, _load, StoreSpec.NONE, ImmediateSpec.NONE, PopSpec.NONE, _push, Operator.NONE); } - private ByteCode(Class> _class, ImmediateSpec _immediate, PushSpec _push) { + ByteCode(Class> _class, ImmediateSpec _immediate, PushSpec _push) { this(_class, LoadSpec.NONE, StoreSpec.NONE, _immediate, PopSpec.NONE, _push, Operator.NONE); } - private ByteCode(Class> _class, PopSpec _pop, PushSpec _push) { + ByteCode(Class> _class, PopSpec _pop, PushSpec _push) { this(_class, LoadSpec.NONE, StoreSpec.NONE, ImmediateSpec.NONE, _pop, _push, Operator.NONE); } - private ByteCode(Class> _class, PopSpec _pop, PushSpec _push, Operator _operator) { + ByteCode(Class> _class, PopSpec _pop, PushSpec _push, Operator _operator) { this(_class, LoadSpec.NONE, StoreSpec.NONE, ImmediateSpec.NONE, _pop, _push, _operator); } - private ByteCode(Class> _class, PopSpec _pop) { + ByteCode(Class> _class, PopSpec _pop) { this(_class, LoadSpec.NONE, StoreSpec.NONE, ImmediateSpec.NONE, _pop, PushSpec.NONE, Operator.NONE); } - private ByteCode() { + ByteCode() { this(null, LoadSpec.NONE, StoreSpec.NONE, ImmediateSpec.NONE, PopSpec.NONE, PushSpec.NONE, Operator.NONE); } @@ -737,7 +731,7 @@ public ImmediateSpec getImmediate() { return (immediate); } - public static ByteCode get(int _idx) { + static ByteCode get(int _idx) { return (values()[_idx]); } @@ -754,32 +748,17 @@ public boolean usesDouble() { final PushSpec push = getPush(); final PopSpec pop = getPop(); - if ((push == PushSpec.D) || (pop == PopSpec.D) || (pop == PopSpec.DD) || (pop == PopSpec.AID)) { - return true; - } + return (push == PushSpec.D) || (pop == PopSpec.D) || (pop == PopSpec.DD) || (pop == PopSpec.AID); - return false; } - public Instruction newInstruction(MethodModel _methodModel, ByteReader byteReader, boolean _isWide) { + Instruction newInstruction(MethodModel _methodModel, ByteReader byteReader, boolean _isWide) { Instruction newInstruction = null; if (constructor != null) { try { newInstruction = (Instruction) constructor.newInstance(_methodModel, byteReader, _isWide); newInstruction.setLength(byteReader.getOffset() - newInstruction.getThisPC()); - } catch (final SecurityException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final IllegalArgumentException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final InstantiationException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final IllegalAccessException e) { - // TODO Auto-generated catch block - e.printStackTrace(); - } catch (final InvocationTargetException e) { + } catch (final SecurityException | InvocationTargetException | IllegalAccessException | InstantiationException | IllegalArgumentException e) { // TODO Auto-generated catch block e.printStackTrace(); } @@ -804,7 +783,7 @@ public static Instruction create(MethodModel _methodModel, ByteReader _byteReade return (newInstruction); } - public Operator getOperator() { + Operator getOperator() { return (operator); } @@ -819,10 +798,10 @@ public StoreSpec getStore() { public static class CompositeInstruction extends Instruction{ - protected BranchSet branchSet; + final BranchSet branchSet; - public CompositeInstruction(MethodModel method, ByteCode _byteCode, Instruction _firstChild, Instruction _lastChild, - BranchSet _branchSet) { + CompositeInstruction(MethodModel method, ByteCode _byteCode, Instruction _firstChild, Instruction _lastChild, + BranchSet _branchSet) { super(method, _byteCode, -1); branchSet = _branchSet; setChildren(_firstChild, _lastChild); @@ -879,60 +858,60 @@ public BranchSet getBranchSet() { } public static class CompositeIfInstruction extends CompositeInstruction{ - public CompositeIfInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { + CompositeIfInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_IF, _firstChild, _lastChild, _branchSet); } } public static class CompositeIfElseInstruction extends CompositeInstruction{ - public CompositeIfElseInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { + CompositeIfElseInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_IF_ELSE, _firstChild, _lastChild, _branchSet); } } public static class CompositeForSunInstruction extends CompositeInstruction{ - public CompositeForSunInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { + CompositeForSunInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_FOR_SUN, _firstChild, _lastChild, _branchSet); } } public static class CompositeWhileInstruction extends CompositeInstruction{ - public CompositeWhileInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { + CompositeWhileInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_WHILE, _firstChild, _lastChild, _branchSet); } } public static class CompositeEmptyLoopInstruction extends CompositeInstruction{ - public CompositeEmptyLoopInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { + CompositeEmptyLoopInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_EMPTY_LOOP, _firstChild, _lastChild, _branchSet); } } public static class CompositeDoWhileInstruction extends CompositeInstruction{ - protected CompositeDoWhileInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, - BranchSet _branchSet) { + CompositeDoWhileInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, + BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_DO_WHILE, _firstChild, _lastChild, _branchSet); } } public static class CompositeForEclipseInstruction extends CompositeInstruction{ - protected CompositeForEclipseInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, - BranchSet _branchSet) { + CompositeForEclipseInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, + BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_FOR_ECLIPSE, _firstChild, _lastChild, _branchSet); } } public static class CompositeArbitraryScopeInstruction extends CompositeInstruction{ - protected CompositeArbitraryScopeInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, - BranchSet _branchSet) { + CompositeArbitraryScopeInstruction(MethodModel method, Instruction _firstChild, Instruction _lastChild, + BranchSet _branchSet) { super(method, ByteCode.COMPOSITE_ARBITRARY_SCOPE, _firstChild, _lastChild, _branchSet); } } public static abstract class OperatorInstruction extends Instruction{ - protected OperatorInstruction(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { + OperatorInstruction(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { super(_methodPoolEntry, code, reader, _wide); } @@ -950,7 +929,7 @@ public static abstract class BinaryOperator extends OperatorInstruction implemen return (getLastChild()); } - protected BinaryOperator(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { + BinaryOperator(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { super(_methodPoolEntry, code, reader, _wide); } } @@ -960,23 +939,23 @@ public static abstract class UnaryOperator extends OperatorInstruction implement return (getFirstChild()); } - protected UnaryOperator(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { + UnaryOperator(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { super(_methodPoolEntry, code, reader, _wide); } } public static abstract class CastOperator extends UnaryOperator{ - protected CastOperator(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { + CastOperator(MethodModel _methodPoolEntry, ByteCode code, ByteReader reader, boolean _wide) { super(_methodPoolEntry, code, reader, _wide); } } public static abstract class Branch extends Instruction{ - protected int offset; + int offset; - protected boolean breakOrContinue; + boolean breakOrContinue; - protected Instruction target; + Instruction target; public int getAbsolute() { return (getThisPC() + getOffset()); @@ -986,11 +965,11 @@ private int getOffset() { return (offset); } - public Branch(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + Branch(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } - public Branch(MethodModel _methodPoolEntry, ByteCode _byteCode, Instruction _target) { + Branch(MethodModel _methodPoolEntry, ByteCode _byteCode, Instruction _target) { super(_methodPoolEntry, _byteCode, -1); setTarget(_target); } @@ -1017,7 +996,7 @@ public boolean isReverseConditional() { return (isConditional() && isReverse()); } - public boolean isForwardConditional() { + boolean isForwardConditional() { return (isConditional() && isForward()); } @@ -1061,7 +1040,7 @@ public void retarget(Instruction _newTarget) { public static abstract class ConditionalBranch extends Branch{ private BranchSet branchSet; - public ConditionalBranch(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + ConditionalBranch(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } @@ -1118,17 +1097,17 @@ public ConditionalBranch findEndOfConditionalBranchSet(Instruction _extent) { } public static abstract class UnconditionalBranch extends Branch{ - public UnconditionalBranch(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + UnconditionalBranch(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } - public UnconditionalBranch(MethodModel _methodPoolEntry, ByteCode _byteCode, Instruction _target) { + UnconditionalBranch(MethodModel _methodPoolEntry, ByteCode _byteCode, Instruction _target) { super(_methodPoolEntry, _byteCode, _target); } } public static abstract class IfUnary extends ConditionalBranch16 implements Unary{ - public IfUnary(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + IfUnary(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } @@ -1138,7 +1117,7 @@ public IfUnary(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byt } public static abstract class If extends ConditionalBranch16 implements Binary{ - public If(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + If(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } @@ -1152,7 +1131,7 @@ public If(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteRead } public static abstract class ConditionalBranch16 extends ConditionalBranch implements HasOperator{ - public ConditionalBranch16(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + ConditionalBranch16(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); offset = _byteReader.s2(); } @@ -1162,22 +1141,22 @@ public ConditionalBranch16(MethodModel _methodPoolEntry, ByteCode _byteCode, Byt } } - public static abstract class UnconditionalBranch16 extends UnconditionalBranch{ - public UnconditionalBranch16(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + static abstract class UnconditionalBranch16 extends UnconditionalBranch{ + UnconditionalBranch16(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); offset = _byteReader.s2(); } } - public static abstract class Branch32 extends UnconditionalBranch{ - public Branch32(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + static abstract class Branch32 extends UnconditionalBranch{ + Branch32(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); offset = _byteReader.s4(); } } public static abstract class ArrayAccess extends Instruction{ - public ArrayAccess(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + ArrayAccess(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } @@ -1191,13 +1170,13 @@ public Instruction getArrayIndex() { } public static abstract class AccessArrayElement extends ArrayAccess{ - protected AccessArrayElement(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + AccessArrayElement(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } } public static class I_AALOAD extends AccessArrayElement{ - public I_AALOAD(MethodModel _methodPoolEntry, ByteReader _byteReader, boolean _wide) { + I_AALOAD(MethodModel _methodPoolEntry, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, ByteCode.AALOAD, _byteReader, _wide); } @@ -1211,13 +1190,13 @@ public Instruction getValue() { return (getFirstChild().getNextExpr().getNextExpr()); } - protected AssignToArrayElement(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { + AssignToArrayElement(MethodModel _methodPoolEntry, ByteCode _byteCode, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, _byteCode, _byteReader, _wide); } } - public static class I_AASTORE extends AssignToArrayElement{ - public I_AASTORE(MethodModel _methodPoolEntry, ByteReader _byteReader, boolean _wide) { + static class I_AASTORE extends AssignToArrayElement{ + I_AASTORE(MethodModel _methodPoolEntry, ByteReader _byteReader, boolean _wide) { super(_methodPoolEntry, ByteCode.AASTORE, _byteReader, _wide); } @@ -1227,7 +1206,7 @@ public I_AASTORE(MethodModel _methodPoolEntry, ByteReader _byteReader, boolean _ } public static class I_ACONST_NULL extends Instruction implements Constant