igneum/proto-cuda/emu/cuda_runtime.h

// CPU emulation shim of the small slice of the CUDA runtime API that host.cu and the generated
// kernel.cu use. Compile-and-semantics check only, on a Mac without nvcc. Not part of the deliverable.
// Kernel launches: the text "name<<<grid, block>>>(" is rewritten by sed into "emu_launch(name, grid, block, ".
// Each emulated block runs `block` host threads; the 32 lanes of a warp synchronise inside __shfl_xor_sync.
#pragma once
#include <cstdint>
#include <cstdlib>
#include <cstring>
#include <cstdio>
#include <chrono>
#include <thread>
#include <vector>
#include <memory>
#include <mutex>
#include <condition_variable>

#define __device__
#define __global__
#define __host__
#define __forceinline__ inline

struct uint3 { unsigned x, y, z; };
// uint4 and make_uint4 (read-width experiment, 5 October 2026): the wide loads and the scratch slots are 16-byte vectors.
struct uint4 { unsigned x, y, z, w; };
static inline uint4 make_uint4(unsigned x, unsigned y, unsigned z, unsigned w) { uint4 v; v.x = x; v.y = y; v.z = z; v.w = w; return v; }
struct dim3 { unsigned x, y, z; dim3(unsigned x_ = 1, unsigned y_ = 1, unsigned z_ = 1) : x(x_), y(y_), z(z_) {} };
extern thread_local uint3 threadIdx;
extern thread_local uint3 blockIdx;
extern thread_local dim3 blockDim;
extern thread_local dim3 gridDim;

enum cudaError_t { cudaSuccess = 0, cudaErrorInvalidValue = 1, cudaErrorMemoryAllocation = 2 };
enum cudaMemcpyKind { cudaMemcpyHostToHost = 0, cudaMemcpyHostToDevice = 1, cudaMemcpyDeviceToHost = 2, cudaMemcpyDeviceToDevice = 3 };
enum cudaDeviceAttr { cudaDevAttrWarpSize = 10, cudaDevAttrClockRate = 13, cudaDevAttrMemoryClockRate = 36,
                      cudaDevAttrGlobalMemoryBusWidth = 37, cudaDevAttrL2CacheSize = 38, cudaDevAttrMaxThreadsPerMultiProcessor = 39 };
struct cudaDeviceProp { char name[256]; size_t totalGlobalMem; int multiProcessorCount; int major, minor; };
struct cudaFuncAttributes { int numRegs; };
struct cudaEventRec { std::chrono::steady_clock::time_point t; };
typedef cudaEventRec* cudaEvent_t;

inline const char* cudaGetErrorString(cudaError_t e) { return e == cudaSuccess ? "no error" : "emulated error"; }
inline cudaError_t cudaGetLastError() { return cudaSuccess; }
inline cudaError_t cudaMalloc(void** p, size_t n) { *p = std::malloc(n); return *p ? cudaSuccess : cudaErrorMemoryAllocation; }
inline cudaError_t cudaFree(void* p) { std::free(p); return cudaSuccess; }
inline cudaError_t cudaMemcpy(void* d, const void* s, size_t n, cudaMemcpyKind) { std::memcpy(d, s, n); return cudaSuccess; }
inline cudaError_t cudaMemGetInfo(size_t* f, size_t* t) { *f = 8ull << 30; *t = 16ull << 30; return cudaSuccess; }
inline cudaError_t cudaEventCreate(cudaEvent_t* e) { *e = new cudaEventRec(); return cudaSuccess; }
inline cudaError_t cudaEventRecord(cudaEvent_t e, int = 0) { e->t = std::chrono::steady_clock::now(); return cudaSuccess; }
inline cudaError_t cudaEventSynchronize(cudaEvent_t) { return cudaSuccess; }
inline cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t a, cudaEvent_t b) { *ms = std::chrono::duration<float, std::milli>(b->t - a->t).count(); return cudaSuccess; }
inline cudaError_t cudaEventDestroy(cudaEvent_t e) { delete e; return cudaSuccess; }
inline cudaError_t cudaDeviceSynchronize() { return cudaSuccess; }
inline cudaError_t cudaGetDeviceCount(int* c) { *c = 1; return cudaSuccess; }
inline cudaError_t cudaSetDevice(int) { return cudaSuccess; }
inline cudaError_t cudaGetDeviceProperties(cudaDeviceProp* p, int) {
    std::strcpy(p->name, "CPU emulation shim (not a GPU)"); p->totalGlobalMem = 16ull << 30; p->multiProcessorCount = 0; p->major = 0; p->minor = 0; return cudaSuccess;
}
inline cudaError_t cudaDriverGetVersion(int* v) { *v = 0; return cudaSuccess; }
inline cudaError_t cudaRuntimeGetVersion(int* v) { *v = 0; return cudaSuccess; }
inline cudaError_t cudaDeviceGetAttribute(int* v, cudaDeviceAttr a, int) { *v = (a == cudaDevAttrWarpSize) ? 32 : 0; return cudaSuccess; }
template<class T> cudaError_t cudaFuncGetAttributes(cudaFuncAttributes* a, T*) { a->numRegs = 0; return cudaSuccess; }
template<class T> cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor(int* n, T*, int, size_t) { *n = 0; return cudaSuccess; }

// Device intrinsics
inline unsigned int __umulhi(unsigned int a, unsigned int b) { return (unsigned int)(((uint64_t)a * (uint64_t)b) >> 32); }
unsigned int __shfl_xor_sync(unsigned int mask, unsigned int v, int laneMask, int width = 32);
unsigned int __shfl_sync(unsigned int mask, unsigned int v, int srcLane, int width = 32);   // wide loads (lever b) broadcast lane 0

// Warp emulation
struct EmuWarp {
    std::mutex m;
    std::condition_variable cv;
    unsigned arrived = 0, generation = 0, size = 32;
    uint32_t slot[32];
};
extern thread_local EmuWarp* emu_current_warp;

template<class F, class... A> void emu_launch(F f, unsigned grid, unsigned block, A... args) {
    unsigned warps = (block + 31u) / 32u;
    std::vector<std::unique_ptr<EmuWarp>> warpObjs;
    for (unsigned w = 0; w < warps; ++w) {
        warpObjs.emplace_back(new EmuWarp());
        warpObjs.back()->size = (block - 32u * w) < 32u ? (block - 32u * w) : 32u;
    }
    std::vector<std::thread> ts;
    for (unsigned t = 0; t < block; ++t) {
        EmuWarp* w = warpObjs[t / 32u].get();
        ts.emplace_back([=]() {
            threadIdx = uint3{t, 0u, 0u};
            blockDim = dim3(block);
            gridDim = dim3(grid);
            emu_current_warp = w;
            for (unsigned b = 0; b < grid; ++b) { blockIdx = uint3{b, 0u, 0u}; f(args...); }
        });
    }
    for (auto& th : ts) th.join();
}