OpenCL (Open Coumputer Language) for FreeBASIC .

Headers, Bindings, Libraries for use with FreeBASIC, Please include example of use to help ensure they are tested and usable.
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

The Open Coumputer Language for FreeBASIC.

download: OpenCL.zip from: Nov 16, 2021

import lib's 32/64-bit and the OpenCL specs. are included.

Don't forget the "readme" ;-)

Joshy

Simple test:

The default OpenGL shader from shadertoy.com as OpenCL version:

Code: Select all

#include "cl.bi"

dim as string CODE
CODE &= !"kernel void shader (global uchar4 * pixels, uint width, uint height, float time) \n"
CODE &= !"{\n"
CODE &= !"  uint ix = get_global_id(0); \n"
CODE &= !"  uint iy = get_global_id(1); \n"
CODE &= !"  uint i  = ix+width*iy;      \n"
CODE &= !"  float x = ix/(float)width;  \n"
CODE &= !"  float y = iy/(float)height; \n"
CODE &= !"  uchar r = (uchar)(x*255.f); \n"
CODE &= !"  uchar g = (uchar)(y*255.f); \n" 
CODE &= !"  uchar b = (uchar)(255.f * (0.5f + 0.5f*sin(time))); \n" 
CODE &= !"  pixels[i] = (uchar4)(b,g,r,255); \n"
CODE &= !"}\n"

dim as zstring ptr pCode = strptr(CODE)
dim as cl_platform_id platform
' get first platform
dim as cl_int status = clGetPlatformIDs(1,@platform,NULL)


' get first device from platform try a GPU device at first
dim as cl_device_id device
status = clGetDeviceIDs(platform,CL_DEVICE_TYPE_GPU,1,@device,NULL)
if status = CL_INVALID_DEVICE_TYPE then
  ' fallback: try a CPU device also
  status = clGetDeviceIDs(Platform,CL_DEVICE_TYPE_CPU,1,@device,NULL)
end if
if status then
  print "error: can't create OpenCL device !"
  flip : beep : sleep : end 1
end if

' create a device context
dim as cl_context context = clCreateContext(0,1,@device,NULL,NULL,@status)

' create program object in context from sourcecode
dim as cl_program program = clCreateProgramWithSource(context,1,@pCode,NULL,@status)

' compile and link the program
status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL)
if status then
  print "error: can't create compile unand link kernel !"
  dim as size_t nChars ' get size of log
  clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,0,0,@nChars)
  dim as string strLog=space(nChars)
  ' get log from compiler/linker
  clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,nChars,strptr(strLog),NULL)
  print strLog 
  flip : beep : sleep : end 1
end if

' create a kernel object
dim as cl_kernel kernel = clCreateKernel(program, "shader", @status)

' get size of workgroup from device
dim as size_t WorkGroupSize
status = clGetKernelWorkGroupInfo(kernel,device,CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t),@WorkGroupSize,NULL)

' calculate local worksize
'(must be workgroupsize div power of two)
dim as size_t LocalSize = WorkGroupSize
while LocalSize*LocalSize > WorkGroupSize
  LocalSize shr=1
wend

dim as cl_int GlobalWidth=640,GlobalHeight = GlobalWidth/16*9

' calculate global worksize (should be a multiply of local worksize
GlobalWidth -=GlobalWidth  mod LocalSize
GlobalHeight-=GlobalHeight mod LocalSize

screenres GlobalWidth,GlobalHeight,32,2
screenset 1,0

windowtitle "global: " & GlobalWidth & " x " & GlobalHeight & " local: " & LocalSize & " x " & LocalSize & " worksize"

dim as any ptr pixels,img=imagecreate(GlobalWidth,GlobalHeight)
imageinfo img,,,,,pixels
dim as size_t bufSize = GlobalWidth*GlobalHeight*4
dim as size_t globalSizes(...) => {GlobalWidth,GlobalHeight}
dim as size_t localSizes (...) => {LocalSize  ,LocalSize}

' create a memory buffer in the context
dim as cl_mem buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, bufSize, NULL, @status)

' set the parameters of the kernel
status = clSetKernelArg(kernel, 0, sizeof(cl_mem ),@buffer)
status = clSetKernelArg(kernel, 1, sizeof(cl_uint),@GlobalWidth)
status = clSetKernelArg(kernel, 2, sizeof(cl_uint),@GlobalHeight)

' create command queue
dim as cl_command_queue queue = clCreateCommandQueue(context,device,0,@status)
if status then
  print "error: can't create a command queue !"
  flip : beep : sleep : end 1
end if

dim as integer frames,fps
dim as double  tStart=timer
dim as double  tLast =tStart
while inkey()=""
  dim as single runtime = timer()-tStart
  ' set the time parameter
  status = clSetKernelArg(kernel, 3, sizeof(cl_float),@runtime)
  ' execute the "2D" kernel
  status = clEnqueueNDRangeKernel(queue,kernel,2, NULL, @globalSizes(0), @localSizes(0),0,NULL,NULL)
  ' read/copy the device buffer in the image buffer
  status = clEnqueueReadBuffer   (queue,buffer,CL_TRUE,0,bufSize,pixels,0,NULL,NULL)
  ' draw the runtime and fps in the image
  draw string  img,(0,0),"fps: " & fps & " time: " & runtime 
  ' put it on the screen
  put (0,0),img,PSET
  ' make hiden page visible
  flip
  ' count the frames
  frames+=1
  ' update every 60 frame the fps
  if frames mod 60=0 then
    var tNow=timer()
    fps=60/(tNow-tLast)
    tLast=tNow
  end if
wend

' free all resources
clReleaseCommandQueue queue
clReleaseProgram      program
clReleaseMemObject    buffer
clReleaseContext      context
Last edited by D.J.Peters on Oct 12, 2022 17:55, edited 9 times in total.
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

On a linux 32-bit box I must use /usr/lib-i386-linux-gnu as path

Joshy

joshy@P4:~/OpenCL$ fbc -p /usr/lib-i386-linux-gnu ./test01.bas
joshy@P4:~/OpenCL$ ./test01
JohnK
Posts: 279
Joined: Sep 01, 2005 5:20
Location: Earth, usually
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by JohnK »

I copy cl.bi to FreeBASIC\inc folder
then copy lib/win64/libOpenCL.dll.a to FreeBASIC\lib\win64 folder
compile error
...\bin\win64\ld.exe: cannot find -lOpenCL

Do I need to find the static lib? or DLL? Can't find this at
https://github.com/KhronosGroup
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

fixed now

Joshy
Last edited by D.J.Peters on Mar 13, 2016 5:36, edited 1 time in total.
JohnK
Posts: 279
Joined: Sep 01, 2005 5:20
Location: Earth, usually
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by JohnK »

I do have it in this beautiful folder name
C:\Windows\WinSxS\wow64_microsoft-windows-r..xwddmdriver-wow64-c_31bf3856ad364e35_10.0.10586.0_none_3dae054b56911c22

Just for *fun* I copied all dll files to the local directory of the fb OpenCL source, now I get

C:\FreeBASIC-1.05.0-win64\bin\win64\ld.exe: skipping incompatible C:\FreeBASIC-1.05.0-win64\lib\win64/libOpenCL.dll.a when searching for -lOpenCL
./OpenCL.dll: file not recognized: File format not recognized

Like most modern software, I am burning in dependency Hell :-)
St_W
Posts: 1619
Joined: Feb 11, 2009 14:24
Location: Austria
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by St_W »

Most likely you are mixing 32-bit and 64-bit binaries. You can use "depends" for tracking any dependency problems of already compiled binaries. If you only have libraries you can use objdump or file commands. See e.g. http://stackoverflow.com/questions/5665 ... -or-64-bit
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

Sorry for the trouble I used the wrong commandline switch :-(
fbc wrote:D:\fbExamples\fbOpenCL\OpenCL>fbc -gen gcc -arch x86_64 test01.bas
C:\FreeBASIC\bin\win64\ld.exe: skipping incompatible C:\FreeBASIC\lib\win64/libO
penCL.dll.a when searching for -lOpenCL
C:\FreeBASIC\bin\win64\ld.exe: skipping incompatible C:\FreeBASIC\lib\win64/libO
penCL.dll.a when searching for -lOpenCL
C:\FreeBASIC\bin\win64\ld.exe: cannot find -lOpenCL

@JonhK may be the file FreeBASIC/lib/win64/libOpenCL.dll.a isn't a 64-bit import lib
In this case delete FreeBASIC/lib/win64/libOpenCL.dll.a so FreeBASIC can't use it.
Be sure you have OpenCL.dll in same folder as test01.bas and rebuild it.

Joshy
Last edited by D.J.Peters on Mar 13, 2016 3:24, edited 4 times in total.
srvaldez
Posts: 3373
Joined: Sep 25, 2005 21:54

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by srvaldez »

<off topic> dependency walker is lousy at it's purported function, I wish there was a real utility that would show whether an application was 32 or 64 bit and the true dependencies, for 32-bit I use exescope a shareware utility that works quite well, too bad it seems to be discontinued.
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

@JonhK problem fixed in OpenCL.zip
It was a corrupt 64-bit import lib.

You have to overwrite OpenCL.dll.a in your FreeBASIC\lib\win64 folder with the new version.

Joshy
JohnK
Posts: 279
Joined: Sep 01, 2005 5:20
Location: Earth, usually
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by JohnK »

Test(02).exe ~ 2800 - 3200 fps... oooohhhh

Joshy you are my hero.
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

JohnK wrote:Test(02).exe ~ 2800 - 3200 fps
Only for info how many FPS from test01.bas please :-)

Joshy
JohnK
Posts: 279
Joined: Sep 01, 2005 5:20
Location: Earth, usually
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by JohnK »

For Test(01) I get ~360 fps

I know it is not comparable, but the same effect with shadertoy offline was about 2800 fps 8-(
Maybe you are thinking that
FB exports c code
C code is then used for kernels
then FB is now a high performance parallel programming language
D.J.Peters
Posts: 8586
Joined: May 28, 2005 3:28
Contact:

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by D.J.Peters »

You know it ?

FPS from test01.bas mensured primary the bandwidth of memory transfer CPU and GPU.

For example your result of 360 FPS means:
640 pixels * 350 pixels * 4 bytes per pixel = 896 000 bytes * 360 Frames = 322 560 000 bytes per second GPU ( kernel write and read in image)
But wait thats not all PLUS 360 times of PUT(0,0),img = 322 560 000 bytes CPU (read write)
But wait thats not all PLUS 360 times of FLIP 322 560 000 read hidden page and write to the window. CPU (read write)

322 560 000 * 3 (write) + 322 560 000 * 3 (read) = 1 935 360 000 bytes = ~1,8 GB per second

FPS from test02.bas mensured primary the speed of kernel execution.
(without put and flip you would get 50% more FPS to get the real execution time of the kernel)

Joshy
srvaldez
Posts: 3373
Joined: Sep 25, 2005 21:54

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by srvaldez »

hello D.J.Peters
I could not get OpenCL to work on my Windows VM but managed to get it to work on my Mac.
on test01 I get about 1000, on test02 I get between 700 and 1000+, it's hard to tell because the numbers vary so quickly.
btw, I get this warning when compiling test02
test02.bas(85) warning 3(1): Passing different pointer types, at parameter 6 of CLGETPROGRAMBUILDINFO()

[offtopic] I could not find an OpenCL lib on my Mac but I found OpenCL <no extension> in /System/Library/Frameworks/OpenCL.framework/Versions/A/
I copied that file to a temp directory and renamed it to libOpenCL.a and copied that file to the local/lib directory and everything works :)
srvaldez
Posts: 3373
Joined: Sep 25, 2005 21:54

Re: OpenCL (Open Coumputer Language) for FreeBASIC .

Post by srvaldez »

hi D.J.Peters
I made a small change to your fps routine to average the fps
insert right after the fps calculation, where count and fpsacum are declared as integer

Code: Select all

    count+=1:fpsacum+=fps
    fps=fpsacum/count
with that change I get about 1080 on test01 and about 1200 on test02, after running each test about 1 minute
Post Reply