In fact, when you type x * 2, it’s not evaluated immediately. That’s why we need x to be Expr.
What Expr express is not a value, but an instruction.
When we type x * 2, it will call to Expr.__mul__ (operator override), which generates coresponding codes for different backend, take x86_64 for example:
mov eax, 2 # when Expr(2)
mov ebx, [x] # when Expr(x)
mul ebx, eax # when Expr.__mul__ called
Then these codes will be executed directly on the CPU, and thus higher performance than python.
That’s why we need @ti.kernel, it translates python codes to native machine codes.