FractalViewport Conductor is a hybrid demonstration that makes a live window and then splits its work into two coordinated engines: one engine draws, the other computes. The drawing engine is classic OpenGL running through a Windows WGL context. The compute engine is OpenCL running on a GPU, and the two engines share the same pixel buffer so the image can move from computation to display without detouring through the CPU. The program begins by creating a Win32 window and preparing an OpenGL context with double buffering. It then loads the minimal OpenGL buffer functions needed to manage a pixel buffer object and feed it directly into a texture. After the graphics setup is stable, the program searches for a GPU device that supports OpenCL and also supports OpenGL sharing. If such a device exists, it creates an OpenCL context that is explicitly tied to the active OpenGL context, allowing both systems to reference the same buffer.

Once initialization is complete, the program enters an event loop that alternates between handling user input and rendering frames. Each frame, the compute engine temporarily takes ownership of the shared pixel buffer, runs a kernel that evaluates the Mandelbrot set for every pixel, and writes a color result into the buffer. The kernel maps screen coordinates into a complex plane region controlled by a center point, a width scale, and an iteration detail level. Pixels that remain stable under repeated iteration are colored dark, while pixels that escape are colored with a smooth gradient based on how quickly they diverge. After computation, ownership of the buffer is returned to OpenGL. OpenGL then uploads the pixel buffer into a texture and draws a full screen quad textured with the result, producing a fluid fractal image.

Interaction turns the renderer into an exploratory instrument. Left click zooms into the fractal at the cursor location, right click resets the view, and keyboard input can exit cleanly. A small Zorro entry wrapper is included so the demo can be launched once from a Zorro run cycle while still behaving like an interactive graphical application.

Code
// mandelbrot_gl_cl_demo.cpp
// Win32 + WGL(OpenGL) display + OpenCL compute (CL/GL sharing)
//
// Build notes (MSVC):
//   link: OpenGL32.lib  OpenCL.lib  User32.lib  Gdi32.lib
//
// If you want this as a Zorro64 C++ strategy DLL, you can add:
//   #include <zorro.h>
// and change entry point to DLLFUNC int main() { ...; return quit("done"); }
// For now it's a plain Win32 EXE-style main(WINARGS).

#include <windows.h>
#include <stdio.h>
#include <math.h>
#include <stddef.h>
#include <zorro.h>

#include <CL/cl.h>
#include <CL/cl_gl.h>     // cl_khr_gl_sharing

#include <GL/gl.h>

// ------------------------- Globals -------------------------
static HWND   gHwnd = 0;
static HDC    gHdc  = 0;
static HGLRC  gHgl  = 0;

static int    gW = 640;
static int    gH = 480;

static double m_x     = 0.344142;
static double m_y     = 0.075094;
static double m_width = 0.017813;
static int    gDetail = 200;

// ------------------------- WinProc forward -------------------------
LRESULT CALLBACK WndProc(HWND hWnd, UINT msg, WPARAM wParam, LPARAM lParam);

// ===========================================================
//                Minimal OpenGL function loading
// ===========================================================

#ifndef GL_ARRAY_BUFFER
#define GL_ARRAY_BUFFER 0x8892
#endif
#ifndef GL_PIXEL_UNPACK_BUFFER
#define GL_PIXEL_UNPACK_BUFFER 0x88EC
#endif
#ifndef GL_DYNAMIC_DRAW
#define GL_DYNAMIC_DRAW 0x88E8
#endif

#ifndef APIENTRY
#define APIENTRY __stdcall
#endif
#ifndef APIENTRYP
#define APIENTRYP APIENTRY *
#endif

typedef void (APIENTRYP PFNGLGENBUFFERSPROC)(GLsizei, GLuint*);
typedef void (APIENTRYP PFNGLBINDBUFFERPROC)(GLenum, GLuint);
typedef void (APIENTRYP PFNGLBUFFERDATAPROC)(GLenum, ptrdiff_t, const void*, GLenum);
typedef void (APIENTRYP PFNGLDELETEBUFFERSPROC)(GLsizei, const GLuint*);

static PFNGLGENBUFFERSPROC    p_glGenBuffers    = 0;
static PFNGLBINDBUFFERPROC    p_glBindBuffer    = 0;
static PFNGLBUFFERDATAPROC    p_glBufferData    = 0;
static PFNGLDELETEBUFFERSPROC p_glDeleteBuffers = 0;

static void* gl_get_proc(const char* name)
{
  void* p = (void*)wglGetProcAddress(name);
  if(!p) {
    HMODULE ogl = GetModuleHandleA("opengl32.dll");
    if(ogl) p = (void*)GetProcAddress(ogl, name);
  }
  return p;
}

static int gl_load_ext()
{
  p_glGenBuffers    = (PFNGLGENBUFFERSPROC)gl_get_proc("glGenBuffers");
  p_glBindBuffer    = (PFNGLBINDBUFFERPROC)gl_get_proc("glBindBuffer");
  p_glBufferData    = (PFNGLBUFFERDATAPROC)gl_get_proc("glBufferData");
  p_glDeleteBuffers = (PFNGLDELETEBUFFERSPROC)gl_get_proc("glDeleteBuffers");

  if(!p_glGenBuffers || !p_glBindBuffer || !p_glBufferData || !p_glDeleteBuffers)
    return 0;
  return 1;
}

// ===========================================================
//                       OpenGL objects
// ===========================================================

static GLuint gPBO = 0;
static GLuint gTex = 0;

static void gl_release_all()
{
  if(gTex) {
    glDeleteTextures(1, &gTex);
    gTex = 0;
  }
  if(gPBO) {
    p_glDeleteBuffers(1, &gPBO);
    gPBO = 0;
  }

  if(gHgl) { wglMakeCurrent(NULL, NULL); wglDeleteContext(gHgl); gHgl = 0; }
  if(gHdc && gHwnd) { ReleaseDC(gHwnd, gHdc); gHdc = 0; }
}

static int gl_init_wgl(HWND hwnd)
{
  gHwnd = hwnd;
  gHdc = GetDC(hwnd);
  if(!gHdc) return 0;

  PIXELFORMATDESCRIPTOR pfd;
  ZeroMemory(&pfd, sizeof(pfd));
  pfd.nSize      = sizeof(pfd);
  pfd.nVersion   = 1;
  pfd.dwFlags    = PFD_DRAW_TO_WINDOW | PFD_SUPPORT_OPENGL | PFD_DOUBLEBUFFER;
  pfd.iPixelType = PFD_TYPE_RGBA;
  pfd.cColorBits = 32;
  pfd.cDepthBits = 16;
  pfd.iLayerType = PFD_MAIN_PLANE;

  int pf = ChoosePixelFormat(gHdc, &pfd);
  if(pf == 0) return 0;
  if(!SetPixelFormat(gHdc, pf, &pfd)) return 0;

  gHgl = wglCreateContext(gHdc);
  if(!gHgl) return 0;
  if(!wglMakeCurrent(gHdc, gHgl)) return 0;

  if(!gl_load_ext()) {
    printf("\nOpenGL buffer functions not available (need VBO/PBO support).");
    return 0;
  }

  // Setup a basic 2D projection for immediate-mode quad
  glDisable(GL_DEPTH_TEST);
  glViewport(0, 0, gW, gH);

  // Create PBO for RGBA pixels
  p_glGenBuffers(1, &gPBO);
  p_glBindBuffer(GL_PIXEL_UNPACK_BUFFER, gPBO);
  p_glBufferData(GL_PIXEL_UNPACK_BUFFER, (ptrdiff_t)(gW * gH * 4), 0, GL_DYNAMIC_DRAW);
  p_glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);

  // Create texture
  glGenTextures(1, &gTex);
  glBindTexture(GL_TEXTURE_2D, gTex);
  glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
  glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
  glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, gW, gH, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0);
  glBindTexture(GL_TEXTURE_2D, 0);

  return 1;
}

// ===========================================================
//                    OpenCL (GL sharing)
// ===========================================================

static int gCL_Ready = 0;

static cl_platform_id   gCL_Platform = 0;
static cl_device_id     gCL_Device   = 0;
static cl_context       gCL_Context  = 0;
static cl_command_queue gCL_Queue    = 0;
static cl_program       gCL_Program  = 0;
static cl_kernel        gCL_K_Mandel = 0;

static cl_mem gCL_PBO = 0; // CL view of GL PBO

static const char* gCL_Source =
"__kernel void mandelbrot(__global uchar4* out, int width, int height,             \n"
"  double mx, double my, double mwidth, int detail)                                 \n"
"{                                                                                 \n"
"  int xpix = get_global_id(0);                                                    \n"
"  int ypix = get_global_id(1);                                                    \n"
"  if(xpix >= width || ypix >= height) return;                                     \n"
"                                                                                  \n"
"  double x0 = mx + ((double)xpix) * mwidth / (double)width;                       \n"
"  double y0 = my + ((double)(height - 1 - ypix)) * mwidth / (double)width;        \n"
"                                                                                  \n"
"  double zx = 0.0, zy = 0.0;                                                      \n"
"  int times = 0;                                                                  \n"
"  int inset = 1;                                                                  \n"
"  while(inset && times < detail)                                                  \n"
"  {                                                                               \n"
"    times++;                                                                      \n"
"    double zxs = zx*zx;                                                           \n"
"    double zys = zy*zy;                                                           \n"
"    zy = 2.0*zx*zy + y0;                                                          \n"
"    zx = zxs - zys + x0;                                                          \n"
"    if(zxs + zys >= 4.0) inset = 0;                                               \n"
"  }                                                                               \n"
"                                                                                  \n"
"  uchar4 c;                                                                       \n"
"  if(inset) {                                                                     \n"
"    c = (uchar4)(0,0,0,255);                                                      \n"
"  } else {                                                                        \n"
"    float t = (float)times / (float)detail;                                       \n"
"    float r = clamp(1.5f - fabs(4.f*t - 3.f), 0.f, 1.f);                          \n"
"    float g = clamp(1.5f - fabs(4.f*t - 2.f), 0.f, 1.f);                          \n"
"    float b = clamp(1.5f - fabs(4.f*t - 1.f), 0.f, 1.f);                          \n"
"    c = (uchar4)((uchar)(255*r),(uchar)(255*g),(uchar)(255*b),255);               \n"
"  }                                                                               \n"
"                                                                                  \n"
"  out[ypix*width + xpix] = c;                                                     \n"
"}                                                                                 \n";

static void cl_release_all()
{
  if(gCL_PBO)     { clReleaseMemObject(gCL_PBO);     gCL_PBO = 0; }
  if(gCL_K_Mandel){ clReleaseKernel(gCL_K_Mandel);   gCL_K_Mandel = 0; }
  if(gCL_Program) { clReleaseProgram(gCL_Program);   gCL_Program = 0; }
  if(gCL_Queue)   { clReleaseCommandQueue(gCL_Queue);gCL_Queue = 0; }
  if(gCL_Context) { clReleaseContext(gCL_Context);   gCL_Context = 0; }

  gCL_Device = 0;
  gCL_Platform = 0;
  gCL_Ready = 0;
}

// Find a GPU device that supports cl_khr_gl_sharing (best effort)
static int cl_pick_device_with_glshare(cl_platform_id* outP, cl_device_id* outD)
{
  cl_uint nPlatforms = 0;
  if(clGetPlatformIDs(0, 0, &nPlatforms) != CL_SUCCESS || nPlatforms == 0)
    return 0;

  cl_platform_id platforms[8];
  if(nPlatforms > 8) nPlatforms = 8;
  if(clGetPlatformIDs(nPlatforms, platforms, &nPlatforms) != CL_SUCCESS)
    return 0;

  for(cl_uint p=0; p<nPlatforms; p++)
  {
    cl_uint nDev = 0;
    if(clGetDeviceIDs(platforms[p], CL_DEVICE_TYPE_GPU, 0, 0, &nDev) != CL_SUCCESS || nDev == 0)
      continue;

    cl_device_id devs[8];
    if(nDev > 8) nDev = 8;
    if(clGetDeviceIDs(platforms[p], CL_DEVICE_TYPE_GPU, nDev, devs, &nDev) != CL_SUCCESS)
      continue;

    for(cl_uint d=0; d<nDev; d++)
    {
      char ext[8192];
      size_t sz = 0;
      if(clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, sizeof(ext), ext, &sz) != CL_SUCCESS)
        continue;

      // Must contain cl_khr_gl_sharing for Acquire/Release GL objects
      if(strstr(ext, "cl_khr_gl_sharing"))
      {
        *outP = platforms[p];
        *outD = devs[d];
        return 1;
      }
    }
  }

  return 0;
}

static int cl_init_glshare()
{
  cl_int err = CL_SUCCESS;

  cl_platform_id P = 0;
  cl_device_id   D = 0;

  if(!cl_pick_device_with_glshare(&P, &D)) {
    printf("\nOpenCL: no GPU device with cl_khr_gl_sharing found.");
    return 0;
  }

  gCL_Platform = P;
  gCL_Device   = D;

  // Create an OpenCL context that shares with the CURRENT WGL context
  cl_context_properties props[] = {
    CL_GL_CONTEXT_KHR,   (cl_context_properties)wglGetCurrentContext(),
    CL_WGL_HDC_KHR,      (cl_context_properties)wglGetCurrentDC(),
    CL_CONTEXT_PLATFORM, (cl_context_properties)gCL_Platform,
    0
  };

  gCL_Context = clCreateContext(props, 1, &gCL_Device, 0, 0, &err);
  if(err != CL_SUCCESS || !gCL_Context) { cl_release_all(); return 0; }

  gCL_Queue = clCreateCommandQueue(gCL_Context, gCL_Device, 0, &err);
  if(err != CL_SUCCESS || !gCL_Queue) { cl_release_all(); return 0; }

  gCL_Program = clCreateProgramWithSource(gCL_Context, 1, &gCL_Source, 0, &err);
  if(err != CL_SUCCESS || !gCL_Program) { cl_release_all(); return 0; }

  // Enable double precision if needed by your device; some require -cl-fast-relaxed-math etc.
  // We keep it empty; if build fails, print log.
  err = clBuildProgram(gCL_Program, 1, &gCL_Device, 0, 0, 0);
  if(err != CL_SUCCESS)
  {
    char logbuf[8192];
    size_t logsz = 0;
    clGetProgramBuildInfo(gCL_Program, gCL_Device, CL_PROGRAM_BUILD_LOG, sizeof(logbuf), logbuf, &logsz);
    printf("\nOpenCL build failed:\n%s", logbuf);
    cl_release_all();
    return 0;
  }

  gCL_K_Mandel = clCreateKernel(gCL_Program, "mandelbrot", &err);
  if(err != CL_SUCCESS || !gCL_K_Mandel) { cl_release_all(); return 0; }

  // Create CL memory object from GL PBO
  gCL_PBO = clCreateFromGLBuffer(gCL_Context, CL_MEM_WRITE_ONLY, gPBO, &err);
  if(err != CL_SUCCESS || !gCL_PBO) { cl_release_all(); return 0; }

  gCL_Ready = 1;
  printf("\nOpenCL: GL-sharing enabled.");
  return 1;
}

// ===========================================================
//                      Render (CL -> GL)
// ===========================================================

static void RenderFrame()
{
  if(!gCL_Ready) return;

  size_t global[2];
  global[0] = (size_t)gW;
  global[1] = (size_t)gH;

  // Optional local size (often helps). If it fails on some devices, set NULL instead.
  size_t local[2];
  local[0] = 16;
  local[1] = 16;

  cl_int err = CL_SUCCESS;

  // Acquire the GL PBO for CL
  err = clEnqueueAcquireGLObjects(gCL_Queue, 1, &gCL_PBO, 0, 0, 0);
  if(err != CL_SUCCESS) return;

  int arg = 0;
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(cl_mem), &gCL_PBO);
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(int),    &gW);
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(int),    &gH);
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(double), &m_x);
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(double), &m_y);
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(double), &m_width);
  clSetKernelArg(gCL_K_Mandel, arg++, sizeof(int),    &gDetail);

  err = clEnqueueNDRangeKernel(gCL_Queue, gCL_K_Mandel, 2, 0, global, local, 0, 0, 0);
  if(err != CL_SUCCESS)
  {
    // Fallback: try without local size
    err = clEnqueueNDRangeKernel(gCL_Queue, gCL_K_Mandel, 2, 0, global, 0, 0, 0, 0);
  }

  clEnqueueReleaseGLObjects(gCL_Queue, 1, &gCL_PBO, 0, 0, 0);
  clFinish(gCL_Queue);

  // Upload PBO -> texture (GPU->GPU)
  p_glBindBuffer(GL_PIXEL_UNPACK_BUFFER, gPBO);
  glBindTexture(GL_TEXTURE_2D, gTex);
  glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, gW, gH, GL_RGBA, GL_UNSIGNED_BYTE, 0);
  p_glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);

  // Draw fullscreen quad (compat immediate mode)
  glClear(GL_COLOR_BUFFER_BIT);
  glEnable(GL_TEXTURE_2D);
  glBindTexture(GL_TEXTURE_2D, gTex);

  glBegin(GL_QUADS);
    glTexCoord2f(0,0); glVertex2f(-1,-1);
    glTexCoord2f(1,0); glVertex2f( 1,-1);
    glTexCoord2f(1,1); glVertex2f( 1, 1);
    glTexCoord2f(0,1); glVertex2f(-1, 1);
  glEnd();

  glBindTexture(GL_TEXTURE_2D, 0);

  SwapBuffers(gHdc);
}

// ===========================================================
//                         WinMain
// ===========================================================

int WINAPI WinMain(HINSTANCE hInst, HINSTANCE, LPSTR, int)
{
  const char* szClass = "MandelCLGLClass";
  UnregisterClassA(szClass, hInst);

  WNDCLASSEXA wc;
  ZeroMemory(&wc, sizeof(wc));
  wc.cbSize = sizeof(wc);
  wc.style = CS_HREDRAW | CS_VREDRAW;
  wc.lpfnWndProc = WndProc;
  wc.hInstance = hInst;
  wc.hCursor = LoadCursor(NULL, IDC_ARROW);
  wc.lpszClassName = szClass;
  RegisterClassExA(&wc);

  RECT r;
  r.left=0; r.top=0; r.right=gW; r.bottom=gH;
  AdjustWindowRect(&r, WS_OVERLAPPEDWINDOW, FALSE);

  HWND hwnd = CreateWindowExA(
    0, szClass, "Mandelbrot (OpenCL compute + OpenGL display)",
    WS_OVERLAPPEDWINDOW,
    100, 100, (r.right-r.left), (r.bottom-r.top),
    0, 0, hInst, 0);

  if(!hwnd) return 0;

  ShowWindow(hwnd, SW_SHOW);
  UpdateWindow(hwnd);

  // Init OpenGL first (we need a current GL context before creating CL shared context)
  if(!gl_init_wgl(hwnd))
  {
    MessageBoxA(hwnd, "OpenGL init failed", "Error", MB_OK);
    gl_release_all();
    return 0;
  }

  if(!cl_init_glshare())
  {
    MessageBoxA(hwnd, "OpenCL GL-sharing init failed", "Error", MB_OK);
    cl_release_all();
    gl_release_all();
    return 0;
  }

  // Message loop + continuous redraw (idle rendering)
  MSG msg;
  ZeroMemory(&msg, sizeof(msg));

  while(msg.message != WM_QUIT)
  {
    while(PeekMessage(&msg, NULL, 0, 0, PM_REMOVE))
    {
      TranslateMessage(&msg);
      DispatchMessage(&msg);
    }

    RenderFrame();
  }

  cl_release_all();
  gl_release_all();
  gHwnd = 0;
  return 0;
}

// ===========================================================
//                         Input / Zoom
// ===========================================================

LRESULT CALLBACK WndProc(HWND hWnd, UINT msg, WPARAM wParam, LPARAM lParam)
{
  switch(msg)
  {
    case WM_RBUTTONDOWN:
      // reset to full fractal
      m_x = -2.5;
      m_y = -2.0;
      m_width = 4.0;
      return 0;

    case WM_LBUTTONDOWN:
    {
      RECT rect;
      GetClientRect(hWnd, &rect);

      int x = (int)(lParam & 0xFFFF);
      int y = (int)((lParam >> 16) & 0xFFFF);

      double zoom = 0.5;

      double ax = (double)x / (double)(rect.right  ? rect.right  : 1);
      double ay = (double)y / (double)(rect.bottom ? rect.bottom : 1);

      m_x += m_width * (ax - 0.5);
      m_y += m_width * (0.5 - ay);
      m_width *= zoom;
      return 0;
    }

    case WM_KEYDOWN:
      if(wParam == VK_ESCAPE || wParam == VK_F12) {
        PostMessage(hWnd, WM_CLOSE, 0, 0);
        return 0;
      }
      return 0;

    case WM_DESTROY:
      PostQuitMessage(0);
      return 0;
  }

  return DefWindowProc(hWnd, msg, wParam, lParam);
}

DLLFUNC int main()
{
  return WinMain(GetModuleHandleA(NULL), NULL, GetCommandLineA(), SW_SHOWDEFAULT);
}

Last edited by TipmyPip; 1 hour ago.
ZorroGPT - https://bit.ly/3Gbsm4S