Taichi splits kernel in several for loops?

[unwise-fish]

I'm trying to add all elements in the field mass to Enclosed_mass in such a way that all the elements prior to i in mass[i] are added together and stored in Enclosed_mass[i]. I can easily do this in python scope, but the equivalent code in taichi scope, which I'd like to think should work, - doesn't. I've got frankly no clue why it's not working as expected - maybe taichi runs the full for loop for sum[None] += mass[i] before it then assigns the total sum to each element in Enclosed_mass?

import matplotlib.pyplot as plt
import taichi as ti
ti.init(ti.gpu)
N = 100
rad = ti.field(ti.f32, N)
mass= ti.field(ti.f32, N)
Enclosed_mass = ti.field(ti.f32, N)
sum = ti.field(ti.f32, ())

@ti.kernel
def init():
    for i in rad:
        r = i/N
        rad[i] = r
        mass[i] = 1-r

### Works in python scope
def python_enclosed_masses():
    for i in range(N):
        sum[None] += mass[i]
        Enclosed_mass[i] = sum[None]

### Doesnt work in taichi scope
@ti.kernel
def taichi_enclosed_masses():
    for i in rad:
        sum[None] += mass[i]
        Enclosed_mass[i] = sum[None]

init()
x = rad.to_numpy()
python_enclosed_masses()
y_python = Enclosed_mass.to_numpy()
sum[None] = 0.
taichi_enclosed_masses()
y_taichi = Enclosed_mass.to_numpy()
plt.plot(x, y_python, label='Python [real] solution')
plt.plot(x, y_taichi, label='Taichi [incorrect] solution')
plt.legend()
plt.show()

Apologies if this question seems a bit smooth brained, I just got no clue whats going on here...

[bobcao3]

Because Taichi for loop is parallel, when you do += it's doing an atomic add so summation will work over parallel loop. However keep in mind in parallelized loops, the order of execution in the loop is considered random. Thus the value you read from sum[None] will appear in a random order