The Ramanujan Platform aims to harness the idle computational power of various digital devices to create a powerful distributed network. This platform allows any device with basic arithmetic capabilities and internet connectivity to contribute to the network's computational power.
20
stars
64
commits
Java
primary language
Sep 5, 2026
updated
This project aims to utilize the untapped computation power of digital devices. Any device that has a CPU which can do basic arithmetic operations and can connect to the internet contribute to the computation power of the network.
Apollo guidance computer had a CPU which was only as powerful as a modern scientific calculator. Modern smartphones and smart TVs are at-least million times more powerful than the Apollo guidance computer, and most of the time they are idle. This project aims to utilize the idle computation power of these devices.
Any device which has a CPU that can do basic arithmetic operations and can connect to the internet can be added on the network. This includes smartphones, smart TVs, smart-watches, smart speakers, smart refrigerators, smart washing machines, and so on. Currently, the network supports only Android devices, linux devices, Windows devices, and macOS devices.
BOINC is an amazing platform that allows users to donate their computation power to various scientific projects. However, for adding a new kind of computation on the network, the project owner has to write a new client for the BOINC network, and the devices on the network that can contribute to the computation have to download and install the new client.
In this platform, for any new kind of computation, the project owner doesn't have to write a new client. The owner just
have to give the computation code in the ramanujan language. The platform would convert the code to an intermediate
code that the device can just run. The devices on the network don't have to download and install a new client for every
new kind of computation. They just have to download and install the ramanujan client one-time.
Variables in ramanujan are declared using the var keyword. Like in most programming languages, the type of the
variable has to be explicitly given. Currently, the supported types are integer, double.
Example:
var x:integer;
var y:double;
Arrays in ramanujan are declared using the var keyword. The size of the array be both variable and explicitly given. This is a special kind of variable and the data type has to be given as array.
Example:
var arr[100]:array;
var n : double;
n =10;
var arr1[n]: array
Functions in ramanujan are declared using the def keyword. The functions do not have any return type. The parameters
of the functions are both input and output. The parameters have to be declared with the var keyword. The parameters
are passed by reference.
To invoke a function, the exec keyword is used.
Example:
def add(var x:integer, var y:integer, var ans:integer) {
ans = x + y;
}
exec add(1, 2, ans);
Here, the value of x is 1, the value of y is 2, and the value of ans NULL when passed to the function. But, since
the arguments are passed by reference, the value of ans is 3 after the function is executed.
ramanujan supports while loops. The syntax is similar to most programming languages.
Example:
var i:integer;
i = 0;
while(i < 10) {
i = i + 1;
}
ramanujan supports if-else statements. The syntax is similar to most programming languages.
Example:
var x:integer;
x = 10;
if(x < 10) {
x = 0;
} else {
x = 1;
}
There are two methods in Ramanujan which need to be used to distribute the computation over multiple devices:
threadStart:
The threadStart keyword is used to define a new thread. Any new variable can be introduced in the thread, or it
can use the variables defined outside the thread. in global scope The thread can call any function defined in the code.
Syntax:
threadStart(t0) {
// code
}
Here, t0 is the name of the thread. The thread name has to be unique in the code.
Example:
threadStart(t0) {
exec add(1, 2, ans);
}
threadOnEnd:
The threadOnEnd keyword is used to define actions that should be taken when one or more threads complete.
The body code executes only on the last iteration.
Syntax:
threadOnEnd(thread_seperated_thread_names, number_of_iterations) {
// code
}
Example:
threadOnEnd(t0, t1, 5) {
// code
}
Here, t0, t1 are the names of the threads. The last argument
5 is the number of iterations the threads have to run. The code defined in threadOnEnd for the n-1 times (here 4 times)
would spawn the threads again. On the n-th time, the code defined in threadOnEnd would be executed.
T1 T1 T1 T1 T1
/ \ / \ / \ / \ / \
X Y Y Y Y Z
\ / \ / \ / \ / \ /
T0 T0 T0 T0 T0
Here, In Y nodes, its just joining the thread and would do nothing, but on the last iteration, it would execute the code
defined in threadOnEnd as node Z.
threadParallelismCycle:
The threadParallelismCycle keyword is used to define actions that should be taken after every cycle of parallelism.
Unlike threadOnEnd, which only executes its body on the final iteration, threadParallelismCycle executes its body
after each and every completed cycle (including the last one).
Syntax:
threadParallelismCycle(thread_names, number_of_iterations) {
// code – runs after every cycle
}
Example:
threadParallelismCycle(t0, t1, 5) {
// this block runs after each of the 5 cycles
}
The DAG structure for threadParallelismCycle(t0, t1, 3) { body }:
T0 T0 T0
/ \ / \ / \
X body body body (final)
\ / \ / \ /
T1 T1 T1
Each body node runs after every pair of T0+T1 threads completes. After each body (except the last), the threads
are re-spawned for the next cycle.
Important: Do not combine
threadParallelismCycleandthreadOnEndon the same thread set.threadParallelismCyclealready runs its body after the final cycle too, sothreadOnEndis redundant. When added on the same threads,threadOnEndinjects an extra empty join node as a second successor of every first-cycle thread (instead of the expected single cycle-body successor), and overwrites the internal re-spawn map entries thatthreadParallelismCycleset up. This breaks the cycle chain and creates an unintended parallel execution path.If you want a dedicated block that only runs after the last cycle, use
threadOnEndalone (withoutthreadParallelismCycle). If you want code that runs after every cycle (including the last), usethreadParallelismCyclealone.
Example of complex threading:
T2 T3
/ \ / \
N0 N1 N2
\ / \ /
T1 T4
Following code would help here:
threadStart(t1) {
}
threadStart(t2) {
}
threadOnEnd(t1, t2, 1) {
threadStart(t3) {
}
threadStart(t4) {
}
}
threadOnEnd(t3, t4, 1) {
}
Example of a distributed gradient descent with Particle Swarm Optimization in ramanujan language:
def getSquared(var xPow:integer, var yPow:integer, var ans:integer) {
if(xPow < yPow) {
ans = yPow - xPow;
} else {
ans = xPow - yPow;
}
}
def getAvg(var arr:array, var originalArr:array, var avgF:integer) {
var index,ans1,tmpAvg1,tmpAvg2:integer;
avgF = 0;
index = 0;
while(index < 100) {
tmpAvg1 = arr[index];
tmpAvg2 = originalArr[index];
exec getSquared(tmpAvg1,tmpAvg2, ans1);
avgF = avgF + ans1;
index = index + 1;
}
avgF = avgF / 100;
}
def getTestArr(var xTest:integer, var yTest:integer, var testArrTest:array) {
var it:integer;
it = 0;
while(it < 100) {
testArrTest[it] = xTest * it + yTest;
it = it + 1;
}
}
var train[100]:array;
var i:integer;
i = 0;
while(i < 100) {
train[i] = i * 1.9 + 33;
i = i + 1;
}
def mainCode(var train : array, var x1:double, var y1:double) {
var x1,y1,j,avg,diff1,diff2x,diff2y,tmp:double;
j = 0;
var testArr[100]:array;
var slope:double;
var nexty,nextx:double;
testArr[1] = 1;
while(j < 15000) {
exec getTestArr(x1,y1,testArr);
exec getAvg(testArr, train, diff1);
tmp = x1 + 0.0001;
exec getTestArr(tmp,y1,testArr);
exec getAvg(testArr, train, diff2x);
slope = (diff2x - diff1) / 0.0001;
nextx = x1 - slope * 0.1;
tmp = y1 + 0.0001;
exec getTestArr(x1,tmp,testArr);
exec getAvg(testArr, train, diff2y);
slope = (diff2y - diff1) / 0.0001;
nexty = y1 - slope * 0.50;
x1 = nextx;
y1 = nexty;
j = j + 1;
}
}
var x1[100][10],y1[100][10]:array;
x1[0][0] = 0;
y1[0][0] = 0;
var ansX1,ansy1 :double;
ansX1 = 0;
ansy1 = 0;
var iteration[10]:array;
i = 0;
while(i < 10) {
iteration[i] = 0;
i = i + 1;
}
def getBest(var train:array, var best:integer, var x1:array, var y1:array, var iteration:integer) {
best = 0;
var index:integer;
var bestM:double;
bestM = 1000000000;
index = 0;
while(index < 10) {
var testArr[100]:array;
testArr[0] = 0;
var testX1,testY1:double;
testX1 = x1[index][iteration];
testY1 = y1[index][iteration];
exec getTestArr(testX1,testY1,testArr);
var avg:double;
avg = 0;
exec getAvg(testArr, train, avg);
if(avg < bestM) {
bestM = avg;
best = index;
}
index = index + 1;
}
}
def posRun(var thread:integer, var train:array, var x1:array, var y1:array, var iteration :array) {
var currentIter:integer;
currentIter = iteration[thread];
if(currentIter == 0) {
x1[thread][currentIter]=thread;
y1[thread][currentIter]=thread;
} else {
var best :integer;
best=0;
var thisIter:integer;
thisIter=currentIter;
currentIter = currentIter-1;
exec getBest(train, best, x1, y1, currentIter);
if(x1[thread][currentIter] < x1[best][currentIter]) {
x1[thread][thisIter] = x1[thread][currentIter]+(x1[best][currentIter]-x1[thread][currentIter])/2;
} else {
x1[thread][thisIter] = x1[thread][currentIter]-(x1[thread][currentIter]-x1[best][currentIter])/2;
}
if(y1[thread][currentIter] < y1[best][currentIter]) {
y1[thread][thisIter] = y1[thread][currentIter]+(y1[best][currentIter]-y1[thread][currentIter])/2;
} else {
y1[thread][thisIter] = y1[thread][currentIter]-(y1[thread][currentIter]-y1[best][currentIter])/2;
}
currentIter=thisIter;
}
var x,y:double;
x=x1[thread][currentIter];
y=y1[thread][currentIter];
exec mainCode(train, x, y);
x1[thread][currentIter]=x;
y1[thread][currentIter]=y;
}
threadStart(t0) {
exec posRun(0, train, x1, y1,iteration);
iteration[0]=iteration[0]+1;
}
threadStart(t1) {
exec posRun(1, train, x1, y1,iteration);
iteration[1]=iteration[1]+1;
}
threadStart(t2) {
exec posRun(2, train, x1, y1,iteration);
iteration[2]=iteration[2]+1;
}
threadStart(t3) {
exec posRun(3, train, x1, y1,iteration);
iteration[3]=iteration[3]+1;
}
threadStart(t4) {
exec posRun(4, train, x1, y1,iteration);
iteration[4]=iteration[4]+1;
}
threadStart(t5) {
exec posRun(5, train, x1, y1,iteration);
iteration[5]=iteration[5]+1;
}
threadStart(t6) {
exec posRun(6, train, x1, y1,iteration);
iteration[6]=iteration[6]+1;
}
threadStart(t7) {
exec posRun(7, train, x1, y1,iteration);
iteration[7]=iteration[7]+1;
}
threadStart(t8) {
exec posRun(8, train, x1, y1,iteration);
iteration[8]=iteration[8]+1;
}
threadStart(t9) {
exec posRun(9, train, x1, y1,iteration);
iteration[9]=iteration[9]+1;
}
threadOnEnd(t0,t1,t2,t3,t4,t5,t6,t7,t8,t9,5) {
var best:integer;
best=0;
exec getBest(train, best, x1, y1, 4);
ansX1=x1[best][4];
ansy1=y1[best][4];
}
The ramanujan language is faster than the Python3 language. The above single node code runs in ~350 ms. The same heuristic
in Python3 takes ~410 ms. The device it was tested on was a MacBook Air M3 : 8GB RAM, Apple M3 chip.
Ramanujan now supports a subset of Python syntax. Python code is converted to Ramanujan's intermediate representation using Python's AST module.
The conversion from Python code to Ramanujan's intermediate code follows this flow:
Python Source Code (.py file)
↓
[PythonAstInvoker]
Writes code to temp file, invokes: python3 -c "import ast, ast2json, json; ..."
↓
Python's ast.parse() generates AST
↓
ast2json.ast2json() converts AST to JSON
↓
JSON string returned to Java
↓
[JsonAstParser]
Parses JSON into Java AST node objects (ModuleNode, AssignNode, IfNode, etc.)
↓
[PythonAstToRuleEngineInputConverter]
Traverses AST nodes and generates RuleEngineInput structures
(Variables, Operations, Commands, Conditions, etc.)
↓
RuleEngineInput (Ramanujan Intermediate Code)
↓
Execution by Ramanujan Engine
PythonAstInvoker: Creates a temporary .py file, then runs a Python snippet that:
ast, ast2json, and jsonast.parse()ast2json.ast2json()JsonAstParser: Uses Jackson to parse the JSON and creates corresponding Java AST node objects.
PythonAstToRuleEngineInputConverter: Walks the AST tree and generates:
Variable objects for variable declarationsArray objects for list/array declarationsOperation objects for arithmetic and assignmentsCondition objects for comparisonsCommand objects that link everything in execution orderIf/While blocks for control flowFunctionCall objects for function invocationsVariables are automatically declared on first assignment. Type is inferred from the assigned value:
x = 5y = 3.14 or any arithmetic operation resultx = 5 # Integer
y = 3.14 # Double
z = x + y # Double (arithmetic result)
Arrays must be created using list comprehensions initialized with 0. This is the only supported form:
# 1D array initialized with zeros
arr = [0 for _ in range(100)]
# 2D array (matrix)
matrix = [[0 for _ in range(10)] for _ in range(10)]
# n-dimensional arrays with nested comprehensions
tensor = [[[0 for _ in range(5)] for _ in range(10)] for _ in range(3)]
# Variable dimensions are supported
n = 100
arr = [0 for _ in range(n)]
# Array element access and assignment
arr[0] = 10
value = arr[i]
matrix[x][y] = 5
Note: Arrays cannot be initialized with explicit values like [1, 2, 3] - only [0 for _ in range(n)] form is allowed.
Supported operators: +, -, *, /
result = a + b * c - d / e
x += 5 # x = x + 5
y -= 3 # y = y - 3
z *= 2 # z = z * 2
w /= 4 # w = w / 4
Supported: <, <=, >, >=, ==, !=
if x > 5:
pass
while count <= 100:
pass
if x > 10:
y = 1
else:
y = 0
Note: elif is not supported. Use nested if-else instead:
# Instead of elif, use nested if-else:
if x > 10:
y = 1
else:
if x > 5:
y = 2
else:
y = 0
i = 0
while i < 100:
# loop body
i += 1
Functions are defined using def and called directly. Arguments are passed by reference.
def calculate(a, b, result):
result = a + b * 2
# Function call
calculate(x, y, answer)
Functions can return values using tuple unpacking:
def get_coords():
x = 10
y = 20
return x, y
# Tuple unpacking to receive return values
a, b = get_coords()
Single return values:
def compute(x):
result = x * 2
return result
value = compute(5)
Returning an array element directly is NOT supported:
# NOT SUPPORTED
def get_element(arr, index):
return arr[index] # ❌ Cannot return array element directly
Workaround: Assign to a variable first, then return:
# SUPPORTED
def get_element(arr, index):
value = arr[index]
return value # ✓ Return variable works
Passing a function call as an argument to another function is NOT supported:
# NOT SUPPORTED
result = outer_func(inner_func(x)) # ❌ Nested function calls not allowed
Workaround: Use intermediate variables:
# SUPPORTED
temp = inner_func(x)
result = outer_func(temp) # ✓ Works with intermediate variable
Returning the result of a function call directly is NOT supported:
# NOT SUPPORTED
def wrapper(x):
return compute(x) # ❌ Cannot return function call directly
Workaround: Assign to a variable first:
# SUPPORTED
def wrapper(x):
result = compute(x)
return result # ✓ Return variable works
Only simple function names are supported (no method chains or computed function references):
# NOT SUPPORTED
obj.method() # ❌ Method calls on objects
funcs[0]() # ❌ Function from array
getattr(obj, 'func')() # ❌ Dynamic function access
for loops are NOT currently supported. Use while loops instead:
# NOT SUPPORTED
for i in range(10): # ❌
pass
# SUPPORTED - Use while loop
i = 0
while i < 10: # ✓
i += 1
Object-oriented programming is not yet supported:
# NOT SUPPORTED
class MyClass: # ❌
pass
Boolean expressions with and/or/not are NOT supported:
# NOT SUPPORTED
if x > 5 and y < 10: # ❌
pass
if not flag: # ❌
pass
Workaround: Use nested if statements:
# SUPPORTED
if x > 5:
if y < 10: # ✓ Nested if for 'and' logic
pass
Dynamic list operations are not supported:
# NOT SUPPORTED
arr.append(5) # ❌ append
arr.pop() # ❌ pop
len(arr) # ❌ len function
arr + other # ❌ list concatenation
Strings are NOT supported:
# NOT SUPPORTED
s = "hello" # ❌ string assignment
s = "hello" + "world" # ❌ string concatenation
Importing modules is not supported:
# NOT SUPPORTED
import math # ❌
from collections import deque # ❌
Try-except blocks are not supported:
# NOT SUPPORTED
try: # ❌
pass
except:
pass
Generator expressions and iterator protocols are not supported.
Power (**) and modulo (%) operators are NOT currently supported:
# NOT SUPPORTED
x = 2 ** 10 # ❌ power operator
y = n % 10 # ❌ modulo operator
The elif keyword is NOT supported. Use nested if-else:
# NOT SUPPORTED
if x > 10:
y = 1
elif x > 5: # ❌ elif not allowed
y = 2
# SUPPORTED - Use nested if-else
if x > 10:
y = 1
else:
if x > 5: # ✓
y = 2
# Gradient descent example in Python syntax
def get_squared(x_pow, y_pow):
if x_pow < y_pow:
ans = y_pow - x_pow
else:
ans = x_pow - y_pow
return ans
def get_test_arr(x, y, test_arr):
it = 0
while it < 100:
test_arr[it] = x * it + y
it = it + 1
# Initialize training data
train = [0 for _ in range(100)]
i = 0
while i < 100:
train[i] = i * 1.9 + 33
i = i + 1
# Main computation
x1 = 0.0
y1 = 0.0
j = 0
test_arr = [0 for _ in range(100)]
while j < 1000:
get_test_arr(x1, y1, test_arr)
# ... gradient computation logic
j = j + 1
Ramanujan supports offloading compute-intensive work to the GPU via OpenCL. Any Python function
whose name matches the pattern funcName_GPU_N (where N is a positive integer) is automatically
compiled to an OpenCL C kernel during translation and dispatched to the GPU at runtime.
The function name encodes the number of NDRange dimensions:
def funcName_GPU_N(dataArg1, dataArg2, ..., dataArgK, rangeDim1, ..., rangeDimN):
# body – use rangeDimK as the per-work-item index
| Part | Role | What the translator does |
|---|---|---|
N in _GPU_N | Number of NDRange dimensions (work_dim) | Read from the function name; no extra argument needed |
dataArg1 … dataArgK | Data arrays (first total_params − N parameters) | Emitted as __global float* kernel parameters |
rangeDim1 … rangeDimN | Work-item index variables (last N parameters) | Emitted as int dim = get_global_id(K); declarations; their call-site values become global_work_size[] for clEnqueueNDRangeKernel |
Requirement: Data args must be arrays. The translator maps them to OpenCL buffers.
def vector_add_GPU_1(a, b, c, gid):
c[gid] = a[gid] + b[gid]
# 1024-element arrays, 1 NDRange dimension
vector_add_GPU_1(a, b, c, 1024)
Generated OpenCL kernel:
__kernel void vector_add(__global float* a, __global float* b, __global float* c) {
int gid = get_global_id(0);
c[gid] = (a[gid] + b[gid]);
}
def matrix_add_GPU_2(a, b, c, row, col):
c[row] = a[row] + b[col]
# 64 × 64 grid (2 NDRange dimensions)
matrix_add_GPU_2(a, b, c, 64, 64)
if/else and while inside the function body are translated to their C equivalents:
def relu_GPU_1(a, out, gid):
if a[gid] > 0:
out[gid] = a[gid]
else:
out[gid] = 0
relu_GPU_1(a, out, 512)
A GPU kernel can call ordinary (non-_GPU_N) Python functions that are defined in the same file.
The translator converts those helpers to OpenCL C device functions and prepends them before the
__kernel declaration so they are visible to the kernel body.
# Plain Python helper – becomes a float device function in the generated OpenCL C
def scale2(x):
r = x * 2
return r
# GPU kernel – calls the helper for every work-item
def apply_scale_GPU_1(a, out, gid):
v = a[gid] # load array element into a local variable first
out[gid] = scale2(v)
apply_scale_GPU_1(a, out, 1024)
Generated OpenCL C:
float scale2(float x) {
float r = (x * 2);
return r;
}
__kernel void apply_scale(__global float* a, __global float* out) {
int gid = get_global_id(0);
float v = a[gid];
out[gid] = scale2(v);
}
| Constraint | Detail |
|---|---|
| Parameter types | All helper parameters are treated as float scalars. Array pointers (__global float*) are not supported in helper functions. |
| Return type | Always float. |
| No subscript as argument | Array element expressions (a[i]) cannot be passed directly as arguments to a helper call. Assign the element to a local variable first: v = a[gid]; out[gid] = scale2(v). |
| No recursion | A function may not call itself. The translator throws a CompilationException if a recursive call is detected at translation time. |
| No GPU–kernel calls | A helper (or the kernel) may not call another _GPU_N function. GPU kernels are dispatched via clEnqueueNDRangeKernel and cannot be invoked as device functions. |
| Scope | Only top-level module functions are eligible as helpers. Nested function definitions are not supported. |
The translator enforces the no-recursion rule at translation time, not at runtime:
# INVALID – will raise CompilationException during translation
def bad_GPU_1(a, gid):
a[gid] = bad_GPU_1(a, gid) # ❌ recursive GPU call
# INVALID – helper self-recursion is also rejected
def factorial(n):
return n * factorial(n - 1) # ❌ recursive helper
def kernel_GPU_1(a, gid):
a[gid] = factorial(a[gid])
These special function names are recognised inside _GPU_N kernel bodies and
translated to OpenCL C intrinsics. They are not available on the host; use
the corresponding Ramanujan host built-ins (FLOOR, EXP, etc.) outside GPU
functions.
Each call mutates its argument in place: FUNC(x) → x = func(x) in the
generated OpenCL C.
| Python call | Generated OpenCL C | Notes |
|---|---|---|
EXP(x) | x = exp(x); | Natural exponential |
LOG(x) | x = log(x); | Natural logarithm |
SQRT(x) | x = sqrt(x); | Square root |
FLOOR(x) | x = floor(x); | Round toward −∞; result stays float |
FLOOR example — locate the grid cell for a particle position:
def p2g_GPU_1(positions, params, gid):
xp = positions[gid * 3]
gxp = (xp - (-1.0)) * 10.0 # map to grid coords
FLOOR(gxp) # gxp = floor(gxp) in OpenCL C
i0 = gxp # integer grid index (stored as float)
nibble = PACKED_NIBBLE(packed, index)
PACKED_NIBBLE numerically converts an exact integer-valued packed float to
OpenCL uint, shifts by index * 4, masks with 0xF, and returns the nibble as
a float. Valid nibble indexes are 0 through 5 for Ramanujan's six-weights-per-
float Phi-3 format. Packed values must not exceed 0xFFFFFF, which is exactly
representable in float32.
((float)(((uint)packed >> ((uint)index * 4u)) & 15u))
This avoids floating-point divide and floor() during 4-bit dequantization.
It is an expression intrinsic and may be used on the right-hand side of an
assignment inside _GPU_N functions.
ATOMIC_ADD_F(arr, idx, delta)
Atomically adds delta (a float) to arr[idx] using a compare-and-swap
loop. Required whenever multiple work-items may scatter into the same array
slot concurrently (e.g., particle-to-grid scattering in MPM).
OpenCL 1.2 has no native atomic_add for float. The translator emits a
CAS loop on the int-reinterpreted bits using atomic_cmpxchg:
/* Generated for: ATOMIC_ADD_F(arr, idx, delta) */
{
__global volatile int* _aAddr = (__global volatile int*)(&arr[(int)(idx)]);
int _aOld, _aNew;
do {
_aOld = *_aAddr;
_aNew = as_int(as_float(_aOld) + (delta));
} while (atomic_cmpxchg(_aAddr, _aOld, _aNew) != _aOld);
}
The {} block scope lets you call ATOMIC_ADD_F multiple times in the same
kernel without variable-name conflicts.
Requirements:
__global float* data argument (not a local variable).atomic_cmpxchg on __global int* is a
core 1.2 feature on all platforms including Apple Metal-backed OpenCL).ATOMIC_ADD_F example — particle-to-grid mass scatter:
def p2g_GPU_1(positions, g_mass, g_vel, params, gid):
# ... compute weight w and grid node index gnode ...
wm = w * params[0] # weighted mass
ATOMIC_ADD_F(g_mass, gnode, wm) # safe concurrent scatter
ATOMIC_ADD_F(g_vel, gnode * 3, wm * vx)
ATOMIC_ADD_F(g_vel, gnode * 3 + 1, wm * vy)
ATOMIC_ADD_F(g_vel, gnode * 3 + 2, wm * vz)
Note on P2G in MPM: The above P2G kernel also requires
FLOORto locate grid nodes (see above). BothFLOORandATOMIC_ADD_Fare needed before P2G can fully run on GPU.
| Platform | Requirement |
|---|---|
| macOS | OpenCL is part of the system framework – no extra install needed |
| Linux | sudo apt install ocl-icd-opencl-dev opencl-headers |
| Windows | Install GPU vendor drivers: NVIDIA CUDA Toolkit, AMD ROCm, or Intel OpenCL SDK |
cl_kernel.double → float before upload and float → double after read-back (OpenCL kernels operate on float).stderr and execution returns immediately.GPU_ENABLED macro must be set at compile time (via -DENABLE_GPU=ON). Builds without it contain no OpenCL code and have no OpenCL runtime dependency.GPU_SYNCBy default, GPU kernels dispatched with _GPU_N functions are non-blocking: the OpenCL
command is queued but the CPU continues immediately. This allows many kernel launches to
be batched together without the CPU stalling after every one — which is the key to high GPU
throughput.
However, any time the host (CPU) side needs to read back a value that a GPU kernel has
written, you must explicitly drain the GPU queue for that array first. The built-in
GPU_SYNC does exactly that.
GPU_SYNC(array)
GPU_SYNC(array) issues a blocking clEnqueueReadBuffer for the given array, flushing
all previously enqueued GPU work and copying the updated data back to the host buffer. It
is a no-op for arrays that are not GPU-backed (e.g., host-only arrays).
GPU_SYNC| Situation | Action |
|---|---|
Reading an array element in a Python host while loop after a GPU kernel has written it | Call GPU_SYNC(array) once before the loop |
Passing a GPU-written array to a host function or exec call | Call GPU_SYNC(array) before the call |
dump array /path after GPU kernel(s) wrote it | Call GPU_SYNC(array) before dump |
| Using a GPU-written array only as input to the next GPU kernel (no host read) | No GPU_SYNC needed — GPU→GPU is handled automatically |
# 120+ GPU kernels dispatched with no CPU stalls …
layernorm_GPU_1(hidden, ln_g, ln_b, h_ln, n_seq)
matmul_bias_GPU_2(h_ln, c_attn_w, c_attn_b, qkv, kp, n_seq, 2304)
# … more kernels …
matmul_bias_GPU_2(h_ff, c_fc_proj_w, c_fc_proj_b, h_out_buf, kp, n_seq, 768)
# CPU needs to read `hidden` and `h_out_buf` in the next loop → sync first
GPU_SYNC(hidden)
GPU_SYNC(h_out_buf)
_i = 0
while _i < n_seq * 768:
hidden[_i] = hidden[_i] + h_out_buf[_i]
_i = _i + 1
Without the GPU_SYNC calls, hidden and h_out_buf would still hold stale values from
before the last GPU kernels ran, producing silently wrong results.
RETURNBy default, after execution completes arrChangeMap() reports every array that was modified
(comparing each element against its pre-execution snapshot). For large models with hundreds of
weight arrays this means the JNI result map can contain far more data than the caller needs.
The RETURN built-in lets user code explicitly name the arrays that should appear in the result.
If RETURN is called at any point during execution, only the listed arrays are included in
arrChangeMap(); all other modified arrays are silently dropped. If RETURN is never called
the behaviour is unchanged — all modified arrays are returned.
RETURN(arr1, arr2, ..., arrN)
RETURN accepts any number of array arguments and marks each with an internal flag at runtime.
It is a no-op for the computation itself (it does not stop execution or modify any values).
hidden = [0 for _ in range(768)]
kv_cache = [0 for _ in range(4096)]
weights = [0 for _ in range(131072)] # large weight buffer
# ... kernel calls that write to all three arrays ...
# Only return the outputs the caller actually needs
RETURN(hidden, kv_cache)
Without RETURN, weights and every other modified array would be serialised back through JNI
even though the caller only needs hidden and kv_cache.
Script contains RETURN(...)? | What arrChangeMap() returns |
|---|---|
| No | All modified arrays with their changed indexes (unchanged default) |
| Yes | Only the listed arrays, still reporting only changed indexes |
RETURN vs GPU_SYNCGPU_SYNC and RETURN are independent and complementary:
| Purpose | When to use | |
|---|---|---|
GPU_SYNC(arr) | Flush the GPU queue and read arr back to the CPU | Before the host reads a GPU-written array |
RETURN(arr1, ...) | Filter which arrays are included in the final result | To reduce serialisation overhead when only a subset of arrays is needed |
RELEASE_MEM and LOAD_MEMEvery array a _GPU_N kernel touches gets its own GPU buffer, and buffers are never freed
automatically while the process is running. For large models (e.g. dozens of transformer
layers, each with several 4-bit packed weight/scale arrays plus K/V caches) this means all
per-layer buffers stay resident on the GPU simultaneously. On memory-constrained devices
(e.g. ~5 GB unified memory) this can exhaust GPU/unified memory and cause an OOM crash.
RELEASE_MEM and LOAD_MEM give user code explicit control over when a buffer actually
occupies GPU memory:
RELEASE_MEM(array) # Frees the array's GPU buffer immediately (host data is untouched)
LOAD_MEM(array) # (Re)allocates the GPU buffer and uploads the current host data
RELEASE_MEM(array) calls clReleaseMemObject on the array's buffer and clears it, so the
memory is returned to the driver right away. The host-side (CPU) copy of the array is never
touched, so the data itself is not lost.LOAD_MEM(array) allocates a fresh GPU buffer for the array from its current host data.
It is the correct counterpart to RELEASE_MEM — plain GPU_LOAD only writes into an
already-existing buffer, so it cannot bring back an array that has been released.Mandatory, not optional. A
_GPU_Nkernel dispatch performs no implicit buffer allocation or upload of its own — it only doesclSetKernelArg+clEnqueueNDRangeKernelagainst whatever GPU buffer the array already has. This means:
LOAD_MEM(array)is required before the first time an array is passed as a data argument to any_GPU_Nkernel, and again after anyRELEASE_MEM(array)on it. Skipping this is not silently "slow" — it is a hard error: the kernel logsmissing LOAD_MEM(array) before first useand the dispatch is skipped entirely (the kernel's outputs are left unchanged).GPU_LOAD(array)is required any time host (CPU) code mutates an array that already has a live GPU buffer (e.g.l_idx_arr[0] = l_idx) and the next kernel dispatch needs to see that new value.GPU_LOADonly writes into an existing buffer — it never allocates one — so it must not be used as a substitute forLOAD_MEMon an array that hasn't been loaded yet (or was just released).
RELEASE_MEM is intended for immutable arrays (4-bit packed weights, scale tables, RoPE
cos/sin tables, etc.). Never release an array that a GPU kernel writes to (activation
buffers, KV caches, etc.) without first reading its current value back with GPU_SYNC —
RELEASE_MEM does not sync, so releasing a GPU-resident write target discards whatever the
GPU hasn't yet written back to the host copy.
Release/reload has a cost — every call is buffer churn.
RELEASE_MEM+LOAD_MEMon the same array is aclReleaseMemObject+clCreateBufferpair, not a full data copy, for weight arrays uploaded withCL_MEM_USE_HOST_PTR(seeisBinaryLoadedin the "Direct CSV Population" section above) — on unified-memory devices (e.g. Apple Silicon) this is cheap since no bytes actually move. It is still real per-call overhead, though, so only cycle an array throughRELEASE_MEM/LOAD_MEMwhen you specifically need to cap peak GPU memory (e.g. to fit a multi-layer model on a memory-constrained device); don't do it reflexively for arrays that are about to be reused within the same layer or the same kernel dispatch.
If you don't need to cap peak GPU memory mid-run, the simplest use is to release every large buffer exactly once, right before the process is about to end (or before the very last use of a single-pass computation), instead of releasing/reloading inside a hot loop:
# ... 32 transformer layers, prefill + decode loop, all reusing the same
# per-layer weight/cache buffers on every iteration ...
# Generation is finished — nothing on the GPU is needed anymore before the
# host reads back `generated_tokens`. Release every large buffer once,
# instead of thrashing release/reload inside the decode loop.
RELEASE_MEM(l0_qkv_packed)
RELEASE_MEM(l0_qkv_scales)
RELEASE_MEM(l0_k_cache)
RELEASE_MEM(l0_v_cache)
# ... one RELEASE_MEM call per large buffer ...
GPU_SYNC(generated_tokens)
When peak GPU memory itself is the constraint (e.g. a multi-layer transformer that would
otherwise keep every layer's weights resident simultaneously), LOAD_MEM an immutable
layer's weights right before that layer's kernels run, and RELEASE_MEM them again
immediately after — in every place that layer is used, including inside a repeated decode
loop. Because weight arrays are read-only inputs (never a GPU write target) and are typically
isBinaryLoaded (zero-copy CL_MEM_USE_HOST_PTR), the repeated release/reload is just cheap
buffer-object churn, not a data copy, and at any instant only one layer's weights occupy GPU
memory:
# ── Layer N ──
LOAD_MEM(lN_qkv_packed)
LOAD_MEM(lN_qkv_scales)
LOAD_MEM(lN_o_packed)
LOAD_MEM(lN_o_scales)
LOAD_MEM(lN_gate_up_packed)
LOAD_MEM(lN_gate_up_scales)
LOAD_MEM(lN_down_packed)
LOAD_MEM(lN_down_scales)
rmsnorm_GPU_1(h_state, lN_ln1_g, h_ln1, n_seq)
matmul_4bit_GPU_2(h_ln1, lN_qkv_packed, lN_qkv_scales, qkv_buf, kp_qkv, n_seq, 9216)
# ... rest of layer N's kernels ...
residual_add_GPU_2(h_state, h_out_buf, n_seq, 3072)
RELEASE_MEM(lN_qkv_packed)
RELEASE_MEM(lN_qkv_scales)
RELEASE_MEM(lN_o_packed)
RELEASE_MEM(lN_o_scales)
RELEASE_MEM(lN_gate_up_packed)
RELEASE_MEM(lN_gate_up_scales)
RELEASE_MEM(lN_down_packed)
RELEASE_MEM(lN_down_scales)
Note that KV caches (lN_k_cache/lN_v_cache) are deliberately not part of this
per-layer cycle — they are write targets updated by every layer every step, so they must stay
GPU-resident across the whole run (they are only released once, at the very end, alongside the
other shared buffers per Pattern 1). See
ramanujan-test-codes/phi3/phi3_transformer_stack_4bit.py
for the full worked example across all 32 layers, in both the prefill pass and the decode loop.
GPU_SYNC vs the previous implicit sync modelBefore GPU_SYNC was introduced, every _GPU_N call automatically issued a blocking
clEnqueueReadBuffer + clFinish after the kernel, preventing any batching. The overhead
measured ~11% of total inference time on macOS (visible as IOKit → IOGPU → clFinish in
profiler traces). With the explicit model, the GPU queue is drained only at the
necessary points, and all other kernel dispatches remain asynchronous.
Ramanujan now supports a subset of Python syntax through AST-based conversion (see Python Support section above). The platform is actively evolving to support more Python features progressively.
Python support is being actively developed on the Ramanujan platform. More and more features are being added continuously to bring the full power of Python to distributed computing:
Coming Shortly: Object-Oriented Programming (OOP) support
Progressive Additions: We will progressively add all Python features to the Ramanujan platform, including:
and, or, not)**) and modulo (%) operatorsfor loops and iteratorselif statementstry/except)Long-term Vision: Full Python3 ecosystem compatibility
The goal is to make Python code seamlessly executable on the distributed Ramanujan platform while maintaining performance and enabling parallel computation across devices.
Use mvn clean install. Following is the dependency hierarchy:
cd ramanujan-native/native
mkdir build
cd build
cmake ..
cmake --build .
cd ramanujan-native/native
mkdir build-gpu
cd build-gpu
cmake -DENABLE_GPU=ON ..
cmake --build .
Passing -DENABLE_GPU=ON sets the GPU_ENABLED preprocessor macro, links OpenCL, and activates
OpenCL kernel dispatch for _GPU-suffixed functions. Standard builds (-DENABLE_GPU=OFF, the
default) compile no OpenCL code and have no dependency on any OpenCL runtime.
Dockerfile is provided to containerize all the necessary services.
prayog device server on the network:For executing code file:
java -jar <developer-console-path>/target/developer-console-1.0-SNAPSHOT-fat.jar execute <path-to-code-file>
The translation module (middleware-translation) currently requires the following Python dependencies to convert Python code to Ramanujan intermediate code:
pip install ast2jsonFor more information about third-party licenses, see THIRD_PARTY_LICENSES.md.
To set up the Ramanujan developer console and required dependencies, run the provided installer script:
# Make the installer executable
chmod +x install_ramanujan.sh
# Run the installer
./install_ramanujan.sh
RAMANUJAN_WS environment variable.libjsoncpp if needed.rj to your shell profile for easy usage.After installation, restart your terminal or run:
source ~/.zshrc # or ~/.bashrc, ~/.bash_profile, depending on your shell
You can now run the developer console with:
rj <path-to-code-file>

The code is submitted to DevConsole process which would be present in the path: developer-console/target/developer-console-1.0-SNAPSHOT-fat.jar.
The process would submit the code to the Middleware server, which would be responsible for converting the submitted code
to the intermediate code. The Middleware server would then work with the orchestrator server to process the required DAG.
Once the code gets converted to the intermediate code, the Middleware server would return back an asyncId which the
DevConsole would use to get the result of the code execution.

The Middleware server after creating the intermediate code, would start submitting the DAG-nodes to the orchestrator server. Whenever a DAG-node gets processed, the Middleware server would submit children - DAG elements to the orchestrator server. Once, the whole DAG is computed, the Middleware server sets the status of the whole processing as SUCCESS.
Pseudocode of the DAG computation is as follows:
DagElementQueue = new Queue();
DagElementQueue.add(rootDagElement);
isDone = false;
while(!isDone) {
asyncTasks = []
children = []
while(DagElementQueue is not empty) {
DagElement currentElement = DagElementQueue.poll();
if(currentElement is not processed) {
asyncTask = asyncCode(currentElement.process(), callback={
if(currentElement has children) {
for(DagElement child : currentElement.children) {
children.add(child);
}
}
})
asyncTasks.add(asyncTask);
}
}
for(asyncTask : asyncTasks) {
asyncTask.wait();
}
if(children is empty) {
isDone = true;
} else {
for(DagElement child : children) {
DagElementQueue.add(child);
}
}
}

Middleware end of the orchestrator:
The Middleware server would submit a computation task to the Orchestrator server. The Orchestrator server would instantly
return back an asyncId that the Middleware server would use to get the result of the computation. The Orchestrator server
would then check for an available device, and will assign the computation task to the device.
Device end of the orchestrator:
Any device-client that is available on the Ramanujan platform, would keep pinging to the Orchestrator server. Currently,
the pings contains the deviceId, but in future it would send in the device-stats as well. The Orchestrator server would
keep track of the devices that are available on the network. On the ping API request, the Orchestrator server would return
back a suitable computation task mapped to the device. The device would then execute the task and return back the result to the
Orchestrator server.
This service helps the system to resume in case the Middleware server goes down. This server becomes a part of a PubSub system. It can be part of Apache-Kafka, GCP-PubSub, and a local PubSub system. The producers and consumer in the server code can be extended to use any of the PubSub systems.
In addition to this, the Middleware servers are stateless and do not have the track of the asyncIds. The Kafka-Manager server keeps the track of the asyncIds. The messages on the PubSub queue contains the asyncId and the status of the computation. The consumer of the Kafka-Manager server would ping the Middleware server with the asyncId to get the result of the computation. This helps the Middleware server to move forward with the DAG computation.
The Ramanujan native interpreter (ramanujan-native) is the core execution engine responsible for running the intermediate code on devices. Written in C++ for maximum performance, it provides a sophisticated bytecode-like execution environment that can run on any device with minimal overhead.

The interpreter is designed with a two-phase architecture: first, it converts JSON-based intermediate representation into an optimized in-memory graph of execution objects, then it executes the command chain. This design separates parsing overhead from execution, enabling efficient repeated execution.
When the Processor.process() method is invoked, it receives:
The RuleEngineInput contains vectors of all program constructs:
| Component | Description |
|---|---|
variables | Scalar variables (integer, double) |
arrays | Multi-dimensional arrays |
constants | Immutable values |
commands | Execution units (each command wraps an operation, condition, function call, etc.) |
operations | Arithmetic and assignment operations (+, -, *, /, =) |
conditions | Comparison and logical operations (<, >, ==, !=, &&, ||) |
ifBlocks | Conditional branching structures |
whileBlocks | Loop structures |
functionCalls | Function definitions and call sites |
ID Map Creation:
The createMap() method iterates through all input elements and calls getInternalAnalogy() on each. This factory method converts input objects (parsed from JSON) into Rule Engine (RE) objects optimized for execution:
Variable → VariableRE
Array → ArrayRE
Command → CommandRE
Operation → OperationRE
Condition → ConditionRE
If → IfRE
While → WhileRE
FunctionCall → FunctionCallRE
The resulting unordered_map<string, RuleEngineInputUnits*> allows O(1) lookup of any program element by its unique ID.
After creating the ID map, references between objects must be resolved. During JSON parsing, relationships are stored as string IDs (e.g., a command's nextId field contains the ID of the next command). The graph fixing phase converts these string references to actual object pointers.
fixGraph(map):
Iterates through all RE objects and calls setFields(map) on each. This method:
Example for CommandRE.setFields():
nextCommandRE = dynamic_cast<CommandRE*>(getFromMap(map, command->nextId));
operationCommand = dynamic_cast<OperationRE*>(getFromMap(map, command->operation));
ifCommandRE = dynamic_cast<IfRE*>(getFromMap(map, command->ifBlocks));
whileCommandRE = dynamic_cast<WhileRE*>(getFromMap(map, command->whileId));
fixOperator(map, operations):
For each operation, creates a CachedOperationFunctioning object. This is a Strategy Pattern implementation where the appropriate operator implementation is selected once and cached:
| Operator | Implementation Class |
|---|---|
+ | AddImpl |
- | MinusImpl |
* | MultiplyImpl |
/ | DivideImpl |
= | AssignImpl |
The cached functor stores direct pointers to operands, eliminating lookup overhead during execution.
fixConditions(map, conditions):
Similarly creates CachedConditionFunctioning objects for conditions:
| Condition | Implementation Class |
|---|---|
< | LessThanImpl |
> | GreaterThanImpl |
<= | LessThanEqualToImpl |
>= | GreaterThanEqualToImpl |
== | IsEqualImpl |
!= | NotEqualImpl |
&& | AndImpl |
|| | OrImpl |
Before execution begins, the interpreter saves the original values of all variables and array elements:
for(RuleEngineInputUnits* variable : variableREs) {
VariableRE* variableRE = (VariableRE*)(variable);
dataFieldOriginalData.insert(make_pair(variableRE->getValPtrPtr(), *variableRE->getValPtrPtr()));
}
This enables the varChangeMap() and arrChangeMap() methods to efficiently compute which values changed during execution—essential for returning results to the calling system.
The main execution loop is elegantly simple:
CommandRE* command = dynamic_cast<CommandRE*>(mapBetweenIdAndRuleInput->at(firstCommandId));
while(command != nullptr) {
command = command->get();
}
Each CommandRE.get() method:
process() on that unitnextCommandRE)Command Types and Their Execution:
void OperationRE::process() {
operationFunctioning->set(); // Execute cached operation
}
The CachedOperationFunctioning directly accesses operand values and stores the result, with no string lookups or type checking at runtime.
void IfRE::process() {
CommandRE* commandRE;
if(conditionFunctioning->operate()) {
commandRE = ifCommandRE; // Execute if-block
} else {
commandRE = elseCommandRE; // Execute else-block
}
while(commandRE != nullptr) {
commandRE = commandRE->get(); // Nested command loop
}
}
Conditional blocks spawn a nested execution loop for their body commands.
void WhileRE::process() {
while(conditionFunctioning->operate()) {
CommandRE* commandRE = whileCommandRE;
while(commandRE != nullptr) {
commandRE = commandRE->get(); // Execute loop body
}
}
}
The outer while loop evaluates the condition; the inner loop executes the body commands.
Function execution is the most complex operation, involving context management for proper recursion support.
The FunctionCommandRE handles both user-defined functions and built-in functions.
Built-in Functions:
A factory function GetFunctionCommandRE() checks the function ID and returns optimized implementations:
| Function | Description |
|---|---|
NINF | Set to negative infinity |
PINF | Set to positive infinity |
RAND | Random number [0, 1) |
ABS | Absolute value |
SIN, COS, TAN | Trigonometric functions |
ASIN, ACOS, ATAN | Inverse trigonometric |
FLOOR, CEIL | Rounding functions |
EXP | Exponential (e^x) |
SQRT | Square root |
POW | Power function |
Built-in functions have minimal overhead—they directly access argument values and apply the operation.
User-Defined Function Execution:
The FunctionCommandRE.process() method orchestrates function calls through these steps:
Parameter Setup (Call-by-Reference)
Local Variable Preservation
Function Body Execution
command = firstCommand;
while(command != nullptr) {
command = command->get();
}
Context Restoration & Result Propagation
Memory Management:
The DataContainerValueFunctionCommandREMemMaintainer provides efficient memory pooling for function call contexts, avoiding repeated heap allocations during recursive calls.
After execution completes, two methods collect changed values:
varChangeMap():
for(RuleEngineInputUnits* variableRE1 : variableREs) {
VariableRE* variableRE = (VariableRE*)variableRE1;
double newVal = *variableRE->getValPtrPtr();
varChangeMap->insert(make_pair(variableRE->id, newVal));
}
arrChangeMap():
Compares current array values against saved originals and returns only changed elements with their indices.
The interpreter uses a unified data container system:
DataContainerValue (abstract)
├── DoublePtr // Scalar variable value
└── ArrayDataContainerValue
└── ArrayValue // Multi-dimensional array storage
├── val[] // Flat double array
├── totalSize // Total element count
└── dimensions // Shape information
This abstraction allows operations and conditions to work uniformly with both variables and array elements.
The interpreter integrates with the rest of the Ramanujan platform through:
The NativeProcessor.cpp file provides the JNI bridge, receiving JSON input from the Java layer and returning results as structured data.
When running execute_inline with CSV data files, the runtime now bypasses the Python interpreter for data population. Previously, every CSV value was converted to a Python assignment statement (arr[r][c] = val) and run through the full AST interpreter — for a 50-million-element CSV, that meant 50 million interpreted statements.
The new approach:
arr = [[0 for _ in range(3072)] for _ in range(9216)]).Double.parseDouble() + ConcurrentHashMap.put(), skipping the interpreter entirely.This makes it practical to load gigabytes of weight data (e.g., 41 GB of neural network weights across 197 CSV files) in seconds rather than hours.
Single-row CSVs (one line of comma-separated values) are automatically declared as 1D arrays ([0 for _ in range(N)]), while multi-row CSVs are declared as 2D arrays ([[0 for _ in range(cols)] for _ in range(rows)]). The population step stores values accordingly:
"0", "1", "2", ..."row_col" format ("0_0", "0_1", ..., "1_0", ...)CSV filenames are automatically converted to valid array names:
../weights/l0_iln.csv → l0_iln).csv extension is removed_This means the array name in your kernel code must match the CSV filename (without path and extension).
dump Command (Array Extraction)The execute_inline query console now supports a dump command for extracting array contents to CSV files:
dump <arrayName> [outputFile]
The command auto-detects whether the array is 1D or 2D:
# Run a kernel, then dump results
printf 'dump normed /tmp/normed_out.csv\nexit\n' | \
java -Xmx4g -jar developer-console-fat.jar execute_inline kernel.py data.csv weights.csv
# Check output
cat /tmp/normed_out.csv
# 0.00288,-0.00078,0.00454,...
All query console commands:
var <name> — print a scalar variablearr <name> <index> — print a single array elementdump <name> [file] — dump entire array to CSVexit — end the sessionThe ramanujan-test-codes/phi3/ directory contains a complete implementation of Microsoft Phi-3-mini-4k-instruct (3.8B parameters) running entirely on Ramanujan's GPU runtime. This serves as a reference for running large neural networks on the platform.
| File | Purpose |
|---|---|
extract_weights.py | Converts safetensors model → 197 CSV weight files (41 GB) |
inference.py | Hybrid NumPy + Ramanujan inference (CPU fallback) |
inference_rj.py | Full Ramanujan inference — every layer runs on GPU |
layer_kernel.py | Single transformer layer GPU kernel (RMS norm, QKV projection, RoPE, attention, FFN) |
embed_kernel.py | Embedding lookup kernel |
head_kernel.py | Final norm + LM head projection kernel |
Rather than loading all 41 GB of weights at once, the orchestrator (inference_rj.py) calls rj once per transformer layer, passing only that layer's weights (~1.3 GB) as CSV files. Between layers, results are extracted via the dump command and passed to the next invocation.
Token → [embed_kernel] → hidden
↓
[layer_kernel × 32] ← weights loaded per-layer
↓
[head_kernel] → logits → argmax → next token
See ramanujan-test-codes/phi3/README.md for full details on weight extraction, model architecture, and usage instructions.
64 commits
Java
66.7%
Python
19.0%
C++
10.4%
JavaScript
2.6%
The Ramanujan Platform aims to harness the idle computational power of various digital devices to create a powerful distributed network. This platform allows any device with basic arithmetic capabilities and internet connectivity to contribute to the network's computational power.
20
stars
64
commits
Java
primary language
Sep 5, 2026
updated
This project aims to utilize the untapped computation power of digital devices. Any device that has a CPU which can do basic arithmetic operations and can connect to the internet contribute to the computation power of the network.
Apollo guidance computer had a CPU which was only as powerful as a modern scientific calculator. Modern smartphones and smart TVs are at-least million times more powerful than the Apollo guidance computer, and most of the time they are idle. This project aims to utilize the idle computation power of these devices.
Any device which has a CPU that can do basic arithmetic operations and can connect to the internet can be added on the network. This includes smartphones, smart TVs, smart-watches, smart speakers, smart refrigerators, smart washing machines, and so on. Currently, the network supports only Android devices, linux devices, Windows devices, and macOS devices.
BOINC is an amazing platform that allows users to donate their computation power to various scientific projects. However, for adding a new kind of computation on the network, the project owner has to write a new client for the BOINC network, and the devices on the network that can contribute to the computation have to download and install the new client.
In this platform, for any new kind of computation, the project owner doesn't have to write a new client. The owner just
have to give the computation code in the ramanujan language. The platform would convert the code to an intermediate
code that the device can just run. The devices on the network don't have to download and install a new client for every
new kind of computation. They just have to download and install the ramanujan client one-time.
Variables in ramanujan are declared using the var keyword. Like in most programming languages, the type of the
variable has to be explicitly given. Currently, the supported types are integer, double.
Example:
var x:integer;
var y:double;
Arrays in ramanujan are declared using the var keyword. The size of the array be both variable and explicitly given. This is a special kind of variable and the data type has to be given as array.
Example:
var arr[100]:array;
var n : double;
n =10;
var arr1[n]: array
Functions in ramanujan are declared using the def keyword. The functions do not have any return type. The parameters
of the functions are both input and output. The parameters have to be declared with the var keyword. The parameters
are passed by reference.
To invoke a function, the exec keyword is used.
Example:
def add(var x:integer, var y:integer, var ans:integer) {
ans = x + y;
}
exec add(1, 2, ans);
Here, the value of x is 1, the value of y is 2, and the value of ans NULL when passed to the function. But, since
the arguments are passed by reference, the value of ans is 3 after the function is executed.
ramanujan supports while loops. The syntax is similar to most programming languages.
Example:
var i:integer;
i = 0;
while(i < 10) {
i = i + 1;
}
ramanujan supports if-else statements. The syntax is similar to most programming languages.
Example:
var x:integer;
x = 10;
if(x < 10) {
x = 0;
} else {
x = 1;
}
There are two methods in Ramanujan which need to be used to distribute the computation over multiple devices:
threadStart:
The threadStart keyword is used to define a new thread. Any new variable can be introduced in the thread, or it
can use the variables defined outside the thread. in global scope The thread can call any function defined in the code.
Syntax:
threadStart(t0) {
// code
}
Here, t0 is the name of the thread. The thread name has to be unique in the code.
Example:
threadStart(t0) {
exec add(1, 2, ans);
}
threadOnEnd:
The threadOnEnd keyword is used to define actions that should be taken when one or more threads complete.
The body code executes only on the last iteration.
Syntax:
threadOnEnd(thread_seperated_thread_names, number_of_iterations) {
// code
}
Example:
threadOnEnd(t0, t1, 5) {
// code
}
Here, t0, t1 are the names of the threads. The last argument
5 is the number of iterations the threads have to run. The code defined in threadOnEnd for the n-1 times (here 4 times)
would spawn the threads again. On the n-th time, the code defined in threadOnEnd would be executed.
T1 T1 T1 T1 T1
/ \ / \ / \ / \ / \
X Y Y Y Y Z
\ / \ / \ / \ / \ /
T0 T0 T0 T0 T0
Here, In Y nodes, its just joining the thread and would do nothing, but on the last iteration, it would execute the code
defined in threadOnEnd as node Z.
threadParallelismCycle:
The threadParallelismCycle keyword is used to define actions that should be taken after every cycle of parallelism.
Unlike threadOnEnd, which only executes its body on the final iteration, threadParallelismCycle executes its body
after each and every completed cycle (including the last one).
Syntax:
threadParallelismCycle(thread_names, number_of_iterations) {
// code – runs after every cycle
}
Example:
threadParallelismCycle(t0, t1, 5) {
// this block runs after each of the 5 cycles
}
The DAG structure for threadParallelismCycle(t0, t1, 3) { body }:
T0 T0 T0
/ \ / \ / \
X body body body (final)
\ / \ / \ /
T1 T1 T1
Each body node runs after every pair of T0+T1 threads completes. After each body (except the last), the threads
are re-spawned for the next cycle.
Important: Do not combine
threadParallelismCycleandthreadOnEndon the same thread set.threadParallelismCyclealready runs its body after the final cycle too, sothreadOnEndis redundant. When added on the same threads,threadOnEndinjects an extra empty join node as a second successor of every first-cycle thread (instead of the expected single cycle-body successor), and overwrites the internal re-spawn map entries thatthreadParallelismCycleset up. This breaks the cycle chain and creates an unintended parallel execution path.If you want a dedicated block that only runs after the last cycle, use
threadOnEndalone (withoutthreadParallelismCycle). If you want code that runs after every cycle (including the last), usethreadParallelismCyclealone.
Example of complex threading:
T2 T3
/ \ / \
N0 N1 N2
\ / \ /
T1 T4
Following code would help here:
threadStart(t1) {
}
threadStart(t2) {
}
threadOnEnd(t1, t2, 1) {
threadStart(t3) {
}
threadStart(t4) {
}
}
threadOnEnd(t3, t4, 1) {
}
Example of a distributed gradient descent with Particle Swarm Optimization in ramanujan language:
def getSquared(var xPow:integer, var yPow:integer, var ans:integer) {
if(xPow < yPow) {
ans = yPow - xPow;
} else {
ans = xPow - yPow;
}
}
def getAvg(var arr:array, var originalArr:array, var avgF:integer) {
var index,ans1,tmpAvg1,tmpAvg2:integer;
avgF = 0;
index = 0;
while(index < 100) {
tmpAvg1 = arr[index];
tmpAvg2 = originalArr[index];
exec getSquared(tmpAvg1,tmpAvg2, ans1);
avgF = avgF + ans1;
index = index + 1;
}
avgF = avgF / 100;
}
def getTestArr(var xTest:integer, var yTest:integer, var testArrTest:array) {
var it:integer;
it = 0;
while(it < 100) {
testArrTest[it] = xTest * it + yTest;
it = it + 1;
}
}
var train[100]:array;
var i:integer;
i = 0;
while(i < 100) {
train[i] = i * 1.9 + 33;
i = i + 1;
}
def mainCode(var train : array, var x1:double, var y1:double) {
var x1,y1,j,avg,diff1,diff2x,diff2y,tmp:double;
j = 0;
var testArr[100]:array;
var slope:double;
var nexty,nextx:double;
testArr[1] = 1;
while(j < 15000) {
exec getTestArr(x1,y1,testArr);
exec getAvg(testArr, train, diff1);
tmp = x1 + 0.0001;
exec getTestArr(tmp,y1,testArr);
exec getAvg(testArr, train, diff2x);
slope = (diff2x - diff1) / 0.0001;
nextx = x1 - slope * 0.1;
tmp = y1 + 0.0001;
exec getTestArr(x1,tmp,testArr);
exec getAvg(testArr, train, diff2y);
slope = (diff2y - diff1) / 0.0001;
nexty = y1 - slope * 0.50;
x1 = nextx;
y1 = nexty;
j = j + 1;
}
}
var x1[100][10],y1[100][10]:array;
x1[0][0] = 0;
y1[0][0] = 0;
var ansX1,ansy1 :double;
ansX1 = 0;
ansy1 = 0;
var iteration[10]:array;
i = 0;
while(i < 10) {
iteration[i] = 0;
i = i + 1;
}
def getBest(var train:array, var best:integer, var x1:array, var y1:array, var iteration:integer) {
best = 0;
var index:integer;
var bestM:double;
bestM = 1000000000;
index = 0;
while(index < 10) {
var testArr[100]:array;
testArr[0] = 0;
var testX1,testY1:double;
testX1 = x1[index][iteration];
testY1 = y1[index][iteration];
exec getTestArr(testX1,testY1,testArr);
var avg:double;
avg = 0;
exec getAvg(testArr, train, avg);
if(avg < bestM) {
bestM = avg;
best = index;
}
index = index + 1;
}
}
def posRun(var thread:integer, var train:array, var x1:array, var y1:array, var iteration :array) {
var currentIter:integer;
currentIter = iteration[thread];
if(currentIter == 0) {
x1[thread][currentIter]=thread;
y1[thread][currentIter]=thread;
} else {
var best :integer;
best=0;
var thisIter:integer;
thisIter=currentIter;
currentIter = currentIter-1;
exec getBest(train, best, x1, y1, currentIter);
if(x1[thread][currentIter] < x1[best][currentIter]) {
x1[thread][thisIter] = x1[thread][currentIter]+(x1[best][currentIter]-x1[thread][currentIter])/2;
} else {
x1[thread][thisIter] = x1[thread][currentIter]-(x1[thread][currentIter]-x1[best][currentIter])/2;
}
if(y1[thread][currentIter] < y1[best][currentIter]) {
y1[thread][thisIter] = y1[thread][currentIter]+(y1[best][currentIter]-y1[thread][currentIter])/2;
} else {
y1[thread][thisIter] = y1[thread][currentIter]-(y1[thread][currentIter]-y1[best][currentIter])/2;
}
currentIter=thisIter;
}
var x,y:double;
x=x1[thread][currentIter];
y=y1[thread][currentIter];
exec mainCode(train, x, y);
x1[thread][currentIter]=x;
y1[thread][currentIter]=y;
}
threadStart(t0) {
exec posRun(0, train, x1, y1,iteration);
iteration[0]=iteration[0]+1;
}
threadStart(t1) {
exec posRun(1, train, x1, y1,iteration);
iteration[1]=iteration[1]+1;
}
threadStart(t2) {
exec posRun(2, train, x1, y1,iteration);
iteration[2]=iteration[2]+1;
}
threadStart(t3) {
exec posRun(3, train, x1, y1,iteration);
iteration[3]=iteration[3]+1;
}
threadStart(t4) {
exec posRun(4, train, x1, y1,iteration);
iteration[4]=iteration[4]+1;
}
threadStart(t5) {
exec posRun(5, train, x1, y1,iteration);
iteration[5]=iteration[5]+1;
}
threadStart(t6) {
exec posRun(6, train, x1, y1,iteration);
iteration[6]=iteration[6]+1;
}
threadStart(t7) {
exec posRun(7, train, x1, y1,iteration);
iteration[7]=iteration[7]+1;
}
threadStart(t8) {
exec posRun(8, train, x1, y1,iteration);
iteration[8]=iteration[8]+1;
}
threadStart(t9) {
exec posRun(9, train, x1, y1,iteration);
iteration[9]=iteration[9]+1;
}
threadOnEnd(t0,t1,t2,t3,t4,t5,t6,t7,t8,t9,5) {
var best:integer;
best=0;
exec getBest(train, best, x1, y1, 4);
ansX1=x1[best][4];
ansy1=y1[best][4];
}
The ramanujan language is faster than the Python3 language. The above single node code runs in ~350 ms. The same heuristic
in Python3 takes ~410 ms. The device it was tested on was a MacBook Air M3 : 8GB RAM, Apple M3 chip.
Ramanujan now supports a subset of Python syntax. Python code is converted to Ramanujan's intermediate representation using Python's AST module.
The conversion from Python code to Ramanujan's intermediate code follows this flow:
Python Source Code (.py file)
↓
[PythonAstInvoker]
Writes code to temp file, invokes: python3 -c "import ast, ast2json, json; ..."
↓
Python's ast.parse() generates AST
↓
ast2json.ast2json() converts AST to JSON
↓
JSON string returned to Java
↓
[JsonAstParser]
Parses JSON into Java AST node objects (ModuleNode, AssignNode, IfNode, etc.)
↓
[PythonAstToRuleEngineInputConverter]
Traverses AST nodes and generates RuleEngineInput structures
(Variables, Operations, Commands, Conditions, etc.)
↓
RuleEngineInput (Ramanujan Intermediate Code)
↓
Execution by Ramanujan Engine
PythonAstInvoker: Creates a temporary .py file, then runs a Python snippet that:
ast, ast2json, and jsonast.parse()ast2json.ast2json()JsonAstParser: Uses Jackson to parse the JSON and creates corresponding Java AST node objects.
PythonAstToRuleEngineInputConverter: Walks the AST tree and generates:
Variable objects for variable declarationsArray objects for list/array declarationsOperation objects for arithmetic and assignmentsCondition objects for comparisonsCommand objects that link everything in execution orderIf/While blocks for control flowFunctionCall objects for function invocationsVariables are automatically declared on first assignment. Type is inferred from the assigned value:
x = 5y = 3.14 or any arithmetic operation resultx = 5 # Integer
y = 3.14 # Double
z = x + y # Double (arithmetic result)
Arrays must be created using list comprehensions initialized with 0. This is the only supported form:
# 1D array initialized with zeros
arr = [0 for _ in range(100)]
# 2D array (matrix)
matrix = [[0 for _ in range(10)] for _ in range(10)]
# n-dimensional arrays with nested comprehensions
tensor = [[[0 for _ in range(5)] for _ in range(10)] for _ in range(3)]
# Variable dimensions are supported
n = 100
arr = [0 for _ in range(n)]
# Array element access and assignment
arr[0] = 10
value = arr[i]
matrix[x][y] = 5
Note: Arrays cannot be initialized with explicit values like [1, 2, 3] - only [0 for _ in range(n)] form is allowed.
Supported operators: +, -, *, /
result = a + b * c - d / e
x += 5 # x = x + 5
y -= 3 # y = y - 3
z *= 2 # z = z * 2
w /= 4 # w = w / 4
Supported: <, <=, >, >=, ==, !=
if x > 5:
pass
while count <= 100:
pass
if x > 10:
y = 1
else:
y = 0
Note: elif is not supported. Use nested if-else instead:
# Instead of elif, use nested if-else:
if x > 10:
y = 1
else:
if x > 5:
y = 2
else:
y = 0
i = 0
while i < 100:
# loop body
i += 1
Functions are defined using def and called directly. Arguments are passed by reference.
def calculate(a, b, result):
result = a + b * 2
# Function call
calculate(x, y, answer)
Functions can return values using tuple unpacking:
def get_coords():
x = 10
y = 20
return x, y
# Tuple unpacking to receive return values
a, b = get_coords()
Single return values:
def compute(x):
result = x * 2
return result
value = compute(5)
Returning an array element directly is NOT supported:
# NOT SUPPORTED
def get_element(arr, index):
return arr[index] # ❌ Cannot return array element directly
Workaround: Assign to a variable first, then return:
# SUPPORTED
def get_element(arr, index):
value = arr[index]
return value # ✓ Return variable works
Passing a function call as an argument to another function is NOT supported:
# NOT SUPPORTED
result = outer_func(inner_func(x)) # ❌ Nested function calls not allowed
Workaround: Use intermediate variables:
# SUPPORTED
temp = inner_func(x)
result = outer_func(temp) # ✓ Works with intermediate variable
Returning the result of a function call directly is NOT supported:
# NOT SUPPORTED
def wrapper(x):
return compute(x) # ❌ Cannot return function call directly
Workaround: Assign to a variable first:
# SUPPORTED
def wrapper(x):
result = compute(x)
return result # ✓ Return variable works
Only simple function names are supported (no method chains or computed function references):
# NOT SUPPORTED
obj.method() # ❌ Method calls on objects
funcs[0]() # ❌ Function from array
getattr(obj, 'func')() # ❌ Dynamic function access
for loops are NOT currently supported. Use while loops instead:
# NOT SUPPORTED
for i in range(10): # ❌
pass
# SUPPORTED - Use while loop
i = 0
while i < 10: # ✓
i += 1
Object-oriented programming is not yet supported:
# NOT SUPPORTED
class MyClass: # ❌
pass
Boolean expressions with and/or/not are NOT supported:
# NOT SUPPORTED
if x > 5 and y < 10: # ❌
pass
if not flag: # ❌
pass
Workaround: Use nested if statements:
# SUPPORTED
if x > 5:
if y < 10: # ✓ Nested if for 'and' logic
pass
Dynamic list operations are not supported:
# NOT SUPPORTED
arr.append(5) # ❌ append
arr.pop() # ❌ pop
len(arr) # ❌ len function
arr + other # ❌ list concatenation
Strings are NOT supported:
# NOT SUPPORTED
s = "hello" # ❌ string assignment
s = "hello" + "world" # ❌ string concatenation
Importing modules is not supported:
# NOT SUPPORTED
import math # ❌
from collections import deque # ❌
Try-except blocks are not supported:
# NOT SUPPORTED
try: # ❌
pass
except:
pass
Generator expressions and iterator protocols are not supported.
Power (**) and modulo (%) operators are NOT currently supported:
# NOT SUPPORTED
x = 2 ** 10 # ❌ power operator
y = n % 10 # ❌ modulo operator
The elif keyword is NOT supported. Use nested if-else:
# NOT SUPPORTED
if x > 10:
y = 1
elif x > 5: # ❌ elif not allowed
y = 2
# SUPPORTED - Use nested if-else
if x > 10:
y = 1
else:
if x > 5: # ✓
y = 2
# Gradient descent example in Python syntax
def get_squared(x_pow, y_pow):
if x_pow < y_pow:
ans = y_pow - x_pow
else:
ans = x_pow - y_pow
return ans
def get_test_arr(x, y, test_arr):
it = 0
while it < 100:
test_arr[it] = x * it + y
it = it + 1
# Initialize training data
train = [0 for _ in range(100)]
i = 0
while i < 100:
train[i] = i * 1.9 + 33
i = i + 1
# Main computation
x1 = 0.0
y1 = 0.0
j = 0
test_arr = [0 for _ in range(100)]
while j < 1000:
get_test_arr(x1, y1, test_arr)
# ... gradient computation logic
j = j + 1
Ramanujan supports offloading compute-intensive work to the GPU via OpenCL. Any Python function
whose name matches the pattern funcName_GPU_N (where N is a positive integer) is automatically
compiled to an OpenCL C kernel during translation and dispatched to the GPU at runtime.
The function name encodes the number of NDRange dimensions:
def funcName_GPU_N(dataArg1, dataArg2, ..., dataArgK, rangeDim1, ..., rangeDimN):
# body – use rangeDimK as the per-work-item index
| Part | Role | What the translator does |
|---|---|---|
N in _GPU_N | Number of NDRange dimensions (work_dim) | Read from the function name; no extra argument needed |
dataArg1 … dataArgK | Data arrays (first total_params − N parameters) | Emitted as __global float* kernel parameters |
rangeDim1 … rangeDimN | Work-item index variables (last N parameters) | Emitted as int dim = get_global_id(K); declarations; their call-site values become global_work_size[] for clEnqueueNDRangeKernel |
Requirement: Data args must be arrays. The translator maps them to OpenCL buffers.
def vector_add_GPU_1(a, b, c, gid):
c[gid] = a[gid] + b[gid]
# 1024-element arrays, 1 NDRange dimension
vector_add_GPU_1(a, b, c, 1024)
Generated OpenCL kernel:
__kernel void vector_add(__global float* a, __global float* b, __global float* c) {
int gid = get_global_id(0);
c[gid] = (a[gid] + b[gid]);
}
def matrix_add_GPU_2(a, b, c, row, col):
c[row] = a[row] + b[col]
# 64 × 64 grid (2 NDRange dimensions)
matrix_add_GPU_2(a, b, c, 64, 64)
if/else and while inside the function body are translated to their C equivalents:
def relu_GPU_1(a, out, gid):
if a[gid] > 0:
out[gid] = a[gid]
else:
out[gid] = 0
relu_GPU_1(a, out, 512)
A GPU kernel can call ordinary (non-_GPU_N) Python functions that are defined in the same file.
The translator converts those helpers to OpenCL C device functions and prepends them before the
__kernel declaration so they are visible to the kernel body.
# Plain Python helper – becomes a float device function in the generated OpenCL C
def scale2(x):
r = x * 2
return r
# GPU kernel – calls the helper for every work-item
def apply_scale_GPU_1(a, out, gid):
v = a[gid] # load array element into a local variable first
out[gid] = scale2(v)
apply_scale_GPU_1(a, out, 1024)
Generated OpenCL C:
float scale2(float x) {
float r = (x * 2);
return r;
}
__kernel void apply_scale(__global float* a, __global float* out) {
int gid = get_global_id(0);
float v = a[gid];
out[gid] = scale2(v);
}
| Constraint | Detail |
|---|---|
| Parameter types | All helper parameters are treated as float scalars. Array pointers (__global float*) are not supported in helper functions. |
| Return type | Always float. |
| No subscript as argument | Array element expressions (a[i]) cannot be passed directly as arguments to a helper call. Assign the element to a local variable first: v = a[gid]; out[gid] = scale2(v). |
| No recursion | A function may not call itself. The translator throws a CompilationException if a recursive call is detected at translation time. |
| No GPU–kernel calls | A helper (or the kernel) may not call another _GPU_N function. GPU kernels are dispatched via clEnqueueNDRangeKernel and cannot be invoked as device functions. |
| Scope | Only top-level module functions are eligible as helpers. Nested function definitions are not supported. |
The translator enforces the no-recursion rule at translation time, not at runtime:
# INVALID – will raise CompilationException during translation
def bad_GPU_1(a, gid):
a[gid] = bad_GPU_1(a, gid) # ❌ recursive GPU call
# INVALID – helper self-recursion is also rejected
def factorial(n):
return n * factorial(n - 1) # ❌ recursive helper
def kernel_GPU_1(a, gid):
a[gid] = factorial(a[gid])
These special function names are recognised inside _GPU_N kernel bodies and
translated to OpenCL C intrinsics. They are not available on the host; use
the corresponding Ramanujan host built-ins (FLOOR, EXP, etc.) outside GPU
functions.
Each call mutates its argument in place: FUNC(x) → x = func(x) in the
generated OpenCL C.
| Python call | Generated OpenCL C | Notes |
|---|---|---|
EXP(x) | x = exp(x); | Natural exponential |
LOG(x) | x = log(x); | Natural logarithm |
SQRT(x) | x = sqrt(x); | Square root |
FLOOR(x) | x = floor(x); | Round toward −∞; result stays float |
FLOOR example — locate the grid cell for a particle position:
def p2g_GPU_1(positions, params, gid):
xp = positions[gid * 3]
gxp = (xp - (-1.0)) * 10.0 # map to grid coords
FLOOR(gxp) # gxp = floor(gxp) in OpenCL C
i0 = gxp # integer grid index (stored as float)
nibble = PACKED_NIBBLE(packed, index)
PACKED_NIBBLE numerically converts an exact integer-valued packed float to
OpenCL uint, shifts by index * 4, masks with 0xF, and returns the nibble as
a float. Valid nibble indexes are 0 through 5 for Ramanujan's six-weights-per-
float Phi-3 format. Packed values must not exceed 0xFFFFFF, which is exactly
representable in float32.
((float)(((uint)packed >> ((uint)index * 4u)) & 15u))
This avoids floating-point divide and floor() during 4-bit dequantization.
It is an expression intrinsic and may be used on the right-hand side of an
assignment inside _GPU_N functions.
ATOMIC_ADD_F(arr, idx, delta)
Atomically adds delta (a float) to arr[idx] using a compare-and-swap
loop. Required whenever multiple work-items may scatter into the same array
slot concurrently (e.g., particle-to-grid scattering in MPM).
OpenCL 1.2 has no native atomic_add for float. The translator emits a
CAS loop on the int-reinterpreted bits using atomic_cmpxchg:
/* Generated for: ATOMIC_ADD_F(arr, idx, delta) */
{
__global volatile int* _aAddr = (__global volatile int*)(&arr[(int)(idx)]);
int _aOld, _aNew;
do {
_aOld = *_aAddr;
_aNew = as_int(as_float(_aOld) + (delta));
} while (atomic_cmpxchg(_aAddr, _aOld, _aNew) != _aOld);
}
The {} block scope lets you call ATOMIC_ADD_F multiple times in the same
kernel without variable-name conflicts.
Requirements:
__global float* data argument (not a local variable).atomic_cmpxchg on __global int* is a
core 1.2 feature on all platforms including Apple Metal-backed OpenCL).ATOMIC_ADD_F example — particle-to-grid mass scatter:
def p2g_GPU_1(positions, g_mass, g_vel, params, gid):
# ... compute weight w and grid node index gnode ...
wm = w * params[0] # weighted mass
ATOMIC_ADD_F(g_mass, gnode, wm) # safe concurrent scatter
ATOMIC_ADD_F(g_vel, gnode * 3, wm * vx)
ATOMIC_ADD_F(g_vel, gnode * 3 + 1, wm * vy)
ATOMIC_ADD_F(g_vel, gnode * 3 + 2, wm * vz)
Note on P2G in MPM: The above P2G kernel also requires
FLOORto locate grid nodes (see above). BothFLOORandATOMIC_ADD_Fare needed before P2G can fully run on GPU.
| Platform | Requirement |
|---|---|
| macOS | OpenCL is part of the system framework – no extra install needed |
| Linux | sudo apt install ocl-icd-opencl-dev opencl-headers |
| Windows | Install GPU vendor drivers: NVIDIA CUDA Toolkit, AMD ROCm, or Intel OpenCL SDK |
cl_kernel.double → float before upload and float → double after read-back (OpenCL kernels operate on float).stderr and execution returns immediately.GPU_ENABLED macro must be set at compile time (via -DENABLE_GPU=ON). Builds without it contain no OpenCL code and have no OpenCL runtime dependency.GPU_SYNCBy default, GPU kernels dispatched with _GPU_N functions are non-blocking: the OpenCL
command is queued but the CPU continues immediately. This allows many kernel launches to
be batched together without the CPU stalling after every one — which is the key to high GPU
throughput.
However, any time the host (CPU) side needs to read back a value that a GPU kernel has
written, you must explicitly drain the GPU queue for that array first. The built-in
GPU_SYNC does exactly that.
GPU_SYNC(array)
GPU_SYNC(array) issues a blocking clEnqueueReadBuffer for the given array, flushing
all previously enqueued GPU work and copying the updated data back to the host buffer. It
is a no-op for arrays that are not GPU-backed (e.g., host-only arrays).
GPU_SYNC| Situation | Action |
|---|---|
Reading an array element in a Python host while loop after a GPU kernel has written it | Call GPU_SYNC(array) once before the loop |
Passing a GPU-written array to a host function or exec call | Call GPU_SYNC(array) before the call |
dump array /path after GPU kernel(s) wrote it | Call GPU_SYNC(array) before dump |
| Using a GPU-written array only as input to the next GPU kernel (no host read) | No GPU_SYNC needed — GPU→GPU is handled automatically |
# 120+ GPU kernels dispatched with no CPU stalls …
layernorm_GPU_1(hidden, ln_g, ln_b, h_ln, n_seq)
matmul_bias_GPU_2(h_ln, c_attn_w, c_attn_b, qkv, kp, n_seq, 2304)
# … more kernels …
matmul_bias_GPU_2(h_ff, c_fc_proj_w, c_fc_proj_b, h_out_buf, kp, n_seq, 768)
# CPU needs to read `hidden` and `h_out_buf` in the next loop → sync first
GPU_SYNC(hidden)
GPU_SYNC(h_out_buf)
_i = 0
while _i < n_seq * 768:
hidden[_i] = hidden[_i] + h_out_buf[_i]
_i = _i + 1
Without the GPU_SYNC calls, hidden and h_out_buf would still hold stale values from
before the last GPU kernels ran, producing silently wrong results.
RETURNBy default, after execution completes arrChangeMap() reports every array that was modified
(comparing each element against its pre-execution snapshot). For large models with hundreds of
weight arrays this means the JNI result map can contain far more data than the caller needs.
The RETURN built-in lets user code explicitly name the arrays that should appear in the result.
If RETURN is called at any point during execution, only the listed arrays are included in
arrChangeMap(); all other modified arrays are silently dropped. If RETURN is never called
the behaviour is unchanged — all modified arrays are returned.
RETURN(arr1, arr2, ..., arrN)
RETURN accepts any number of array arguments and marks each with an internal flag at runtime.
It is a no-op for the computation itself (it does not stop execution or modify any values).
hidden = [0 for _ in range(768)]
kv_cache = [0 for _ in range(4096)]
weights = [0 for _ in range(131072)] # large weight buffer
# ... kernel calls that write to all three arrays ...
# Only return the outputs the caller actually needs
RETURN(hidden, kv_cache)
Without RETURN, weights and every other modified array would be serialised back through JNI
even though the caller only needs hidden and kv_cache.
Script contains RETURN(...)? | What arrChangeMap() returns |
|---|---|
| No | All modified arrays with their changed indexes (unchanged default) |
| Yes | Only the listed arrays, still reporting only changed indexes |
RETURN vs GPU_SYNCGPU_SYNC and RETURN are independent and complementary:
| Purpose | When to use | |
|---|---|---|
GPU_SYNC(arr) | Flush the GPU queue and read arr back to the CPU | Before the host reads a GPU-written array |
RETURN(arr1, ...) | Filter which arrays are included in the final result | To reduce serialisation overhead when only a subset of arrays is needed |
RELEASE_MEM and LOAD_MEMEvery array a _GPU_N kernel touches gets its own GPU buffer, and buffers are never freed
automatically while the process is running. For large models (e.g. dozens of transformer
layers, each with several 4-bit packed weight/scale arrays plus K/V caches) this means all
per-layer buffers stay resident on the GPU simultaneously. On memory-constrained devices
(e.g. ~5 GB unified memory) this can exhaust GPU/unified memory and cause an OOM crash.
RELEASE_MEM and LOAD_MEM give user code explicit control over when a buffer actually
occupies GPU memory:
RELEASE_MEM(array) # Frees the array's GPU buffer immediately (host data is untouched)
LOAD_MEM(array) # (Re)allocates the GPU buffer and uploads the current host data
RELEASE_MEM(array) calls clReleaseMemObject on the array's buffer and clears it, so the
memory is returned to the driver right away. The host-side (CPU) copy of the array is never
touched, so the data itself is not lost.LOAD_MEM(array) allocates a fresh GPU buffer for the array from its current host data.
It is the correct counterpart to RELEASE_MEM — plain GPU_LOAD only writes into an
already-existing buffer, so it cannot bring back an array that has been released.Mandatory, not optional. A
_GPU_Nkernel dispatch performs no implicit buffer allocation or upload of its own — it only doesclSetKernelArg+clEnqueueNDRangeKernelagainst whatever GPU buffer the array already has. This means:
LOAD_MEM(array)is required before the first time an array is passed as a data argument to any_GPU_Nkernel, and again after anyRELEASE_MEM(array)on it. Skipping this is not silently "slow" — it is a hard error: the kernel logsmissing LOAD_MEM(array) before first useand the dispatch is skipped entirely (the kernel's outputs are left unchanged).GPU_LOAD(array)is required any time host (CPU) code mutates an array that already has a live GPU buffer (e.g.l_idx_arr[0] = l_idx) and the next kernel dispatch needs to see that new value.GPU_LOADonly writes into an existing buffer — it never allocates one — so it must not be used as a substitute forLOAD_MEMon an array that hasn't been loaded yet (or was just released).
RELEASE_MEM is intended for immutable arrays (4-bit packed weights, scale tables, RoPE
cos/sin tables, etc.). Never release an array that a GPU kernel writes to (activation
buffers, KV caches, etc.) without first reading its current value back with GPU_SYNC —
RELEASE_MEM does not sync, so releasing a GPU-resident write target discards whatever the
GPU hasn't yet written back to the host copy.
Release/reload has a cost — every call is buffer churn.
RELEASE_MEM+LOAD_MEMon the same array is aclReleaseMemObject+clCreateBufferpair, not a full data copy, for weight arrays uploaded withCL_MEM_USE_HOST_PTR(seeisBinaryLoadedin the "Direct CSV Population" section above) — on unified-memory devices (e.g. Apple Silicon) this is cheap since no bytes actually move. It is still real per-call overhead, though, so only cycle an array throughRELEASE_MEM/LOAD_MEMwhen you specifically need to cap peak GPU memory (e.g. to fit a multi-layer model on a memory-constrained device); don't do it reflexively for arrays that are about to be reused within the same layer or the same kernel dispatch.
If you don't need to cap peak GPU memory mid-run, the simplest use is to release every large buffer exactly once, right before the process is about to end (or before the very last use of a single-pass computation), instead of releasing/reloading inside a hot loop:
# ... 32 transformer layers, prefill + decode loop, all reusing the same
# per-layer weight/cache buffers on every iteration ...
# Generation is finished — nothing on the GPU is needed anymore before the
# host reads back `generated_tokens`. Release every large buffer once,
# instead of thrashing release/reload inside the decode loop.
RELEASE_MEM(l0_qkv_packed)
RELEASE_MEM(l0_qkv_scales)
RELEASE_MEM(l0_k_cache)
RELEASE_MEM(l0_v_cache)
# ... one RELEASE_MEM call per large buffer ...
GPU_SYNC(generated_tokens)
When peak GPU memory itself is the constraint (e.g. a multi-layer transformer that would
otherwise keep every layer's weights resident simultaneously), LOAD_MEM an immutable
layer's weights right before that layer's kernels run, and RELEASE_MEM them again
immediately after — in every place that layer is used, including inside a repeated decode
loop. Because weight arrays are read-only inputs (never a GPU write target) and are typically
isBinaryLoaded (zero-copy CL_MEM_USE_HOST_PTR), the repeated release/reload is just cheap
buffer-object churn, not a data copy, and at any instant only one layer's weights occupy GPU
memory:
# ── Layer N ──
LOAD_MEM(lN_qkv_packed)
LOAD_MEM(lN_qkv_scales)
LOAD_MEM(lN_o_packed)
LOAD_MEM(lN_o_scales)
LOAD_MEM(lN_gate_up_packed)
LOAD_MEM(lN_gate_up_scales)
LOAD_MEM(lN_down_packed)
LOAD_MEM(lN_down_scales)
rmsnorm_GPU_1(h_state, lN_ln1_g, h_ln1, n_seq)
matmul_4bit_GPU_2(h_ln1, lN_qkv_packed, lN_qkv_scales, qkv_buf, kp_qkv, n_seq, 9216)
# ... rest of layer N's kernels ...
residual_add_GPU_2(h_state, h_out_buf, n_seq, 3072)
RELEASE_MEM(lN_qkv_packed)
RELEASE_MEM(lN_qkv_scales)
RELEASE_MEM(lN_o_packed)
RELEASE_MEM(lN_o_scales)
RELEASE_MEM(lN_gate_up_packed)
RELEASE_MEM(lN_gate_up_scales)
RELEASE_MEM(lN_down_packed)
RELEASE_MEM(lN_down_scales)
Note that KV caches (lN_k_cache/lN_v_cache) are deliberately not part of this
per-layer cycle — they are write targets updated by every layer every step, so they must stay
GPU-resident across the whole run (they are only released once, at the very end, alongside the
other shared buffers per Pattern 1). See
ramanujan-test-codes/phi3/phi3_transformer_stack_4bit.py
for the full worked example across all 32 layers, in both the prefill pass and the decode loop.
GPU_SYNC vs the previous implicit sync modelBefore GPU_SYNC was introduced, every _GPU_N call automatically issued a blocking
clEnqueueReadBuffer + clFinish after the kernel, preventing any batching. The overhead
measured ~11% of total inference time on macOS (visible as IOKit → IOGPU → clFinish in
profiler traces). With the explicit model, the GPU queue is drained only at the
necessary points, and all other kernel dispatches remain asynchronous.
Ramanujan now supports a subset of Python syntax through AST-based conversion (see Python Support section above). The platform is actively evolving to support more Python features progressively.
Python support is being actively developed on the Ramanujan platform. More and more features are being added continuously to bring the full power of Python to distributed computing:
Coming Shortly: Object-Oriented Programming (OOP) support
Progressive Additions: We will progressively add all Python features to the Ramanujan platform, including:
and, or, not)**) and modulo (%) operatorsfor loops and iteratorselif statementstry/except)Long-term Vision: Full Python3 ecosystem compatibility
The goal is to make Python code seamlessly executable on the distributed Ramanujan platform while maintaining performance and enabling parallel computation across devices.
Use mvn clean install. Following is the dependency hierarchy:
cd ramanujan-native/native
mkdir build
cd build
cmake ..
cmake --build .
cd ramanujan-native/native
mkdir build-gpu
cd build-gpu
cmake -DENABLE_GPU=ON ..
cmake --build .
Passing -DENABLE_GPU=ON sets the GPU_ENABLED preprocessor macro, links OpenCL, and activates
OpenCL kernel dispatch for _GPU-suffixed functions. Standard builds (-DENABLE_GPU=OFF, the
default) compile no OpenCL code and have no dependency on any OpenCL runtime.
Dockerfile is provided to containerize all the necessary services.
prayog device server on the network:For executing code file:
java -jar <developer-console-path>/target/developer-console-1.0-SNAPSHOT-fat.jar execute <path-to-code-file>
The translation module (middleware-translation) currently requires the following Python dependencies to convert Python code to Ramanujan intermediate code:
pip install ast2jsonFor more information about third-party licenses, see THIRD_PARTY_LICENSES.md.
To set up the Ramanujan developer console and required dependencies, run the provided installer script:
# Make the installer executable
chmod +x install_ramanujan.sh
# Run the installer
./install_ramanujan.sh
RAMANUJAN_WS environment variable.libjsoncpp if needed.rj to your shell profile for easy usage.After installation, restart your terminal or run:
source ~/.zshrc # or ~/.bashrc, ~/.bash_profile, depending on your shell
You can now run the developer console with:
rj <path-to-code-file>

The code is submitted to DevConsole process which would be present in the path: developer-console/target/developer-console-1.0-SNAPSHOT-fat.jar.
The process would submit the code to the Middleware server, which would be responsible for converting the submitted code
to the intermediate code. The Middleware server would then work with the orchestrator server to process the required DAG.
Once the code gets converted to the intermediate code, the Middleware server would return back an asyncId which the
DevConsole would use to get the result of the code execution.

The Middleware server after creating the intermediate code, would start submitting the DAG-nodes to the orchestrator server. Whenever a DAG-node gets processed, the Middleware server would submit children - DAG elements to the orchestrator server. Once, the whole DAG is computed, the Middleware server sets the status of the whole processing as SUCCESS.
Pseudocode of the DAG computation is as follows:
DagElementQueue = new Queue();
DagElementQueue.add(rootDagElement);
isDone = false;
while(!isDone) {
asyncTasks = []
children = []
while(DagElementQueue is not empty) {
DagElement currentElement = DagElementQueue.poll();
if(currentElement is not processed) {
asyncTask = asyncCode(currentElement.process(), callback={
if(currentElement has children) {
for(DagElement child : currentElement.children) {
children.add(child);
}
}
})
asyncTasks.add(asyncTask);
}
}
for(asyncTask : asyncTasks) {
asyncTask.wait();
}
if(children is empty) {
isDone = true;
} else {
for(DagElement child : children) {
DagElementQueue.add(child);
}
}
}

Middleware end of the orchestrator:
The Middleware server would submit a computation task to the Orchestrator server. The Orchestrator server would instantly
return back an asyncId that the Middleware server would use to get the result of the computation. The Orchestrator server
would then check for an available device, and will assign the computation task to the device.
Device end of the orchestrator:
Any device-client that is available on the Ramanujan platform, would keep pinging to the Orchestrator server. Currently,
the pings contains the deviceId, but in future it would send in the device-stats as well. The Orchestrator server would
keep track of the devices that are available on the network. On the ping API request, the Orchestrator server would return
back a suitable computation task mapped to the device. The device would then execute the task and return back the result to the
Orchestrator server.
This service helps the system to resume in case the Middleware server goes down. This server becomes a part of a PubSub system. It can be part of Apache-Kafka, GCP-PubSub, and a local PubSub system. The producers and consumer in the server code can be extended to use any of the PubSub systems.
In addition to this, the Middleware servers are stateless and do not have the track of the asyncIds. The Kafka-Manager server keeps the track of the asyncIds. The messages on the PubSub queue contains the asyncId and the status of the computation. The consumer of the Kafka-Manager server would ping the Middleware server with the asyncId to get the result of the computation. This helps the Middleware server to move forward with the DAG computation.
The Ramanujan native interpreter (ramanujan-native) is the core execution engine responsible for running the intermediate code on devices. Written in C++ for maximum performance, it provides a sophisticated bytecode-like execution environment that can run on any device with minimal overhead.

The interpreter is designed with a two-phase architecture: first, it converts JSON-based intermediate representation into an optimized in-memory graph of execution objects, then it executes the command chain. This design separates parsing overhead from execution, enabling efficient repeated execution.
When the Processor.process() method is invoked, it receives:
The RuleEngineInput contains vectors of all program constructs:
| Component | Description |
|---|---|
variables | Scalar variables (integer, double) |
arrays | Multi-dimensional arrays |
constants | Immutable values |
commands | Execution units (each command wraps an operation, condition, function call, etc.) |
operations | Arithmetic and assignment operations (+, -, *, /, =) |
conditions | Comparison and logical operations (<, >, ==, !=, &&, ||) |
ifBlocks | Conditional branching structures |
whileBlocks | Loop structures |
functionCalls | Function definitions and call sites |
ID Map Creation:
The createMap() method iterates through all input elements and calls getInternalAnalogy() on each. This factory method converts input objects (parsed from JSON) into Rule Engine (RE) objects optimized for execution:
Variable → VariableRE
Array → ArrayRE
Command → CommandRE
Operation → OperationRE
Condition → ConditionRE
If → IfRE
While → WhileRE
FunctionCall → FunctionCallRE
The resulting unordered_map<string, RuleEngineInputUnits*> allows O(1) lookup of any program element by its unique ID.
After creating the ID map, references between objects must be resolved. During JSON parsing, relationships are stored as string IDs (e.g., a command's nextId field contains the ID of the next command). The graph fixing phase converts these string references to actual object pointers.
fixGraph(map):
Iterates through all RE objects and calls setFields(map) on each. This method:
Example for CommandRE.setFields():
nextCommandRE = dynamic_cast<CommandRE*>(getFromMap(map, command->nextId));
operationCommand = dynamic_cast<OperationRE*>(getFromMap(map, command->operation));
ifCommandRE = dynamic_cast<IfRE*>(getFromMap(map, command->ifBlocks));
whileCommandRE = dynamic_cast<WhileRE*>(getFromMap(map, command->whileId));
fixOperator(map, operations):
For each operation, creates a CachedOperationFunctioning object. This is a Strategy Pattern implementation where the appropriate operator implementation is selected once and cached:
| Operator | Implementation Class |
|---|---|
+ | AddImpl |
- | MinusImpl |
* | MultiplyImpl |
/ | DivideImpl |
= | AssignImpl |
The cached functor stores direct pointers to operands, eliminating lookup overhead during execution.
fixConditions(map, conditions):
Similarly creates CachedConditionFunctioning objects for conditions:
| Condition | Implementation Class |
|---|---|
< | LessThanImpl |
> | GreaterThanImpl |
<= | LessThanEqualToImpl |
>= | GreaterThanEqualToImpl |
== | IsEqualImpl |
!= | NotEqualImpl |
&& | AndImpl |
|| | OrImpl |
Before execution begins, the interpreter saves the original values of all variables and array elements:
for(RuleEngineInputUnits* variable : variableREs) {
VariableRE* variableRE = (VariableRE*)(variable);
dataFieldOriginalData.insert(make_pair(variableRE->getValPtrPtr(), *variableRE->getValPtrPtr()));
}
This enables the varChangeMap() and arrChangeMap() methods to efficiently compute which values changed during execution—essential for returning results to the calling system.
The main execution loop is elegantly simple:
CommandRE* command = dynamic_cast<CommandRE*>(mapBetweenIdAndRuleInput->at(firstCommandId));
while(command != nullptr) {
command = command->get();
}
Each CommandRE.get() method:
process() on that unitnextCommandRE)Command Types and Their Execution:
void OperationRE::process() {
operationFunctioning->set(); // Execute cached operation
}
The CachedOperationFunctioning directly accesses operand values and stores the result, with no string lookups or type checking at runtime.
void IfRE::process() {
CommandRE* commandRE;
if(conditionFunctioning->operate()) {
commandRE = ifCommandRE; // Execute if-block
} else {
commandRE = elseCommandRE; // Execute else-block
}
while(commandRE != nullptr) {
commandRE = commandRE->get(); // Nested command loop
}
}
Conditional blocks spawn a nested execution loop for their body commands.
void WhileRE::process() {
while(conditionFunctioning->operate()) {
CommandRE* commandRE = whileCommandRE;
while(commandRE != nullptr) {
commandRE = commandRE->get(); // Execute loop body
}
}
}
The outer while loop evaluates the condition; the inner loop executes the body commands.
Function execution is the most complex operation, involving context management for proper recursion support.
The FunctionCommandRE handles both user-defined functions and built-in functions.
Built-in Functions:
A factory function GetFunctionCommandRE() checks the function ID and returns optimized implementations:
| Function | Description |
|---|---|
NINF | Set to negative infinity |
PINF | Set to positive infinity |
RAND | Random number [0, 1) |
ABS | Absolute value |
SIN, COS, TAN | Trigonometric functions |
ASIN, ACOS, ATAN | Inverse trigonometric |
FLOOR, CEIL | Rounding functions |
EXP | Exponential (e^x) |
SQRT | Square root |
POW | Power function |
Built-in functions have minimal overhead—they directly access argument values and apply the operation.
User-Defined Function Execution:
The FunctionCommandRE.process() method orchestrates function calls through these steps:
Parameter Setup (Call-by-Reference)
Local Variable Preservation
Function Body Execution
command = firstCommand;
while(command != nullptr) {
command = command->get();
}
Context Restoration & Result Propagation
Memory Management:
The DataContainerValueFunctionCommandREMemMaintainer provides efficient memory pooling for function call contexts, avoiding repeated heap allocations during recursive calls.
After execution completes, two methods collect changed values:
varChangeMap():
for(RuleEngineInputUnits* variableRE1 : variableREs) {
VariableRE* variableRE = (VariableRE*)variableRE1;
double newVal = *variableRE->getValPtrPtr();
varChangeMap->insert(make_pair(variableRE->id, newVal));
}
arrChangeMap():
Compares current array values against saved originals and returns only changed elements with their indices.
The interpreter uses a unified data container system:
DataContainerValue (abstract)
├── DoublePtr // Scalar variable value
└── ArrayDataContainerValue
└── ArrayValue // Multi-dimensional array storage
├── val[] // Flat double array
├── totalSize // Total element count
└── dimensions // Shape information
This abstraction allows operations and conditions to work uniformly with both variables and array elements.
The interpreter integrates with the rest of the Ramanujan platform through:
The NativeProcessor.cpp file provides the JNI bridge, receiving JSON input from the Java layer and returning results as structured data.
When running execute_inline with CSV data files, the runtime now bypasses the Python interpreter for data population. Previously, every CSV value was converted to a Python assignment statement (arr[r][c] = val) and run through the full AST interpreter — for a 50-million-element CSV, that meant 50 million interpreted statements.
The new approach:
arr = [[0 for _ in range(3072)] for _ in range(9216)]).Double.parseDouble() + ConcurrentHashMap.put(), skipping the interpreter entirely.This makes it practical to load gigabytes of weight data (e.g., 41 GB of neural network weights across 197 CSV files) in seconds rather than hours.
Single-row CSVs (one line of comma-separated values) are automatically declared as 1D arrays ([0 for _ in range(N)]), while multi-row CSVs are declared as 2D arrays ([[0 for _ in range(cols)] for _ in range(rows)]). The population step stores values accordingly:
"0", "1", "2", ..."row_col" format ("0_0", "0_1", ..., "1_0", ...)CSV filenames are automatically converted to valid array names:
../weights/l0_iln.csv → l0_iln).csv extension is removed_This means the array name in your kernel code must match the CSV filename (without path and extension).
dump Command (Array Extraction)The execute_inline query console now supports a dump command for extracting array contents to CSV files:
dump <arrayName> [outputFile]
The command auto-detects whether the array is 1D or 2D:
# Run a kernel, then dump results
printf 'dump normed /tmp/normed_out.csv\nexit\n' | \
java -Xmx4g -jar developer-console-fat.jar execute_inline kernel.py data.csv weights.csv
# Check output
cat /tmp/normed_out.csv
# 0.00288,-0.00078,0.00454,...
All query console commands:
var <name> — print a scalar variablearr <name> <index> — print a single array elementdump <name> [file] — dump entire array to CSVexit — end the sessionThe ramanujan-test-codes/phi3/ directory contains a complete implementation of Microsoft Phi-3-mini-4k-instruct (3.8B parameters) running entirely on Ramanujan's GPU runtime. This serves as a reference for running large neural networks on the platform.
| File | Purpose |
|---|---|
extract_weights.py | Converts safetensors model → 197 CSV weight files (41 GB) |
inference.py | Hybrid NumPy + Ramanujan inference (CPU fallback) |
inference_rj.py | Full Ramanujan inference — every layer runs on GPU |
layer_kernel.py | Single transformer layer GPU kernel (RMS norm, QKV projection, RoPE, attention, FFN) |
embed_kernel.py | Embedding lookup kernel |
head_kernel.py | Final norm + LM head projection kernel |
Rather than loading all 41 GB of weights at once, the orchestrator (inference_rj.py) calls rj once per transformer layer, passing only that layer's weights (~1.3 GB) as CSV files. Between layers, results are extracted via the dump command and passed to the next invocation.
Token → [embed_kernel] → hidden
↓
[layer_kernel × 32] ← weights loaded per-layer
↓
[head_kernel] → logits → argmax → next token
See ramanujan-test-codes/phi3/README.md for full details on weight extraction, model architecture, and usage instructions.
64 commits
Java
66.7%
Python
19.0%
C++
10.4%
JavaScript
2.6%