Add a specification for the C-ABI design

[joshbainbridge]

Having a C-ABI will allow exposure to more languages. This is with the broader aim to improve adoption of the library. #87.

This review is for the initial design specification prior to any development.

For easier reading, you can enable the formatted display using this option on the Files changed tab:

[joshbainbridge]

@fpsunflower and @hills, I've updated the spec with Options A (original macros) and Option B (function tables). Would be great to get your thoughts. Now that I've seen them both laid out and done a quick POC, I'm leaning towards Option B. The code at least appears a lot cleaner to maintain and extent. But still open to either option.

[fpsunflower]

Do we really want to keep the ability to pick the sampler at runtime? That seems like an extra cost that will prevent a lot of inlining and complicates eventually porting the code to other GPU languages like GLSL or HLSL (due to how the opaque sampler type gets re-interpreted).

In Unreal our sampler API looks like this (tweaked slightly to map to the equivalent concepts like "domains" in OpenQMC). The sampler API is purely stateless, you have to make a new domain if you want your sampler to return something new.

#define SAMPLER_RANDOM    0
#define SAMPLER_LATTICE   1
#define SAMPLER_SOBOL     2
#define SAMPLER_PMJ       3
// more sampler types here ...

// if user did not choose, default to best sampler
#ifndef SAMPLER_TYPE
#define SAMPLER_TYPE SAMPLER_SOBOL
#endif

struct Sampler {
    // The internals could be opaque, but keeping the type concrete is unavoidable in HLSL.
    // There is no requirement to guarantee the size or layout as the library is header only.
    // The type is small enough to be passed around by value.
    uint Index, Seed;
};

Sampler Create(uint3 PixelAndFrame, uint SampleIndex); // progressive sampler
Sampler Create(uint3 PixelAndFrame, uint SampleIndex, uint SampleCount); // fixed sample count sampler

// Domain update
Sampler DomainCreate(Sampler base, uint Key);
Sampler DomainSplit(Sampler base, uint Key, uint Size, uint Index);
Sampler DomainDistrib(Sampler base, uint Key, uint Index);
Sampler DomainChain(Sampler base, uint Key, uint Index);

// Draw samples, either the raw bits, or floats in [0,1):
uint  GetBits1D(Sampler s);
uint2 GetBits2D(Sampler s);
uint3 GetBits3D(Sampler s);
uint4 GetBits4D(Sampler s);

float  Get1D(Sampler s);
float2 Get2D(Sampler s);
float3 Get3D(Sampler s);
float4 Get4D(Sampler s);

The whole library can be a single header (say openqmc.h). Instead of picking at runtime, you decide before you include openqmc.h which sampler you want. If you want to compare samplers (something you probably only do once when first integrating the library), you just rebuild the application. In any case, since the code is all header only, keeping the sizeof() the core Sampler type fixed is not a hard requirement, though we could certainly make the effort to do this to keep things more consistent).

The last missing bit is how to upload the GPU resource. Like I mentioned in the other thread, I think all that openqmc should provide is a pointer to the CPU data and leave it up to the user to upload to the GPU with the API they are using (CUDA, Vulkan, Metal, etc ...). Since its a global resource, we could give it a predictable name/size (depending on the compile-time switch) so that the user code doesn't need to change when flipping between samplers.

Maybe something like:

#define CACHE_TYPE uin16_t
#define CACHE_SIZE 65536 // or whatever makes sense
extern CACHE_TYPE cache_data[CACHE_SIZE];

On the CPU, no extra code is required (can directly access cache_data[]), on the GPU you need something like:

__global__ CACHE_TYPE* cache_data; // cuda
Buffer<uint> cache_data; // HLSL
/// etc ...

to define the resource, but then the actual access into the array looks the same.

[joshbainbridge]

Thanks @fpsunflower. I've been thinking about this while on holiday. These are good suggestions. Before going through some of the points individually, I've updated the specification to align somewhat more with what you describe:

• Option A: This is still very simple, but now with no more macros. When compiling the OpenQMC library the sampler is chosen.
• Option B: This has the same API and data structure as Option A. When compiling the client application the sampler is chosen. Uses a vtable.

What is nice is that we can even start with A, and then add B later as an opt-in feature without breaking the API or ABI.

Now to address a few of your points. I would 100% agree that we should expect application developers to pick a sampler at compile time. There is little practical use in making selection an option that is exposed to end users.

That said, Option B does allow selecting a sampler at runtime. However, that is by chance, and a by-product of providing a single sampler API (not per-sampler) that also does not require the liboqmc.so to be recompiled. This could be useful for, say, an OpenQMC package that is installed via a package manager.

This all hinges on my initial assumption for the C-ABI, that this was going to be the same C++ backend, but with C bindings. As a result, I do not think it can be just a header file as that would expose the C++ code to the calling code, which may not be able to parse C++.

Perhaps I'm wrong there and we could just make this all header-only? Would that prevent some users from using the C-ABI?

If it were just header-only, then the design you pointed out would be exactly what we want. But Option A is close to this. It still allows a developer to recompile and test different samplers. But the define is set as a compiler flag of OpenQMC. That would work for C, Rust, and Zig.

Looking forward, if we were to start to support shading languages like GLSL or HLSL, it would be nice if this was a compatible API. Although for us to make that step, perhaps the biggest challenge would be supporting new (multiple) backend implementations (non C++).

Lastly, on the subject of the cache memory. The current design is intended to allow the caller to manage the memory. The API provides the size of the cache and a function to initialise a block of memory on the CPU. Once that is initialised, then the caller can copy that block of memory to a new location. For example, when using CUDA, the caller is expected to memcpy that data from host to device. Then the pointer to the device-side allocation is passed into the sampler constructor.

It isn't super clear, but in the 'trace' example program, we allocate the memory here. And that routes through to cudaMallocManaged(). This then does the transfer from host to device automatically as a convenience. But making two allocations, one on host and another on device, and using cudaMemcpy after calling initialiseCache() would be the long-form way of writing it out.

This is all needed as we don't always statically store all the data for a sampler. It is up to the sampler implementation. But for example, the pmj samplers generate the data live when initialiseCache() is called. So the CPU-side data block needs to be initialised as a sampler could (but does not have to) do some initialisation work.

Apologies for the many words. Lots of good discussion points though. I think we are getting close to the final design and perhaps the crux of the decisions that need to be made.