Practical differences cuda / opencl

Once again, I find myself porting CUDA code to OpenCL. And I had to remind myself of the process, and the differences between the two. Let's document, so future Bram will catch up quickly again.

The major difference is of course compatibility. OpenCL is supported by Nvidia, Amd, Intel, etc. CUDA is only supported by Nvidia. It's the whole reason of this porting exercise.

The CUDA-kernel can be pre-compiled by the developer using nvcc, and then be shipped as a PTX file. The OpenCL-kernel typically ships as C-like source.

The function definition of a CUDA-kernel gets prefixed with __kernel whereas an OpenCL uses a __global__ prefix.

A CUDA-kernel uses blockIdx and blockDim to determine which sub-part of the input data it is supposed to process. An OpenCL kernel uses get_global_id() instead.

Using 16 bit floating point arithmetic is as easy as #include <cuda_fp16.h> and using __half. On OpenCL you need to check for an extension and add #pragma OPENCL EXTENSION cl_khr_fp16 : enable

In CUDA, constant input data is prefixed with __constant__ which the host then sets with the cuMemcpyHtoDAsync() function after getting the symbol using the cuModuleGetGlobal() function. In OpenCL if the constant input data is large, you prefix it with __global__ and the host sets it with the clCreateBuffer() and clEnqueueWriteBuffer()functions.

Half float literals need a constructor in CUDA: __half(1.0) but can use a suffix in OpenCL, where 1.0h can be used.

Many operators need a function call in CUDA. The reciprocal in CUDA is hrcp(x) and a simple <= on a half float becomes __hlt(a,b) which makes it less legible.

A minor diffence is that CUDA's atomicAdd is atomic_add in OpenCL.

CUDA has asserts but OpenCL does not.

Minecraft Server

When it comes to video games, I don't think Minecraft has been surpassed. Like all the best games, it has a simple premise: mine blocks in an infinite world, and use them to craft other blocks. But from this simple fundement, a master piece arose.

Even after the purchase by Microsoft, the game is still available for Linux, as Mojang maintains a Java version of their game. Lately, I have been looking into running a server, and intend to write custom code, server-side, to see what can be done with that. Minecraft lets you modify the client code. But this post investigates the modification of the Minecraft server.

Minecraft modding has a long history. But it appears that my best entry point would be SpigotMC.

After executing BuildTools.jar, I ended up with a file called spigot-1.19.3.jar which I copied to my minecraft server root directory. Instead of running minecraft_server.1.19.3.jar I now run the spigot jar file instead.

To make sure newer minecraft clients can connect, you can put ViaVersion.jar in the plugins directory.

Next up, is finding out how to write custom server code. So far I have found Kody Simpson's Youtube series an excellent resource.

Bukkit is a Java Interface, and it is implemented by CraftBukkit.

CraftBukkit is implemented using net.minecraft.server code.

This means that the same object can be referenced and used at three different levels. For instance, the player exists as Bukkit's Player, as CraftBukkit's CraftPlayer and as Mojang's EntityPlayer.

Only the Bukkit interface remains stable between releases, and can be used on different minecraft server versions. Also, the NMS (net.minecraft.server from Mojang) code is decompiled, and obfuscated.

Does coding an NPC require NMS? Or can it be done in SpigotMC? If not, what about PaperMC?

To deal with the obfuscation of n.m.s. there are the so called Mojang mappings. Initially I was unable to use those with Java 19. But using a development version of SpecialSource I was able to use n.m.s. with Java 19.

When getting confused in API terminology, this is a good reminder.

I noticed that using the remapped API, really slowed down my iteration speed. Building my plugin went from a few seconds, to 1 minute, because the remapper is slow.

The first restults from my experiments: The Little Vex That Could. A construction-buddy that builds castles from IKEA flat-packs.

Espressif IoT Development Framework

So far, I have been programming my µcontrollers using the Arduino IDE. As I progressed with my TurboLEDz experiments, I came across different µcontrollers, including the ESP32 series from Espressif.

Espressif actually has their own SDK, which targets the command line. As I feel quite at home on the command-line, I decided to switch over to the Espressif SDK, which is called ESP IDF, or ESP IoT Development Framework.

This environment is based on cmake and python tools that wrap it. Nowadays, it recommends to use it with vscode IDE, but frankly, I found it to work better without.

You mainly interface with the build environment using the idf.py tool. The typical sequence would be:

    $ idf.py set-target esp32s2
    $ idf.py build
    $ idf.py flash
    $ idf.py monitor
  

And what I really appreciate from ESP IDF is that you can code in pure C, without any C++ stuff involved. No need for .ino files, either. Plain, simple C. I love it.

Speaking of Internet of Things... I've always found it to be a nebulous term. But I guess it pertains to devices that don't have a screen, keyboard, mouse, but do have Ethernet/WiFi? And thanks to the march of progress, those tiny little devices are typically as capable as desktop computers were two decades ago?

My favourite µcontrollers are RISCV based. And some devices really blur the line between µcontroller and single-board-computer. Like the Pine64's Ox64 which is powerful enough to run Linux (but witout the protected memory, I think.)

I also like the ESP32 with two cores, a lot. It enabled me to create this voice-changer where the two cores can tag-team: one core is computing the Fourier Transformations while the other is doing I2S, and then they switch. So one core does everything for the odd-numbered batches with samples, and one core does everything for the even-numbered batches. Nice low-latency concurrency. Yay! But that project was all in Arduino IDE.

For my fist IDF project, I wrote a single-header-file piece of code for use with Adafruit's LED backpack. See my ledbp.h gist. It sits on top of Espressif's driver/i2c.h code.

The $8 Linux computer (part 2.)

Caveat Emptor! At the time of writing, ethernet, wifi and SD-Card are not working. It's good to realize this before jumping into this.

Things have improved a little, since writing PART 1 of my experience. Instead of a patched bflb-mcu-tool to flash, it is now possible to use a modified picoprobe [DOWNLOAD HERE] with the Dev Cube flashing tool.

Also, there is a patch for the PIN definitions so that we can capture the UART of the D1 core (which is a T-HEAD C906) on pins GPIO17/GPIO16.

With the fixed pin definitions we can hook up a USB to UART dongle, like the CH340G, to the second UART as shown below.

In this picture, the red dongle is to capture the output from the linux console running on the T-HEAD C906 core. The pink dongle is to program the Ox64. Both of these need to be connected to your PC. The Ox64 itself needs no connection to the PC, other than a power feed into the Micro-USB.

So, once we wired this all up, flashed to binaries, what do we get?

We get root!

And that's after booting in a blistering 0.22 seconds!

Full boot log.

I love Linux on a diet. So mean and lean.

Next up: Can we get networking?

The $8 linux computer (part 1.)

Yes, you read that right... for $8 you can run linux on a RISC-V computer. And not just a stripped down kernel without an MMU. No, the Ox64 comes with capable RISC-V cores. Heck, and not just one core, but three cores (all different.) Exciting times!

Caveat Emptor! At the time of writing, ethernet, wifi and SD-Card are not working. It's good to realize this before jumping into this.

It looks like a microcontroller, but behaves like a UNIX machine. Sweet, or what? Getting it up and running, though, is far from trivial. There are so many hurdles to take!

Hurdle 1: You cannot program it using USB, like you would an Arduino, ESP32, Stamp-C3 or rPi Pico. Nope... this needs to be programmed via the UART TX/RX pins. In the picture above, you can see them hooked up to GPIO14 / GPIO15 and GND.

Hudle 2: The flashing tool Dev Cube will crash on launch when using Ubuntu on Wayland: you have to switch to Xorg instead.

$ ./BLDevCube-ubuntu 
Warning: Ignoring XDG_SESSION_TYPE=wayland on Gnome. Use QT_QPA_PLATFORM=wayland to run on Wayland anyway.
/usr/lib/x86_64-linux-gnu/gio/modules/libdconfsettings.so: undefined symbol: g_log_structured_standard
Failed to load module: /usr/lib/x86_64-linux-gnu/gio/modules/libdconfsettings.so
GLib-GIO-Message: Using the 'memory' GSettings backend.  Your settings will not be saved or shared with other applications.

(BLDevCube-ubuntu:54646): GLib-GIO-ERROR **: Settings schema 'org.gnome.settings-daemon.plugins.xsettings' does not contain a key named 'antialiasing'
Trace/breakpoint trap (core dumped)
  

So, UART programming then... fine, we'll just use one of those little CH340G dongles. I got one off Amazon, but it would not work... every time the handshake would fail in Dev Cube (hurdle 3).

Luckily, some people found out that you can program it with an RP2040 device, like the rPi Pico. I had a KB2040 from Adafruit lying around, so used that.

Which puts us at hurdle 4: How to turn a KB2040 into a serial bridge between PC and Ox64? Hook up GPIO-4 of the RP2040 to GPIO-15 of the Ox64. Hook up GPIO-5 of the RP2040 to GPIO-14 of the Ox64. Also, connect the grounds, of course. See below how to hook up for flashing the M0.\

Then we need to load our RP2040 with picoprobe which we can do by holding BOOT while powering up, and feeding it the uf2 file with picoprobe.

Ok, so now we can flash it then? Unfortunately, no, hurdle 5: DevCube is still not able to flash the Ox64, but we can use bflb-mcu-tool instead.

Ok, but with bflb-mcu-tool, we are all set! Are we? Introducing hurdle 6: we need to modify the tool before it will work for us. This means stripping DTR and RTS stuff from the tool by editing the ~/.local/lib/python3.10/site-packages/bflb_mcu_tool/libs/bflb_interface_uart.py file. Here is my copy.

So with the Ox64 powered via MicroUSB, and the RP2040 connected to my PC using USB-C, I can finally flash something. For that, I took the low_load_bl808_m0.bin from the bl808 linux guide and flashed it with:

$ bflb-mcu-tool --chipname=bl808 --port=/dev/ttyACM2 --baudrate=2000000 --firmware=out/low_load_bl808_m0.bin
    

After which I now get output from my Ox64, via my KB2040 to my PC:

dynamic memory init success,heap size = 156 Kbyte 
E907 start...
mtimer clk:1000000
psram clk init ok!
    

Ok, a modest start, but the Ox64 is talking to my PC now! Yay! Come back to my blog at a later time to see if I will manage to run Linux on the Ox64! Because as I understand it, I also need to flash something to the D0 and then flash the actual OS as well.

UPDATE: Here is PART 2.

TurboLEDz Voice Changer.

I've released the 6th product under the TurboLEDz name. This time it is a voice changer. It makes you sound like a bad asss, by doing real time FFT, frequency manupulation and iFFT to go back to the time domain. It's built with an I2S mic and I2S DAC/Amplifier. To make the μcontroller capable of keeping up with the samples, I've programmed it as multi-threaded, dual core. So if one core is busy doing math, the other is still available to do IO. And the cores will tag-team. One core for the even numbered chunks, the other core for the odd numbered chunks. It works really well, as you can see in the video.

Super Capacitors

I bought three super capacitors from Adafruit, model DDLE2R5LGN701KAA5S. They are rated at 700 Farad at 2.5V. (They have very little stock left, at time of writing, so if you want them too, act quickly!)

If I place them in series, I will end up with 233 Farad at the very maximum of 7.5V. However, charging to the max 7.5V is dangerous: one of the capacitors can easily overcharge and explode, that way. Unless you balance the capacitors.

Balancing a capacitor bank is not trivial. But there seems to be an alternative, as offered by Patrick of EngineeringShock.

Patrick's advice is to charge the capacitors to 80% of their nominal Voltage, so that even if there is a mismatch in internal resistance between the capacitors in series, none of them will overcharge. I like this approach, as it is simpler, does not require extra circuitry, at the cost of lost capacitance.

So if we charge a 7.5V bank to 80%, we get a charging voltage of 6 Volts. Let's make that our target. When we have have our bank charged to 6V, we can feed that into a voltage regulator, and as the bank drains to 3.3V, we should be able to ride from 6V to 3.3V and run a μcontroller off that, or some other low-power electronics.

And this leads me to my next objective: how to generate a steady 6V charging voltage? I would like to use solar panels for this. So let's dive into photovoltaics.

Reading the specification of a PV panel, there is the nominal voltage, which we can't really rely on. Then there is the Voc (Open Circuit Voltage) which we will measure if there is no current flowing. And also the Max Power Voltage, which is lower.

I need to read up on how to make sure that a photovoltaic panel output a steady 6V, and never any higher, at a tight tolerance. PV panels typically use charge controllers for that, but I want something simple. Maybe an opamp?

UPDATE

I am taking a different approach. I've decided to charge the capacitors in parallel, not in series, at a steady 2.5V.

To achieve this, I take input from a solar panel (with wildly varying voltage, and current) and convert that with a buck-converter to 2.5V that can be used to charge the capacitors.

I then use the capacitors as input for a boost-converter that brings a varying low voltage (It is only 2.5V when fully charged, and they discharge linearly) to a steady 3.3V. This I can then use in a client circuit, typically a μcontroller with some peripherals attached.