Skip to content

GitLab

Notice to GitKraken users: A vulnerability has been found in the SSH key generation of GitKraken versions 7.6.0 to 8.0.0 (https://www.gitkraken.com/blog/weak-ssh-key-fix). If you use GitKraken and have generated a SSH key using one of these versions, please remove it both from your local workstation and from your LRZ GitLab profile.

21.10.2021, 9:00 - 11:00: Due to updates GitLab may be unavailable for some minutes between 09:00 and 11:00.

Commit e0ffa528 authored by Benedikt Zoennchen's avatar Benedikt Zoennchen
Browse files

add native java support if open cl is missing, add alert message for the gui and console as well.

parent b030e6a4
Hide whitespace changes
Inline Side-by-side
VadereGui/resources/messages.properties
View file @ e0ffa528
......@@ -201,6 +201,9 @@ AdjustPanel.lblVelocity.text=FPS
AdjustPanel.lblTime=Time
AdjustPanel.lblStep.text=Timestep
ProjectView.warning.lwjgl.title=Warning
ProjectView.warning.opencl.title=Warning
ProjectView.warning.opencl.text=OpenCL acceleration disabled.\nOpenCL is not supported on your machine.\nDriver updates might resolve the issue.
ProjectView.mntmOutputToSceneario.text=Generate scenario
ProjectView.mntmRunOutput.text=Run offline again
ProjectView.OpenInExplorer.text=Open in Explorer
......
VadereGui/resources/messages_de_DE.properties
View file @ e0ffa528
......@@ -63,6 +63,9 @@ ProjectView.JSONSwitch.link=<u>Wechsle zum <b>{0}</b> Modus</u>
ProjectView.JSONDisplay.label=Keine Auswahl
ProjectView.btnAbout.tooltip=\u00dcber
ProjectView.about.title=\u00dcber
ProjectView.warning.lwjgl.title=Warnung
ProjectView.warning.opencl.title=Warnung
ProjectView.warning.opencl.text=OpenCL-Beschleunigung deaktiviert.\nEs konnte keine OpenCL Unterst\u00FCtzung gefunden werden.\nTreiber-Aktualisierung kann das Problem beheben.
ProjectView.title=Vadere
ProjectView.version= Version {0}
ProjectView.license.text=Dieses Software ist unter der {0} lizenziert
......
VadereGui/src/org/vadere/gui/projectview/VadereApplication.java
View file @ e0ffa528
......@@ -6,12 +6,15 @@ import org.vadere.gui.projectview.view.ProjectView;
import org.vadere.util.io.IOUtils;
import org.vadere.util.logging.Logger;
import org.vadere.util.logging.StdOutErrLog;
import org.vadere.util.opencl.CLUtils;
import java.io.IOException;
import java.util.prefs.BackingStoreException;
import java.util.prefs.InvalidPreferencesFormatException;
import java.util.prefs.Preferences;
import javax.swing.*;
/**
* Entry point for the Vadere GUI.
*
......
VadereGui/src/org/vadere/gui/projectview/view/ProjectView.java
View file @ e0ffa528
package org.vadere.gui.projectview.view;
import org.jetbrains.annotations.NotNull;
import org.vadere.gui.components.utils.Messages;
import org.vadere.gui.postvisualization.control.Player;
import org.vadere.gui.projectview.VadereApplication;
......@@ -44,6 +45,7 @@ import org.vadere.simulator.projects.SingleScenarioFinishedListener;
import org.vadere.simulator.projects.VadereProject;
import org.vadere.util.io.IOUtils;
import org.vadere.util.logging.Logger;
import org.vadere.util.opencl.CLUtils;
import java.awt.*;
import java.awt.event.ActionEvent;
......@@ -256,9 +258,26 @@ public class ProjectView extends JFrame implements ProjectFinishedListener, Sing
frame.setSize(1200, 800);
frame.openLastUsedProject(model);
checkDependencies(frame);
});
}
private static void checkDependencies(@NotNull final JFrame frame) {
try {
if(!CLUtils.isOpenCLSupported()) {
JOptionPane.showMessageDialog(frame,
Messages.getString("ProjectView.warning.opencl.text"),
Messages.getString("ProjectView.warning.opencl.title"),
JOptionPane.WARNING_MESSAGE);
}
} catch (UnsatisfiedLinkError linkError) {
JOptionPane.showMessageDialog(frame,
"[LWJGL]: " + linkError.getMessage(),
Messages.getString("ProjectView.warning.lwjgl.title"),
JOptionPane.WARNING_MESSAGE);
}
}
private void openLastUsedProject(final ProjectViewModel model) {
String lastUsedProjectPath =
Preferences.userNodeForPackage(VadereApplication.class).get("last_used_project", null);
......
VadereSimulator/src/org/vadere/simulator/entrypoints/cmd/VadereConsole.java
View file @ e0ffa528
......@@ -18,6 +18,9 @@ import org.vadere.simulator.entrypoints.cmd.commands.SuqSubCommand;
import org.vadere.simulator.utils.scenariochecker.ScenarioChecker;
import org.vadere.util.logging.Logger;
import org.vadere.util.logging.StdOutErrLog;
import org.vadere.util.opencl.CLUtils;
import javax.swing.*;
/**
......@@ -35,14 +38,20 @@ public class VadereConsole {
try {
StdOutErrLog.addStdOutErrToLog();
if (!CLUtils.isOpenCLSupported()) {
System.out.println("Warning: OpenCL acceleration disabled, since no OpenCL support could be found!");
}
Namespace ns = parser.parseArgs(args);
SubCommandRunner sRunner = ns.get("func");
sRunner.run(ns, parser);
} catch (UnsatisfiedLinkError linkError) {
System.err.println("[LWJGL]: " + linkError.getMessage());
} catch (ArgumentParserException e) {
parser.handleError(e);
System.exit(1);
} catch (Exception e) {
logger.error("topographyError in command:" + e.getMessage());
System.err.println("topographyError in command:" + e.getMessage());
e.printStackTrace();
System.exit(1);
}
......
VadereSimulator/src/org/vadere/simulator/models/density/CLGaussianFilter.java
View file @ e0ffa528
......@@ -5,6 +5,7 @@ import java.io.IOException;
import java.util.function.BiFunction;
import org.vadere.util.logging.Logger;
import org.vadere.util.math.Convolution;
import org.vadere.util.opencl.CLConvolution;
import org.vadere.util.opencl.OpenCLException;
......@@ -14,7 +15,7 @@ class CLGaussianFilter extends GaussianFilter {
private Logger logger = Logger.getLogger(CLConvolution.class);
CLGaussianFilter(final Rectangle2D scenarioBounds, final double scale, final BiFunction<Integer, Integer, Float> f,
final boolean normalize) throws IOException, OpenCLException {
final boolean normalize) throws OpenCLException {
super(scenarioBounds, scale, f, normalize);
this.convolution = new CLConvolution(matrixWidth, matrixHeight, kernelWidth, kernel);
this.convolution.init();
......@@ -22,8 +23,12 @@ class CLGaussianFilter extends GaussianFilter {
@Override
public void filterImage() {
try {
outputMatrix = this.convolution.convolve(inputMatrix);
long ms = System.currentTimeMillis();
outputMatrix = convolution.convolve(inputMatrix);
ms = System.currentTimeMillis() - ms;
logger.infof("filtering required " + ms + "[ms]");
} catch (OpenCLException e) {
logger.error(e.getMessage());
e.printStackTrace();
......
VadereSimulator/src/org/vadere/simulator/models/density/IGaussianFilter.java
View file @ e0ffa528
package org.vadere.simulator.models.density;
import org.jetbrains.annotations.NotNull;
import org.vadere.simulator.models.potential.timeCostFunction.loading.IPedestrianLoadingStrategy;
import org.vadere.state.attributes.scenario.AttributesAgent;
import org.vadere.state.scenario.Agent;
import org.vadere.state.scenario.Topography;
import org.vadere.util.logging.Logger;
import org.vadere.util.opencl.CLUtils;
import org.vadere.util.opencl.OpenCLException;
import java.awt.geom.Rectangle2D;
......@@ -83,11 +85,11 @@ public interface IGaussianFilter {
*/
void destroy();
static <E extends Agent> IGaussianFilter create(final Rectangle2D scenarioBounds,
Collection<E> pedestrians, final double scale,
static <E extends Agent> IGaussianFilter create(@NotNull final Rectangle2D scenarioBounds,
@NotNull Collection<E> pedestrians, final double scale,
final double standardDerivation,
final AttributesAgent attributesPedestrian,
final IPedestrianLoadingStrategy loadingStrategy) {
@NotNull final AttributesAgent attributesPedestrian,
@NotNull final IPedestrianLoadingStrategy loadingStrategy) {
return create(scenarioBounds, pedestrians, scale, standardDerivation, attributesPedestrian, loadingStrategy,
Type.OpenCL);
}
......@@ -95,11 +97,11 @@ public interface IGaussianFilter {
/*
* Factory-methods
*/
static <E extends Agent> IGaussianFilter create(final Rectangle2D scenarioBounds,
Collection<E> pedestrians, final double scale,
static <E extends Agent> IGaussianFilter create(@NotNull final Rectangle2D scenarioBounds,
@NotNull Collection<E> pedestrians, final double scale,
final double standardDerivation,
final AttributesAgent attributesPedestrian,
final IPedestrianLoadingStrategy loadingStrategy, final Type type) {
@NotNull final AttributesAgent attributesPedestrian,
@NotNull final IPedestrianLoadingStrategy loadingStrategy, final Type type) {
double scaleFactor = attributesPedestrian.getRadius() * 2
* attributesPedestrian.getRadius() * 2
......@@ -111,52 +113,71 @@ public interface IGaussianFilter {
(centerI, i) -> (float) (Math.sqrt(scaleFactor) * Math.exp(-((centerI - i) / scale)
* ((centerI - i) / scale) / (2 * standardDerivation * standardDerivation)));
IGaussianFilter clFilter;
switch (type) {
case OpenCL: {
try {
clFilter = new CLGaussianFilter(scenarioBounds, scale, f, false);
} catch (IOException | OpenCLException e) {
e.printStackTrace();
logger.warn(e.getClass().getName() + " while initializing OpenCL: " + e.getMessage() + " OpenCL can not be used. Native Java implementation will be used which might cause performance issues.");
clFilter = new JGaussianFilter(scenarioBounds, scale, f, false);
}
} break;
default:
clFilter = new JGaussianFilter(scenarioBounds, scale, f, false);
}
IGaussianFilter clFilter = getFilter(scenarioBounds, f, type, scale);
return new PedestrianGaussianFilter(pedestrians, clFilter, loadingStrategy);
}
static IGaussianFilter create(
final Topography scenario, final double scale,
final boolean scenarioHasBoundary, final double standardDerivation) {
/**
* Returns the desired filter if possible. If the filter <tt>type</tt> is not supported
* it returns a default java implementation of the filter which should on any device / platform
* / machine.
*
* @param scenarioBounds the bound the filter is working on
* @param f the function for generating the kernel values
* @param type the type of the filter e.g. OpenCL (GPU filter) or native java filter
* @param scale the scale of the filtered grid which directly maps to the size of the matrices which take part in the convolution
*
* @return the desired filter or (if it is not supported) a native java implementation
*/
static IGaussianFilter getFilter(@NotNull final Rectangle2D scenarioBounds,
@NotNull final BiFunction<Integer, Integer, Float> f,
@NotNull final Type type,
final double scale) {
IGaussianFilter clFilter;
switch (type) {
case OpenCL: {
if(CLUtils.isOpenCLSupported()) {
try {
clFilter = new CLGaussianFilter(scenarioBounds, scale, f, false);
} catch (OpenCLException e) {
e.printStackTrace();
logger.warn("Error while initializing OpenCL: " + e.getMessage());
clFilter = new JGaussianFilter(scenarioBounds, scale, f, false);
}
catch (UnsatisfiedLinkError linkError) {
linkError.printStackTrace();
logger.warn("Linking error (native lib problem) while initializing OpenCL: " + linkError.getMessage());
clFilter = new JGaussianFilter(scenarioBounds, scale, f, false);
}
}
else {
clFilter = new JGaussianFilter(scenarioBounds, scale, f, false);
}
} break;
default:
clFilter = new JGaussianFilter(scenarioBounds, scale, f, false);
}
return clFilter;
}
static IGaussianFilter create(@NotNull final Topography scenario,
final double scale,
final boolean scenarioHasBoundary,
final double standardDerivation) {
return create(scenario, scale, scenarioHasBoundary, standardDerivation, Type.OpenCL);
}
static IGaussianFilter create(
final Topography scenario, final double scale,
final boolean scenarioHasBoundary, final double standardDerivation, final Type type) {
static IGaussianFilter create(final Topography scenario,
final double scale,
final boolean scenarioHasBoundary,
final double standardDerivation,
final Type type) {
double varianz = standardDerivation * standardDerivation;
BiFunction<Integer, Integer, Float> f = (centerI, i) -> (float) ((1.0 / (2 * Math.PI * varianz))
* Math.exp(-((centerI - i) / scale) * ((centerI - i) / scale) / (2 * varianz)));
IGaussianFilter clFilter;
switch (type) {
case OpenCL: {
try {
clFilter = new CLGaussianFilter(scenario.getBounds(), scale, f, true);
} catch (IOException | OpenCLException e) {
e.printStackTrace();
logger.warn(e.getClass().getName() + " while initializing OpenCL: " + e.getMessage() + " OpenCL can not be used. Native Java implementation will be used which might cause performance issues.");
clFilter = new JGaussianFilter(scenario.getBounds(), scale, f, true);
}
} break;
default:
clFilter = new JGaussianFilter(scenario.getBounds(), scale, f, true);
}
IGaussianFilter clFilter = getFilter(scenario.getBounds(), f, type, scale);
return new ObstacleGaussianFilter(scenario, clFilter);
}
......
VadereSimulator/src/org/vadere/simulator/models/density/JGaussianFilter.java
View file @ e0ffa528
......@@ -14,7 +14,10 @@ public class JGaussianFilter extends GaussianFilter {
@Override
public void filterImage() {
long ms = System.currentTimeMillis();
outputMatrix = Convolution.convolveSeperate(inputMatrix, kernel, kernel, matrixWidth, matrixHeight, kernelWidth);
ms = System.currentTimeMillis() - ms;
IGaussianFilter.logger.infof("filtering required " + ms + "[ms]");
}
@Override
......
VadereUtils/src/org/vadere/util/opencl/CLConvolution.java
View file @ e0ffa528
......@@ -49,18 +49,16 @@ import static org.lwjgl.system.MemoryUtil.NULL;
import static org.lwjgl.system.MemoryUtil.memUTF8;
/**
* This class offers the convolution operation via OpenCL i.e. the use of the GPU to
* accelerate the computation. Successive convolutions can be computed by reusing memory if the size of
* the involved matrices do not change, that is, {@link CLConvolution#clearCL()} can be called
* after multiple convolutions are computed.
*
* @author Benedikt Zoennchen
*/
public class CLConvolution {
public class CLConvolution extends CLOperation {
private static Logger log = Logger.getLogger(CLConvolution.class);
// CL ids
private long clPlatform;
private long clDevice;
private long clContext;
private long clQueue;
private long clProgram;
// CL Memory
private long clInput;
private long clOutput;
......@@ -106,11 +104,24 @@ public class CLConvolution {
this(KernelType.Separate, matrixWidth, matrixHeight, kernelWidth, kernel);
}
/**
* Default constructor.
*
* @param type kernel type e.g. separated kernel
* @param matrixWidth input matrix width
* @param matrixHeight input matrix height
* @param kernelWidth kernel width which is also the kernel height
* @param kernel a 1D error representing the kernel (this can be a 1D kernel or a 2D kernel depending on the <tt>type</tt>)
*
* @throws OpenCLException if there is some OpenCL problem e.g. it is not supported
* @throws UnsatisfiedLinkError if native libraries for LWJGL are missing
*/
public CLConvolution(
@NotNull final KernelType type,
final int matrixWidth,
final int matrixHeight,
final int kernelWidth, @NotNull final float[] kernel) throws OpenCLException {
final int kernelWidth,
@NotNull final float[] kernel) throws OpenCLException, UnsatisfiedLinkError {
this.type = type;
this.matrixHeight = matrixHeight;
this.matrixWidth = matrixWidth;
......@@ -126,7 +137,7 @@ public class CLConvolution {
init();
}
public void init() throws OpenCLException {
public void init() throws OpenCLException, BootstrapMethodError {
initCallbacks();
initCL();
buildProgram();
......@@ -151,7 +162,18 @@ public class CLConvolution {
}
}
public float[] convolve(final float[] input) throws OpenCLException {
/**
* Executes the convolution operation this might be a 2D convolution realized by one 2D kernel
* or by two 1D kernels or just one 1D convolution which depends on the constructor arguments of this
* class.
*
* @param input the input matrix
*
* @return the output matrix of same dimension as the input matrix
*
* @throws OpenCLException if there is any problem with OpenCL e.g. no OpenCL support
*/
public float[] convolve(final float[] input) throws OpenCLException {
// 1. write input to native-c-like-memory
CLUtils.toFloatBuffer(input, hostScenario);
......@@ -273,64 +295,17 @@ public class CLConvolution {
}
}
public void clearCL() throws OpenCLException {
/**
* Works like a C++ destructor, i.e. frees host and GPU / device memory.
* This has to be called if no more convolutions are computed. Note that
* successive convolutions can be computed by reusing memory if the size of
* the involved matrices do not change.
*
* @throws OpenCLException if there is any problem with OpenCL e.g. no OpenCL support
*/
public void clearCL() throws OpenCLException {
clearMemory();
CLInfo.checkCLError(clReleaseCommandQueue(clQueue));
CLInfo.checkCLError(clReleaseProgram(clProgram));
CLInfo.checkCLError(clReleaseContext(clContext));
contextCB.free();
programCB.free();
}
// private helpers
private void initCallbacks() {
contextCB = CLContextCallback.create((errinfo, private_info, cb, user_data) ->
{
log.debug("[LWJGL] cl_context_callback" + "\tInfo: " + memUTF8(errinfo));
});
programCB = CLProgramCallback.create((program, user_data) ->
{
try {
log.debug("The cl_program [0x"+program+"] was built " + (CLInfo.getProgramBuildInfoInt(program, clDevice, CL_PROGRAM_BUILD_STATUS) == CL_SUCCESS ? "successfully" : "unsuccessfully"));
} catch (OpenCLException e) {
e.printStackTrace();
}
});
}
private void initCL() throws OpenCLException {
try (MemoryStack stack = stackPush()) {
IntBuffer errcode_ret = stack.callocInt(1);
IntBuffer numberOfPlatforms = stack.mallocInt(1);
CLInfo.checkCLError(clGetPlatformIDs(null, numberOfPlatforms));
PointerBuffer platformIDs = stack.mallocPointer(numberOfPlatforms.get(0));
CLInfo.checkCLError(clGetPlatformIDs(platformIDs, numberOfPlatforms));
clPlatform = platformIDs.get(0);
IntBuffer numberOfDevices = stack.mallocInt(1);
CLInfo.checkCLError(clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, null, numberOfDevices));
PointerBuffer deviceIDs = stack.mallocPointer(numberOfDevices.get(0));
CLInfo.checkCLError(clGetDeviceIDs(clPlatform, CL_DEVICE_TYPE_GPU, deviceIDs, numberOfDevices));
clDevice = deviceIDs.get(0);
log.debug("CL_DEVICE_NAME = " + CLInfo.getDeviceInfoStringUTF8(clDevice, CL_DEVICE_NAME));
PointerBuffer ctxProps = stack.mallocPointer(3);
ctxProps.put(CL_CONTEXT_PLATFORM)
.put(clPlatform)
.put(NULL)
.flip();
clContext = clCreateContext(ctxProps, clDevice, contextCB, NULL, errcode_ret);
CLInfo.checkCLError(errcode_ret);
clQueue = clCreateCommandQueue(clContext, clDevice, 0, errcode_ret);
CLInfo.checkCLError(errcode_ret);
}
super.clearCL();
}
private void buildProgram() throws OpenCLException {
......
VadereUtils/src/org/vadere/util/opencl/CLLinkedCell.java
View file @ e0ffa528
......@@ -73,16 +73,9 @@ import static org.lwjgl.system.MemoryUtil.memUTF8;
* This class offers the methods to compute an array based linked-cell which contains 2D-coordinates i.e. {@link VPoint}
* using the GPU (see. green-2007 Building the Grid using Sorting).
*/
public class CLLinkedCell {
public class CLLinkedCell extends CLOperation {
private static Logger log = Logger.getLogger(CLLinkedCell.class);
// CL ids
private long clPlatform;
private long clDevice;
private long clContext;
private long clQueue;
private long clProgram;
// CL Memory
private long clHashes;
private long clIndices;
......@@ -112,10 +105,6 @@ public class CLLinkedCell {
private ByteBuffer source;
private ByteBuffer particleSource;
// CL callbacks
private CLContextCallback contextCB;
private CLProgramCallback programCB;
// CL kernel
private long clBitonicSortLocal;
private long clBitonicSortLocal1;
......@@ -130,7 +119,6 @@ public class CLLinkedCell {
private float iCellSize;
private int[] iGridSize;
private List<VPoint> positionList;
private final int deviceType;
private int[] keys;
private int[] values;
......@@ -145,7 +133,6 @@ public class CLLinkedCell {
// time measurement
private boolean debug = false;
private boolean profiling = false;
private static int MIN_LOCAL_SIZE = 1;
......@@ -163,12 +150,12 @@ public class CLLinkedCell {
final VRectangle bound,
final double cellSize,
final int device) throws OpenCLException {
super(device);
this.numberOfElements = numberOfElements;
this.iGridSize = new int[]{ (int)Math.ceil(bound.getWidth() / cellSize), (int)Math.ceil(bound.getHeight() / cellSize)};
this.numberOfGridCells = this.iGridSize[0] * this.iGridSize[1];
this.bound = bound;
this.iCellSize = (float)cellSize;
this.deviceType = device;
if(debug) {
Configuration.DEBUG.set(true);
......@@ -478,7 +465,7 @@ public class CLLinkedCell {
}
else {
localWorkSize = maxWorkGroupSize;
globalWorkSize = multipleOf(numberOfElements, localWorkSize);
globalWorkSize = CLOperation.multipleOf(numberOfElements, localWorkSize);
}
clGlobalWorkSize.put(0, globalWorkSize);
......@@ -488,14 +475,6 @@ public class CLLinkedCell {
}
}
private long multipleOf(long value, long multiple) {
long result = multiple;
while (result < value) {
result += multiple;
}
return result;
}
public boolean checkMinSupportedLocalSize() throws OpenCLException {
try (MemoryStack stack = stackPush()) {
LongBuffer rBitonicMergeLocal = stack.mallocLong(1);
......@@ -615,30 +594,6 @@ public class CLLinkedCell {
}
}
private long enqueueNDRangeKernel(final String name, long command_queue, long kernel, int work_dim, PointerBuffer global_work_offset, PointerBuffer global_work_size, PointerBuffer local_work_size, PointerBuffer event_wait_list, PointerBuffer event) throws OpenCLException {
if(profiling) {
try (MemoryStack stack = stackPush()) {
PointerBuffer clEvent = stack.mallocPointer(1);
LongBuffer startTime = stack.mallocLong(1);
LongBuffer endTime = stack.mallocLong(1);
long result = clEnqueueNDRangeKernel(command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, event_wait_list, clEvent);
clWaitForEvents(clEvent);
long eventAddr = clEvent.get();
CLInfo.checkCLError(clGetEventProfilingInfo(eventAddr, CL_PROFILING_COMMAND_START, startTime, null));
CLInfo.checkCLError(clGetEventProfilingInfo(eventAddr, CL_PROFILING_COMMAND_END, endTime, null));
clEvent.clear();
// in nanaSec
log.info(name + " event time " + "0x"+eventAddr + ": " + ((double)endTime.get() - startTime.get()) / 1_000_000.0 + " [ms]");
endTime.clear();
startTime.clear();
return result;
}
}
else {
return clEnqueueNDRangeKernel(command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, event_wait_list, event);
}
}
private long getMaxWorkGroupSizeForKernel(long clDevice, long clKernel, long workItemMem) throws OpenCLException {
try (MemoryStack stack = stackPush()) {
LongBuffer pp = stack.mallocLong(1);
......@@ -701,75 +656,15 @@ public class CLLinkedCell {
}
}
private void clearCL() throws OpenCLException {
@Override
protected void clearCL() throws OpenCLException {
CLInfo.checkCLError(clReleaseKernel(clBitonicSortLocal));
CLInfo.checkCLError(clReleaseKernel(clBitonicSortLocal1));
CLInfo.checkCLError(clReleaseKernel(clBitonicMergeGlobal));
CLInfo.checkCLError(clReleaseKernel(clBitonicMergeLocal));