Enum CudaExpr 

pub enum CudaExpr {
    LitInt(i64),
    LitFloat(f64),
    LitBool(bool),
    Var(String),
    ThreadIdx(char),
    BlockIdx(char),
    BlockDim(char),
    GridDim(char),
    SyncThreads,
    AtomicAdd(Box<CudaExpr>, Box<CudaExpr>),
    AtomicSub(Box<CudaExpr>, Box<CudaExpr>),
    AtomicExch(Box<CudaExpr>, Box<CudaExpr>),
    AtomicCas(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>),
    AtomicMax(Box<CudaExpr>, Box<CudaExpr>),
    AtomicMin(Box<CudaExpr>, Box<CudaExpr>),
    BinOp(Box<CudaExpr>, CudaBinOp, Box<CudaExpr>),
    UnOp(CudaUnOp, Box<CudaExpr>),
    Index(Box<CudaExpr>, Box<CudaExpr>),
    Member(Box<CudaExpr>, String),
    PtrMember(Box<CudaExpr>, String),
    Cast(CudaType, Box<CudaExpr>),
    Call(String, Vec<CudaExpr>),
    Ternary(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>),
    Ldg(Box<CudaExpr>),
    ShflDownSync(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>),
    ShflXorSync(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>),
    WarpSize,
    BallotSync(Box<CudaExpr>, Box<CudaExpr>),
    Popc(Box<CudaExpr>),
}

CUDA C++ expression AST node.

Variants

LitInt(i64)

Integer literal: 42

LitFloat(f64)

Float literal: 3.14f

LitBool(bool)

Boolean literal: true / false

Var(String)

Named variable or parameter: x

ThreadIdx(char)

threadIdx.x, threadIdx.y, threadIdx.z

BlockIdx(char)

blockIdx.x, blockIdx.y, blockIdx.z

BlockDim(char)

blockDim.x, blockDim.y, blockDim.z

GridDim(char)

gridDim.x, gridDim.y, gridDim.z

SyncThreads

__syncthreads()

AtomicAdd(Box<CudaExpr>, Box<CudaExpr>)

atomicAdd(addr, val) — atomic addition

AtomicSub(Box<CudaExpr>, Box<CudaExpr>)

atomicSub(addr, val)

AtomicExch(Box<CudaExpr>, Box<CudaExpr>)

atomicExch(addr, val)

AtomicCas(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>)

atomicCAS(addr, compare, val)

AtomicMax(Box<CudaExpr>, Box<CudaExpr>)

atomicMax(addr, val)

AtomicMin(Box<CudaExpr>, Box<CudaExpr>)

atomicMin(addr, val)

BinOp(Box<CudaExpr>, CudaBinOp, Box<CudaExpr>)

Binary operation: a + b

UnOp(CudaUnOp, Box<CudaExpr>)

Unary operation: !a

Index(Box<CudaExpr>, Box<CudaExpr>)

Array subscript: arr[idx]

Member(Box<CudaExpr>, String)

Struct member access: s.field

PtrMember(Box<CudaExpr>, String)

Pointer member access: p->field

Cast(CudaType, Box<CudaExpr>)

C-style cast: (T)expr

Call(String, Vec<CudaExpr>)

Function call: func(args...)

Ternary(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>)

Ternary conditional: cond ? then : else

Ldg(Box<CudaExpr>)

__ldg(&x) — read-only cache load

ShflDownSync(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>)

__shfl_down_sync(mask, var, delta)

ShflXorSync(Box<CudaExpr>, Box<CudaExpr>, Box<CudaExpr>)

__shfl_xor_sync(mask, var, laneMask)

WarpSize

warpSize builtin

BallotSync(Box<CudaExpr>, Box<CudaExpr>)

__ballot_sync(mask, predicate)

Popc(Box<CudaExpr>)

__popc(x) — popcount

Trait Implementations

impl Clone for CudaExpr

fn clone(&self) -> CudaExpr

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for CudaExpr

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl PartialEq for CudaExpr

fn eq(&self, other: &CudaExpr) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl StructuralPartialEq for CudaExpr