private void computeCL(boolean is64bit) {
double minX = transformX(-ww * 0.5) + offsetX;
double maxX = transformX(ww * 0.5) + offsetX;
double minY = transformY(-wh * 0.5) + offsetY;
double maxY = transformY(wh * 0.5) + offsetY;
double rangeX = maxX - minX;
double rangeY = maxY - minY;
kernel2DGlobalWorkSize.put(0, ww).put(1, wh);
// start computation
clSetKernelArg1i(clKernel, 0, ww);
clSetKernelArg1i(clKernel, 1, wh);
if (!is64bit || !isDoubleFPAvailable(deviceCaps)) {
clSetKernelArg1f(clKernel, 2, (float) minX);
clSetKernelArg1f(clKernel, 3, (float) minY);
clSetKernelArg1f(clKernel, 4, (float) rangeX);
clSetKernelArg1f(clKernel, 5, (float) rangeY);
} else {
clSetKernelArg1d(clKernel, 2, minX);
clSetKernelArg1d(clKernel, 3, minY);
clSetKernelArg1d(clKernel, 4, rangeX);
clSetKernelArg1d(clKernel, 5, rangeY);
}
// acquire GL objects, and enqueue a kernel with a probe from the list
int errcode = clEnqueueAcquireGLObjects(clQueue, clTexture, null, null);
checkCLError(errcode);
errcode =
clEnqueueNDRangeKernel(
clQueue, clKernel, 2, null, kernel2DGlobalWorkSize, null, null, null);
checkCLError(errcode);
errcode = clEnqueueReleaseGLObjects(clQueue, clTexture, null, !syncGLtoCL ? syncBuffer : null);
checkCLError(errcode);
if (!syncGLtoCL) {
clEvent = syncBuffer.get(0);
glFenceFromCLEvent = glCreateSyncFromCLeventARB(clContext, clEvent, 0);
}
// block until done (important: finish before doing further gl work)
if (syncGLtoCL) {
errcode = clFinish(clQueue);
checkCLError(errcode);
}
}
private void setKernelConstants() {
clSetKernelArg1p(clKernel, 6, clTexture);
clSetKernelArg1p(clKernel, 7, clColorMap);
clSetKernelArg1i(clKernel, 8, COLOR_MAP_SIZE);
clSetKernelArg1i(clKernel, 9, maxIterations);
}