Repository: CrossGL/crosstl
Pull request #285 (2025-06-18 00:38:49)

Cuda support V0


Patch:
diff --git a/MANIFEST.in b/MANIFEST.in
index e09b5cd..58f676b 100644
--- a/MANIFEST.in
+++ b/MANIFEST.in
@@ -1,8 +1,13 @@
 include README.md
 include LICENSE
+include examples/*.cgl
 include backend/DirectX/*.py
 include backend/Metal/*.py
-include backend/Opengl/*.py
+include backend/GLSL/*.py
+include backend/CUDA/*.py
+include backend/Mojo/*.py
+include backend/Rust/*.py
+include backend/SPIRV/*.py
 include translator/*.py
 include translator/codegen/*.py
 include crosstl.py
diff --git a/README.md b/README.md
index 3355b52..3427bbd 100644
--- a/README.md
+++ b/README.md
@@ -68,6 +68,7 @@ Imagine writing a shader _once_ and deploying it across:
 - ⚙️ **Slang**
 - 🔥 **Mojo**
 - 🦀 **Rust**
+- 🚀 **CUDA**
 
 ...all without changing a single line of code!
 
@@ -87,6 +88,7 @@ CrossTL now supports comprehensive translation to and from these major graphics
 - **Slang** - Real-time shading language
 - **Mojo** - High-performance systems programming language
 - **Rust** - Systems programming language with GPU computing support
+- **CUDA** - NVIDIA's parallel computing platform and programming model
 
 ### Universal Format
 
@@ -208,6 +210,9 @@ rust_code = crosstl.translate('shader.cgl', backend='rust', save_shader='shader.
 
 # Translate to Slang
 slang_code = crosstl.translate('shader.cgl', backend='slang', save_shader='shader.slang')
+
+# Translate to CUDA
+cuda_code = crosstl.translate('shader.cgl', backend='cuda', save_shader='shader.cu')
 ```
 
 #### Converting from HLSL to CrossGL
@@ -359,6 +364,46 @@ crossgl_code = crosstl.translate('shader.mojo', backend='cgl', save_shader='shad
 print(crossgl_code)
 ```
 
+#### Converting from CUDA to CrossGL
+
+1. Write your CUDA shader (e.g., `shader.cu`):
+
+```cuda
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+
+struct VertexInput {
+    float3 position;
+    float2 texCoord;
+};
+
+struct VertexOutput {
+    float4 position;
+    float2 uv;
+};
+
+__global__ void compute_main() {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    // Compute shader logic here
+}
+
+__device__ VertexOutput vertex_main(VertexInput input) {
+    VertexOutput output;
+    output.position = make_float4(input.position, 1.0f);
+    output.uv = input.texCoord;
+    return output;
+}
+```
+
+2. Convert to CrossGL:
+
+```python
+import crosstl
+
+crossgl_code = crosstl.translate('shader.cu', backend='cgl', save_shader='shader.cgl')
+print(crossgl_code)
+```
+
 ## 🎯 Complete Translation Example
 
 Here's a complete example showing how to translate one CrossGL shader to all supported backends:
@@ -377,7 +422,8 @@ backends = {
     'vulkan': '.spirv',
     'mojo': '.mojo',
     'rust': '.rs',
-    'slang': '.slang'
+    'slang': '.slang',
+    'cuda': '.cu'
 }
 
 for backend, ext in backends.items():
@@ -399,7 +445,8 @@ conversions = [
     ('shader.hlsl', 'shader_from_hlsl.cgl'),
     ('shader.glsl', 'shader_from_glsl.cgl'),
     ('shader.rs', 'shader_from_rust.cgl'),
-    ('shader.mojo', 'shader_from_mojo.cgl')
+    ('shader.mojo', 'shader_from_mojo.cgl'),
+    ('shader.cu', 'shader_from_cuda.cgl')
 ]
 
 for input_file, output_file in conversions:
diff --git a/crosstl/_crosstl.py b/crosstl/_crosstl.py
index 4d78b79..4a7abd8 100644
--- a/crosstl/_crosstl.py
+++ b/crosstl/_crosstl.py
@@ -8,6 +8,7 @@ from .translator.codegen import (
     SPIRV_codegen,
     mojo_codegen,
     rust_codegen,
+    cuda_codegen,
 )
 from .translator.ast import ASTNode
 import argparse
@@ -83,6 +84,15 @@ def translate(
 
         lexer = RustLexer(shader_code)
         parser = RustParser(lexer.tokenize())
+    elif (
+        file_path.endswith(".cu")
+        or file_path.endswith(".cuh")
+        or file_path.endswith(".cuda")
+    ):
+        from .backend.CUDA import CudaLexer, CudaParser
+
+        lexer = CudaLexer(shader_code)
+        parser = CudaParser(lexer.tokenize())
     else:
         raise ValueError(f"Unsupported shader file type: {file_path}")
 
@@ -101,6 +111,8 @@ def translate(
             codegen = mojo_codegen.MojoCodeGen()
         elif backend == "rust":
             codegen = rust_codegen.RustCodeGen()
+        elif backend == "cuda":
+            codegen = cuda_codegen.CudaCodeGen()
         else:
             raise ValueError(f"Unsupported backend for CrossGL file: {backend}")
     else:
@@ -131,6 +143,14 @@ def translate(
                 from .backend.Rust.RustCrossGLCodeGen import RustToCrossGLConverter
 
                 codegen = RustToCrossGLConverter()
+            elif (
+                file_path.endswith(".cu")
+                or file_path.endswith(".cuh")
+                or file_path.endswith(".cuda")
+            ):
+                from .backend.CUDA.CudaCrossGLCodeGen import CudaToCrossGLConverter
+
+                codegen = CudaToCrossGLConverter()
             else:
                 raise ValueError(f"Reverse translation not supported for: {file_path}")
         else:
@@ -162,7 +182,7 @@ def main():
         "--backend",
         "-b",
         default="cgl",
-        help="Target backend (metal, directx, opengl, vulkan, mojo, rust, cgl)",
+        help="Target backend (metal, directx, opengl, vulkan, mojo, rust, cuda, cgl)",
     )
     parser.add_argument("--output", "-o", help="Output file path")
     parser.add_argument(
@@ -187,6 +207,7 @@ def main():
                 "vulkan": ".spirv",
                 "mojo": ".mojo",
                 "rust": ".rs",
+                "cuda": ".cu",
                 "cgl": ".cgl",
             }
             output_path = base + ext_map.get(args.backend, ".out")
diff --git a/crosstl/backend/CUDA/CudaAst.py b/crosstl/backend/CUDA/CudaAst.py
new file mode 100644
index 0000000..a2a71da
--- /dev/null
+++ b/crosstl/backend/CUDA/CudaAst.py
@@ -0,0 +1,375 @@
+"""CUDA AST Node definitions"""
+
+
+class ASTNode:
+    """Base class for all AST nodes"""
+
+
+class ShaderNode(ASTNode):
+    """Root node representing a complete CUDA program"""
+
+    def __init__(
+        self,
+        includes=None,
+        functions=None,
+        structs=None,
+        global_variables=None,
+        kernels=None,
+    ):
+        self.includes = includes or []
+        self.functions = functions or []
+        self.structs = structs or []
+        self.global_variables = global_variables or []
+        self.kernels = kernels or []
+
+    def __repr__(self):
+        return f"ShaderNode(includes={self.includes}, functions={self.functions}, structs={self.structs}, global_variables={self.global_variables}, kernels={self.kernels})"
+
+
+class FunctionNode(ASTNode):
+    """Node representing a function declaration"""
+
+    def __init__(
+        self, return_type, name, params, body, qualifiers=None, attributes=None
+    ):
+        self.return_type = return_type
+        self.name = name
+        self.params = params
+        self.body = body
+        self.qualifiers = qualifiers or []  # __global__, __device__, __host__
+        self.attributes = attributes or []
+
+    def __repr__(self):
+        return f"FunctionNode(return_type={self.return_type}, name={self.name}, params={self.params}, body={self.body}, qualifiers={self.qualifiers})"
+
+
+class KernelNode(FunctionNode):
+    """Node representing a CUDA kernel function (marked with __global__)"""
+
+    def __init__(self, return_type, name, params, body, attributes=None):
+        super().__init__(return_type, name, params, body, ["__global__"], attributes)
+
+    def __repr__(self):
+        return f"KernelNode(name={self.name}, params={self.params}, body={self.body})"
+
+
+class KernelLaunchNode(ASTNode):
+    """Node representing a kernel launch: kernel<<<blocks, threads>>>(args)"""
+
+    def __init__(
+        self, kernel_name, blocks, threads, shared_mem=None, stream=None, args=None
+    ):
+        self.kernel_name = kernel_name
+        self.blocks = blocks
+        self.threads = threads
+        self.shared_mem = shared_mem  # Optional shared memory size
+        self.stream = stream  # Optional stream
+        self.args = args or []
+
+    def __repr__(self):
+        return f"KernelLaunchNode(kernel_name={self.kernel_name}, blocks={self.blocks}, threads={self.threads}, args={self.args})"
+
+
+class StructNode(ASTNode):
+    """Node representing a struct declaration"""
+
+    def __init__(self, name, members, attributes=None):
+        self.name = name
+        self.members = members
+        self.attributes = attributes or []
+
+    def __repr__(self):
+        return f"StructNode(name={self.name}, members={self.members})"
+
+
+class VariableNode(ASTNode):
+    """Node representing a variable declaration"""
+
+    def __init__(self, vtype, name, value=None, qualifiers=None):
+        self.vtype = vtype
+        self.name = name
+        self.value = value
+        self.qualifiers = qualifiers or []  # __shared__, __constant__, etc.
+
+    def __repr__(self):
+        return f"VariableNode(vtype={self.vtype}, name={self.name}, value={self.value}, qualifiers={self.qualifiers})"
+
+
+class AssignmentNode(ASTNode):
+    """Node representing an assignment operation"""
+
+    def __init__(self, left, right, operator="="):
+        self.left = left
+        self.right = right
+        self.operator = operator
+
+    def __repr__(self):
+        return f"AssignmentNode(left={self.left}, operator={self.operator}, right={self.right})"
+
+
+class BinaryOpNode(ASTNode):
+    """Node representing a binary operation"""
+
+    def __init__(self, left, op, right):
+        self.left = left
+        self.op = op
+        self.right = right
+
+    def __repr__(self):
+        return f"BinaryOpNode(left={self.left}, op={self.op}, right={self.right})"
+
+
+class UnaryOpNode(ASTNode):
+    """Node representing a unary operation"""
+
+    def __init__(self, op, operand):
+        self.op = op
+        self.operand = operand
+
+    def __repr__(self):
+        return f"UnaryOpNode(op={self.op}, operand={self.operand})"
+
+
+class FunctionCallNode(ASTNode):
+    """Node representing a function call"""
+
+    def __init__(self, name, args):
+        self.name = name
+        self.args = args
+
+    def __repr__(self):
+        return f"FunctionCallNode(name={self.name}, args={self.args})"
+
+
+class AtomicOperationNode(FunctionCallNode):
+    """Node representing a CUDA atomic operation"""
+
+    def __init__(self, operation, args):
+        super().__init__(operation, args)
+        self.operation = operation  # atomicAdd, atomicSub, etc.
+
+    def __repr__(self):
+        return f"AtomicOperationNode(operation={self.operation}, args={self.args})"
+
+
+class SyncNode(ASTNode):
+    """Node representing synchronization operations"""
+
+    def __init__(self, sync_type, args=None):
+        self.sync_type = sync_type  # __syncthreads, __syncwarp
+        self.args = args or []
+
+    def __repr__(self):
+        return f"SyncNode(sync_type={self.sync_type}, args={self.args})"
+
+
+class MemberAccessNode(ASTNode):
+    """Node representing member access (dot or arrow operator)"""
+
+    def __init__(self, object, member, is_pointer=False):
+        self.object = object
+        self.member = member
+        self.is_pointer = is_pointer  # True for ->, False for .
+
+    def __repr__(self):
+        op = "->" if self.is_pointer else "."
+        return f"MemberAccessNode(object={self.object}, member={self.member}, operator={op})"
+
+
+class ArrayAccessNode(ASTNode):
+    """Node representing array access"""
+
+    def __init__(self, array, index):
+        self.array = array
+        self.index = index
+
+    def __repr__(self):
+        return f"ArrayAccessNode(array={self.array}, index={self.index})"
+
+
+class IfNode(ASTNode):
+    """Node representing an if statement"""
+
+    def __init__(self, condition, if_body, else_body=None):
+        self.condition = condition
+        self.if_body = if_body
+        self.else_body = else_body
+
+    def __repr__(self):
+        return f"IfNode(condition={self.condition}, if_body={self.if_body}, else_body={self.else_body})"
+
+
+class ForNode(ASTNode):
+    """Node representing a for loop"""
+
+    def __init__(self, init, condition, update, body):
+        self.init = init
+        self.condition = condition
+        self.update = update
+        self.body = body
+
+    def __repr__(self):
+        return f"ForNode(init={self.init}, condition={self.condition}, update={self.update}, body={self.body})"
+
+
+class WhileNode(ASTNode):
+    """Node representing a while loop"""
+
+    def __init__(self, condition, body):
+        self.condition = condition
+        self.body = body
+
+    def __repr__(self):
+        return f"WhileNode(condition={self.condition}, body={self.body})"
+
+
+class DoWhileNode(ASTNode):
+    """Node representing a do-while loop"""
+
+    def __init__(self, body, condition):
+        self.body = body
+        self.condition = condition
+
+    def __repr__(self):
+        return f"DoWhileNode(body={self.body}, condition={self.condition})"
+
+
+class SwitchNode(ASTNode):
+    """Node representing a switch statement"""
+
+    def __init__(self, expression, cases, default_case=None):
+        self.expression = expression
+        self.cases = cases
+        self.default_case = default_case
+
+    def __repr__(self):
+        return f"SwitchNode(expression={self.expression}, cases={self.cases}, default_case={self.default_case})"
+
+
+class CaseNode(ASTNode):
+    """Node representing a case in a switch statement"""
+
+    def __init__(self, value, body):
+        self.value = value
+        self.body = body
+
+    def __repr__(self):
+        return f"CaseNode(value={self.value}, body={self.body})"
+
+
+class ReturnNode(ASTNode):
+    """Node representing a return statement"""
+
+    def __init__(self, value=None):
+        self.value = value
+
+    def __repr__(self):
+        return f"ReturnNode(value={self.value})"
+
+
+class BreakNode(ASTNode):
+    """Node representing a break statement"""
+
+    def __repr__(self):
+        return "BreakNode()"
+
+
+class ContinueNode(ASTNode):
+    """Node representing a continue statement"""
+
+    def __repr__(self):
+        return "ContinueNode()"
+
+
+class VectorConstructorNode(ASTNode):
+    """Node representing CUDA vector constructor (make_float4, etc.)"""
+
+    def __init__(self, vector_type, args):
+        self.vector_type = vector_type
+        self.args = args
+
+    def __repr__(self):
+        return (
+            f"VectorConstructorNode(vector_type={self.vector_type}, args={self.args})"
+        )
+
+
+class TernaryOpNode(ASTNode):
+    """Node representing a ternary conditional operator"""
+
+    def __init__(self, condition, true_expr, false_expr):
+        self.condition = condition
+        self.true_expr = true_expr
+        self.false_expr = false_expr
+
+    def __repr__(self):
+        return f"TernaryOpNode(condition={self.condition}, true_expr={self.true_expr}, false_expr={self.false_expr})"
+
+
+class CastNode(ASTNode):
+    """Node representing a type cast"""
+
+    def __init__(self, target_type, expression):
+        self.target_type = target_type
+        self.expression = expression
+
+    def __repr__(self):
+        return f"CastNode(target_type={self.target_type}, expression={self.expression})"
+
+
+class PreprocessorNode(ASTNode):
+    """Node representing preprocessor directives"""
+
+    def __init__(self, directive, content):
+        self.directive = directive  # include, define, etc.
+        self.content = content
+
+    def __repr__(self):
+        return f"PreprocessorNode(directive={self.directive}, content={self.content})"
+
+
+class CudaBuiltinNode(ASTNode):
+    """Node representing CUDA built-in variables (threadIdx, blockIdx, etc.)"""
+
+    def __init__(self, builtin_name, component=None):
+        self.builtin_name = builtin_name  # threadIdx, blockIdx, etc.
+        self.component = component  # x, y, z component
+
+    def __repr__(self):
+        if self.component:
+            return f"CudaBuiltinNode(builtin_name={self.builtin_name}, component={self.component})"
+        return f"CudaBuiltinNode(builtin_name={self.builtin_name})"
+
+
+class TextureAccessNode(ASTNode):
+    """Node representing texture memory access"""
+
+    def __init__(self, texture_name, coordinates):
+        self.texture_name = texture_name
+        self.coordinates = coordinates
+
+    def __repr__(self):
+        return f"TextureAccessNode(texture_name={self.texture_name}, coordinates={self.coordinates})"
+
+
+class SharedMemoryNode(VariableNode):
+    """Node representing shared memory variable declaration"""
+
+    def __init__(self, vtype, name, size=None):
+        super().__init__(vtype, name, qualifiers=["__shared__"])
+        self.size = size  # For dynamic shared memory
+
+    def __repr__(self):
+        return (
+            f"SharedMemoryNode(vtype={self.vtype}, name={self.name}, size={self.size})"
+        )
+
+
+class ConstantMemoryNode(VariableNode):
+    """Node representing constant memory variable declaration"""
+
+    def __init__(self, vtype, name, value=None):
+        super().__init__(vtype, name, value, qualifiers=["__constant__"])
+
+    def __repr__(self):
+        return f"ConstantMemoryNode(vtype={self.vtype}, name={self.name}, value={self.value})"
diff --git a/crosstl/backend/CUDA/CudaCrossGLCodeGen.py b/crosstl/backend/CUDA/CudaCrossGLCodeGen.py
new file mode 100644
index 0000000..ef4ed48
--- /dev/null
+++ b/crosstl/backend/CUDA/CudaCrossGLCodeGen.py
@@ -0,0 +1,458 @@
+"""CUDA to CrossGL Code Generator"""
+
+from .CudaAst import *
+
+
+class CudaToCrossGLConverter:
+    """Converts CUDA AST to CrossGL format"""
+
+    def __init__(self):
+        self.indent_level = 0
+        self.output = []
+
+    def generate(self, ast_node):
+        """Generate CrossGL code from CUDA AST"""
+        self.output = []
+        self.indent_level = 0
+        self.visit(ast_node)
+        return "\n".join(self.output)
+
+    def visit(self, node):
+        """Visit a node and generate code"""
+        method_name = f"visit_{type(node).__name__}"
+        visitor = getattr(self, method_name, self.generic_visit)
+        return visitor(node)
+
+    def generic_visit(self, node):
+        """Default visitor for unknown nodes"""
+        if isinstance(node, str):
+            return node
+        elif isinstance(node, list):
+            return [self.visit(item) for item in node]
+        else:
+            return str(node)
+
+    def emit(self, code):
+        """Emit code with proper indentation"""
+        if code.strip():
+            self.output.append("    " * self.indent_level + code)
+        else:
+            self.output.append("")
+
+    def visit_ShaderNode(self, node):
+        """Visit shader node (main program)"""
+        # Always emit the comment first
+        self.emit("// CUDA to CrossGL conversion")
+
+        # Process includes
+        if hasattr(node, "includes") and node.includes:
+            for include in node.includes:
+                self.visit(include)
+            self.emit("")
+
+        # Process structs
+        if hasattr(node, "structs") and node.structs:
+            for struct in node.structs:
+                self.visit(struct)
+                self.emit("")
+
+        # Process global variables
+        if hasattr(node, "global_variables") and node.global_variables:
+            for var in node.global_variables:
+                self.visit(var)
+                self.emit("")
+
+        # Process functions (convert kernels to compute shaders)
+        if hasattr(node, "functions") and node.functions:
+            for func in node.functions:
+                self.emit(f"// Function: {func.name}")
+                self.visit(func)
+                self.emit("")
+
+        # Process kernels (convert to compute shaders)
+        if hasattr(node, "kernels") and node.kernels:
+            for kernel in node.kernels:
+                self.emit(f"// Kernel: {kernel.name}")
+                self.visit_kernel_as_compute_shader(kernel)
+                self.emit("")
+
+    def visit_PreprocessorNode(self, node):
+        """Visit preprocessor directive"""
+        if node.directive == "include":
+            # Convert CUDA includes to CrossGL equivalents
+            if "cuda_runtime.h" in node.content:
+                self.emit("// CUDA runtime functionality built-in")
+            elif "device_launch_parameters.h" in node.content:
+                self.emit("// Device parameters built-in")
+            else:
+                self.emit(f"// {node.directive} {node.content}")
+        else:
+            self.emit(f"// {node.directive} {node.content}")
+
+    def visit_StructNode(self, node):
+        """Visit struct declaration"""
+        self.emit(f"struct {node.name} {{")
+        self.indent_level += 1
+
+        for member in node.members:
+            member_type = self.convert_cuda_type_to_crossgl(member.vtype)
+            self.emit(f"{member_type} {member.name};")
+
+        self.indent_level -= 1
+        self.emit("};")
+
+    def visit_FunctionNode(self, node):
+        """Visit function declaration"""
+        # Skip device functions in CrossGL (they become inline)
+        if "__device__" in node.qualifiers:
+            return
+
+        return_type = self.convert_cuda_type_to_crossgl(node.return_type)
+        params = []
+
+        for param in node.params:
+            param_type = self.convert_cuda_type_to_crossgl(param.vtype)
+            params.append(f"{param_type} {param.name}")
+
+        param_str = ", ".join(params)
+        self.emit(f"{return_type} {node.name}({param_str}) {{")
+
+        self.indent_level += 1
+        for stmt in node.body:
+            self.visit(stmt)
+        self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_kernel_as_compute_shader(self, kernel):
+        """Convert CUDA kernel to CrossGL compute shader"""
+        # Generate compute shader layout
+        self.emit("@compute")
+        self.emit("@workgroup_size(1, 1, 1)  // Default workgroup size")
+
+        # Generate function signature
+        params = []
+        for param in kernel.params:
+            param_type = self.convert_cuda_type_to_crossgl(param.vtype)
+            # Add storage buffer qualifiers for pointer parameters
+            if "*" in param.vtype:
+                params.append(
+                    f"@group(0) @binding({len(params)}) var<storage, read_write> {param.name}: array<{param_type.replace('*', '').strip()}>"
+                )
+            else:
+                params.append(f"{param_type} {param.name}")
+
+        self.emit(f"fn {kernel.name}(")
+        self.indent_level += 1
+        for i, param in enumerate(params):
+            if i == len(params) - 1:
+                self.emit(f"{param}")
+            else:
+                self.emit(f"{param},")
+        self.indent_level -= 1
+        self.emit(") {")
+
+        # Add built-in variable declarations
+        self.indent_level += 1
+        self.emit("let thread_id = gl_GlobalInvocationID;")
+        self.emit("let block_id = gl_WorkGroupID;")
+        self.emit("let thread_local_id = gl_LocalInvocationID;")
+        self.emit("let block_dim = gl_WorkGroupSize;")
+        self.emit("")
+
+        # Process kernel body
+        for stmt in kernel.body:
+            self.visit(stmt)
+
+        self.indent_level -= 1
+        self.emit("}")
+
+    def visit_KernelLaunchNode(self, node):
+        """Visit kernel launch (convert to dispatch call comment)"""
+        self.emit(
+            f"// Kernel launch: {node.kernel_name}<<<{self.visit(node.blocks)}, {self.visit(node.threads)}>>>()"
+        )
+        if node.args:
+            args_str = ", ".join([self.visit(arg) for arg in node.args])
+            self.emit(f"// Arguments: {args_str}")
+
+    def visit_VariableNode(self, node):
+        """Visit variable declaration"""
+        var_type = self.convert_cuda_type_to_crossgl(node.vtype)
+
+        if node.value:
+            value = self.visit(node.value)
+            self.emit(f"var {node.name}: {var_type} = {value};")
+        else:
+            self.emit(f"var {node.name}: {var_type};")
+
+    def visit_SharedMemoryNode(self, node):
+        """Visit shared memory declaration"""
+        # Convert to workgroup memory in CrossGL
+        var_type = self.convert_cuda_type_to_crossgl(node.vtype)
+        if node.size:
+            size = self.visit(node.size)
+            self.emit(f"var<workgroup> {node.name}: array<{var_type}, {size}>;")
+        else:
+            self.emit(f"var<workgroup> {node.name}: {var_type};")
+
+    def visit_ConstantMemoryNode(self, node):
+        """Visit constant memory declaration"""
+        # Convert to uniform buffer in CrossGL
+        var_type = self.convert_cuda_type_to_crossgl(node.vtype)
+        if node.value:
+            value = self.visit(node.value)
+            self.emit(
+                f"@group(0) @binding(0) var<uniform> {node.name}: {var_type} = {value};"
+            )
+        else:
+            self.emit(f"@group(0) @binding(0) var<uniform> {node.name}: {var_type};")
+
+    def visit_AssignmentNode(self, node):
+        """Visit assignment"""
+        left = self.visit(node.left)
+        right = self.visit(node.right)
+        self.emit(f"{left} {node.operator} {right};")
+
+    def visit_BinaryOpNode(self, node):
+        """Visit binary operation"""
+        left = self.visit(node.left)
+        right = self.visit(node.right)
+        return f"({left} {node.op} {right})"
+
+    def visit_UnaryOpNode(self, node):
+        """Visit unary operation"""
+        operand = self.visit(node.operand)
+        if node.op.startswith("post"):
+            return f"({operand}{node.op[4:]})"
+        else:
+            return f"({node.op}{operand})"
+
+    def visit_FunctionCallNode(self, node):
+        """Visit function call"""
+        args = [self.visit(arg) for arg in node.args]
+        args_str = ", ".join(args)
+
+        # Convert CUDA built-in functions
+        func_name = self.convert_cuda_builtin_function(node.name)
+        return f"{func_name}({args_str})"
+
+    def visit_AtomicOperationNode(self, node):
+        """Visit atomic operation"""
+        args = [self.visit(arg) for arg in node.args]
+        args_str = ", ".join(args)
+
+        # Convert CUDA atomic operations to CrossGL equivalents
+        atomic_map = {
+            "atomicAdd": "atomicAdd",
+            "atomicSub": "atomicSub",
+            "atomicMax": "atomicMax",
+            "atomicMin": "atomicMin",
+            "atomicExch": "atomicExchange",
+            "atomicCAS": "atomicCompareExchange",
+        }
+
+        crossgl_func = atomic_map.get(node.operation, node.operation)
+        return f"{crossgl_func}({args_str})"
+
+    def visit_SyncNode(self, node):
+        """Visit synchronization"""
+        if node.sync_type == "__syncthreads":
+            self.emit("workgroupBarrier();")
+        elif node.sync_type == "__syncwarp":
+            self.emit("// Warp sync not directly supported in CrossGL")
+        else:
+            self.emit(f"// {node.sync_type}();")
+
+    def visit_CudaBuiltinNode(self, node):
+        """Visit CUDA built-in variables"""
+        builtin_map = {
+            "threadIdx": "gl_LocalInvocationID",
+            "blockIdx": "gl_WorkGroupID",
+            "gridDim": "gl_NumWorkGroups",
+            "blockDim": "gl_WorkGroupSize",
+            "warpSize": "32",  # Constant in CrossGL
+        }
+
+        base_name = builtin_map.get(node.builtin_name, node.builtin_name)
+        if node.component:
+            return f"{base_name}.{node.component}"
+        else:
+            return base_name
+
+    def visit_MemberAccessNode(self, node):
+        """Visit member access"""
+        obj = self.visit(node.object)
+        if node.is_pointer:
+            # Convert pointer access to direct access in CrossGL
+            return f"{obj}.{node.member}"
+        else:
+            return f"{obj}.{node.member}"
+
+    def visit_ArrayAccessNode(self, node):
+        """Visit array access"""
+        array = self.visit(node.array)
+        index = self.visit(node.index)
+        return f"{array}[{index}]"
+
+    def visit_IfNode(self, node):
+        """Visit if statement"""
+        condition = self.visit(node.condition)
+        self.emit(f"if ({condition}) {{")
+
+        self.indent_level += 1
+        if isinstance(node.if_body, list):
+            for stmt in node.if_body:
+                self.visit(stmt)
+        else:
+            self.visit(node.if_body)
+        self.indent_level -= 1
+
+        if node.else_body:
+            self.emit("} else {")
+            self.indent_level += 1
+            if isinstance(node.else_body, list):
+                for stmt in node.else_body:
+                    self.visit(stmt)
+            else:
+                self.visit(node.else_body)
+            self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_ForNode(self, node):
+        """Visit for loop"""
+        init = self.visit(node.init) if node.init else ""
+        condition = self.visit(node.condition) if node.condition else ""
+        update = self.visit(node.update) if node.update else ""
+
+        self.emit(f"for ({init}; {condition}; {update}) {{")
+
+        self.indent_level += 1
+        if isinstance(node.body, list):
+            for stmt in node.body:
+                self.visit(stmt)
+        else:
+            self.visit(node.body)
+        self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_WhileNode(self, node):
+        """Visit while loop"""
+        condition = self.visit(node.condition)
+        self.emit(f"while ({condition}) {{")
+
+        self.indent_level += 1
+        if isinstance(node.body, list):
+            for stmt in node.body:
+                self.visit(stmt)
+        else:
+            self.visit(node.body)
+        self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_ReturnNode(self, node):
+        """Visit return statement"""
+        if node.value:
+            value = self.visit(node.value)
+            self.emit(f"return {value};")
+        else:
+            self.emit("return;")
+
+    def visit_BreakNode(self, node):
+        """Visit break statement"""
+        self.emit("break;")
+
+    def visit_ContinueNode(self, node):
+        """Visit continue statement"""
+        self.emit("continue;")
+
+    def convert_cuda_type_to_crossgl(self, cuda_type):
+        """Convert CUDA types to CrossGL equivalents"""
+        type_mapping = {
+            # Basic types
+            "void": "void",
+            "bool": "bool",
+            "char": "i8",
+            "unsigned char": "u8",
+            "short": "i16",
+            "unsigned short": "u16",
+            "int": "i32",
+            "unsigned int": "u32",
+            "long": "i64",
+            "unsigned long": "u64",
+            "float": "f32",
+            "double": "f64",
+            "size_t": "u32",
+            # CUDA vector types
+            "float2": "vec2<f32>",
+            "float3": "vec3<f32>",
+            "float4": "vec4<f32>",
+            "double2": "vec2<f64>",
+            "double3": "vec3<f64>",
+            "double4": "vec4<f64>",
+            "int2": "vec2<i32>",
+            "int3": "vec3<i32>",
+            "int4": "vec4<i32>",
+            "uint2": "vec2<u32>",
+            "uint3": "vec3<u32>",
+            "uint4": "vec4<u32>",
+        }
+
+        # Handle pointers
+        if "*" in cuda_type:
+            base_type = cuda_type.replace("*", "").strip()
+            mapped_base = type_mapping.get(base_type, base_type)
+            return f"ptr<{mapped_base}>"
+
+        # Handle arrays
+        if "[" in cuda_type and "]" in cuda_type:
+            parts = cuda_type.split("[")
+            base_type = parts[0].strip()
+            size = parts[1].split("]")[0]
+            mapped_base = type_mapping.get(base_type, base_type)
+            return f"array<{mapped_base}, {size}>"
+
+        return type_mapping.get(cuda_type, cuda_type)
+
+    def convert_cuda_builtin_function(self, func_name):
+        """Convert CUDA built-in functions to CrossGL equivalents"""
+        function_mapping = {
+            # Math functions
+            "sqrtf": "sqrt",
+            "powf": "pow",
+            "sinf": "sin",
+            "cosf": "cos",
+            "tanf": "tan",
+            "logf": "log",
+            "expf": "exp",
+            "fabsf": "abs",
+            "fminf": "min",
+            "fmaxf": "max",
+            "floorf": "floor",
+            "ceilf": "ceil",
+            # Double precision variants
+            "sqrt": "sqrt",
+            "pow": "pow",
+            "sin": "sin",
+            "cos": "cos",
+            "tan": "tan",
+            "log": "log",
+            "exp": "exp",
+            "fabs": "abs",
+            "fmin": "min",
+            "fmax": "max",
+            "floor": "floor",
+            "ceil": "ceil",
+            # Vector functions
+            "make_float2": "vec2<f32>",
+            "make_float3": "vec3<f32>",
+            "make_float4": "vec4<f32>",
+            "make_int2": "vec2<i32>",
+            "make_int3": "vec3<i32>",
+            "make_int4": "vec4<i32>",
+        }
+
+        return function_mapping.get(func_name, func_name)
diff --git a/crosstl/backend/CUDA/CudaLexer.py b/crosstl/backend/CUDA/CudaLexer.py
new file mode 100644
index 0000000..efdee63
--- /dev/null
+++ b/crosstl/backend/CUDA/CudaLexer.py
@@ -0,0 +1,500 @@
+import re
+from typing import Iterator, Tuple, List
+from enum import Enum, auto
+
+# Skip tokens that don't need processing
+SKIP_TOKENS = {"WHITESPACE", "COMMENT_SINGLE", "COMMENT_MULTI"}
+
+# CUDA token definitions - order matters for correct tokenization
+TOKENS = tuple(
+    [
+        # Comments
+        ("COMMENT_SINGLE", r"//.*"),
+        ("COMMENT_MULTI", r"/\*[\s\S]*?\*/"),
+        # CUDA execution configuration (must come before operators)
+        ("KERNEL_LAUNCH_START", r"<<<"),
+        ("KERNEL_LAUNCH_END", r">>>"),
+        # CUDA keywords and qualifiers
+        ("GLOBAL", r"\b__global__\b"),
+        ("DEVICE", r"\b__device__\b"),
+        ("HOST", r"\b__host__\b"),
+        ("SHARED", r"\b__shared__\b"),
+        ("CONSTANT", r"\b__constant__\b"),
+        ("RESTRICT", r"\b__restrict__\b"),
+        ("MANAGED", r"\b__managed__\b"),
+        ("NOINLINE", r"\b__noinline__\b"),
+        ("FORCEINLINE", r"\b__forceinline__\b"),
+        # CUDA built-in variables
+        ("THREADIDX", r"\bthreadIdx\b"),
+        ("BLOCKIDX", r"\bblockIdx\b"),
+        ("GRIDDIM", r"\bgridDim\b"),
+        ("BLOCKDIM", r"\bblockDim\b"),
+        ("WARPSIZE", r"\bwarpSize\b"),
+        # CUDA built-in functions
+        ("SYNCTHREADS", r"\b__syncthreads\b"),
+        ("SYNCWARP", r"\b__syncwarp\b"),
+        ("ATOMICADD", r"\batomicAdd\b"),
+        ("ATOMICSUB", r"\batomicSub\b"),
+        ("ATOMICMAX", r"\batomicMax\b"),
+        ("ATOMICMIN", r"\batomicMin\b"),
+        ("ATOMICEXCH", r"\batomicExch\b"),
+        ("ATOMICCAS", r"\batomicCAS\b"),
+        # Standard C/C++ keywords
+        ("TYPEDEF", r"\btypedef\b"),
+        ("STRUCT", r"\bstruct\b"),
+        ("UNION", r"\bunion\b"),
+        ("ENUM", r"\benum\b"),
+        ("CLASS", r"\bclass\b"),
+        ("NAMESPACE", r"\bnamespace\b"),
+        ("TEMPLATE", r"\btemplate\b"),
+        ("TYPENAME", r"\btypename\b"),
+        ("EXTERN", r"\bextern\b"),
+        ("STATIC", r"\bstatic\b"),
+        ("INLINE", r"\binline\b"),
+        ("CONST", r"\bconst\b"),
+        ("VOLATILE", r"\bvolatile\b"),
+        ("MUTABLE", r"\bmutable\b"),
+        ("VIRTUAL", r"\bvirtual\b"),
+        ("PUBLIC", r"\bpublic\b"),
+        ("PRIVATE", r"\bprivate\b"),
+        ("PROTECTED", r"\bprotected\b"),
+        # Control flow
+        ("IF", r"\bif\b"),
+        ("ELSE", r"\belse\b"),
+        ("FOR", r"\bfor\b"),
+        ("WHILE", r"\bwhile\b"),
+        ("DO", r"\bdo\b"),
+        ("SWITCH", r"\bswitch\b"),
+        ("CASE", r"\bcase\b"),
+        ("DEFAULT", r"\bdefault\b"),
+        ("BREAK", r"\bbreak\b"),
+        ("CONTINUE", r"\bcontinue\b"),
+        ("RETURN", r"\breturn\b"),
+        ("GOTO", r"\bgoto\b"),
+        # CUDA and C++ types
+        ("VOID", r"\bvoid\b"),
+        ("BOOL", r"\bbool\b"),
+        ("CHAR", r"\bchar\b"),
+        ("SHORT", r"\bshort\b"),
+        ("INT", r"\bint\b"),
+        ("LONG", r"\blong\b"),
+        ("FLOAT", r"\bfloat\b"),
+        ("DOUBLE", r"\bdouble\b"),
+        ("SIGNED", r"\bsigned\b"),
+        ("UNSIGNED", r"\bunsigned\b"),
+        ("SIZE_T", r"\bsize_t\b"),
+        # CUDA vector types
+        ("FLOAT2", r"\bfloat2\b"),
+        ("FLOAT3", r"\bfloat3\b"),
+        ("FLOAT4", r"\bfloat4\b"),
+        ("DOUBLE2", r"\bdouble2\b"),
+        ("DOUBLE3", r"\bdouble3\b"),
+        ("DOUBLE4", r"\bdouble4\b"),
+        ("INT2", r"\bint2\b"),
+        ("INT3", r"\bint3\b"),
+        ("INT4", r"\bint4\b"),
+        ("UINT2", r"\buint2\b"),
+        ("UINT3", r"\buint3\b"),
+        ("UINT4", r"\buint4\b"),
+        ("CHAR2", r"\bchar2\b"),
+        ("CHAR3", r"\bchar3\b"),
+        ("CHAR4", r"\bchar4\b"),
+        ("UCHAR2", r"\buchar2\b"),
+        ("UCHAR3", r"\buchar3\b"),
+        ("UCHAR4", r"\buchar4\b"),
+        ("SHORT2", r"\bshort2\b"),
+        ("SHORT3", r"\bshort3\b"),
+        ("SHORT4", r"\bshort4\b"),
+        ("USHORT2", r"\bushort2\b"),
+        ("USHORT3", r"\bushort3\b"),
+        ("USHORT4", r"\bushort4\b"),
+        ("LONG2", r"\blong2\b"),
+        ("LONG3", r"\blong3\b"),
+        ("LONG4", r"\blong4\b"),
+        ("ULONG2", r"\bulong2\b"),
+        ("ULONG3", r"\bulong3\b"),
+        ("ULONG4", r"\bulong4\b"),
+        ("LONGLONG2", r"\blonglong2\b"),
+        ("LONGLONG3", r"\blonglong3\b"),
+        ("LONGLONG4", r"\blonglong4\b"),
+        ("ULONGLONG2", r"\bulonglong2\b"),
+        ("ULONGLONG3", r"\bulonglong3\b"),
+        ("ULONGLONG4", r"\bulonglong4\b"),
+        # CUDA texture types
+        ("TEXTURE", r"\btexture\b"),
+        ("SURFACE", r"\bsurface\b"),
+        ("CUDAARRAY", r"\bcudaArray\b"),
+        ("CUDAARRAYT", r"\bcudaArray_t\b"),
+        # Boolean literals
+        ("TRUE", r"\btrue\b"),
+        ("FALSE", r"\bfalse\b"),
+        ("NULL", r"\bNULL\b"),
+        ("NULLPTR", r"\bnullptr\b"),
+        # Identifiers and literals (must come after keywords)
+        ("IDENTIFIER", r"[a-zA-Z_][a-zA-Z0-9_]*"),
+        ("NUMBER", r"\d+(\.\d+)?([fFdDlLuU]*|[eE][+-]?\d+[fFdDlLuU]*)?"),
+        ("STRING", r'"([^"\\]|\\.)*"'),
+        ("CHAR_LIT", r"'([^'\\]|\\.)'"),
+        # Preprocessor
+        ("PREPROCESSOR", r"#[^\n]*"),
+        # Operators (multi-character first)
+        ("SCOPE", r"::"),
+        ("SHIFT_LEFT", r"<<"),
+        ("SHIFT_RIGHT", r">>"),
+        ("PLUS_EQUALS", r"\+="),
+        ("MINUS_EQUALS", r"-="),
+        ("MULTIPLY_EQUALS", r"\*="),
+        ("DIVIDE_EQUALS", r"/="),
+        ("MODULO_EQUALS", r"%="),
+        ("AND_EQUALS", r"&="),
+        ("OR_EQUALS", r"\|="),
+        ("XOR_EQUALS", r"\^="),
+        ("SHIFT_LEFT_EQUALS", r"<<="),
+        ("SHIFT_RIGHT_EQUALS", r">>="),
+        ("LOGICAL_AND", r"&&"),
+        ("LOGICAL_OR", r"\|\|"),
+        ("EQUAL", r"=="),
+        ("NOT_EQUAL", r"!="),
+        ("LESS_EQUAL", r"<="),
+        ("GREATER_EQUAL", r">="),
+        ("INCREMENT", r"\+\+"),
+        ("DECREMENT", r"--"),
+        ("ARROW", r"->"),
+        ("DOT", r"\."),
+        # Single character operators
+        ("PLUS", r"\+"),
+        ("MINUS", r"-"),
+        ("MULTIPLY", r"\*"),
+        ("DIVIDE", r"/"),
+        ("MODULO", r"%"),
+        ("ASSIGN", r"="),
+        ("LESS_THAN", r"<"),
+        ("GREATER_THAN", r">"),
+        ("LOGICAL_NOT", r"!"),
+        ("BITWISE_AND", r"&"),
+        ("BITWISE_OR", r"\|"),
+        ("BITWISE_XOR", r"\^"),
+        ("BITWISE_NOT", r"~"),
+        ("QUESTION", r"\?"),
+        ("COLON", r":"),
+        # Delimiters
+        ("LBRACE", r"\{"),
+        ("RBRACE", r"\}"),
+        ("LPAREN", r"\("),
+        ("RPAREN", r"\)"),
+        ("LBRACKET", r"\["),
+        ("RBRACKET", r"\]"),
+        ("SEMICOLON", r";"),
+        ("COMMA", r","),
+        # Whitespace (must be last)
+        ("WHITESPACE", r"\s+"),
+    ]
+)
+
+# Keywords mapping for reserved word detection
+KEYWORDS = {
+    "__global__": "GLOBAL",
+    "__device__": "DEVICE",
+    "__host__": "HOST",
+    "__shared__": "SHARED",
+    "__constant__": "CONSTANT",
+    "__restrict__": "RESTRICT",
+    "__managed__": "MANAGED",
+    "__noinline__": "NOINLINE",
+    "__forceinline__": "FORCEINLINE",
+    "threadIdx": "THREADIDX",
+    "blockIdx": "BLOCKIDX",
+    "gridDim": "GRIDDIM",
+    "blockDim": "BLOCKDIM",
+    "warpSize": "WARPSIZE",
+    "__syncthreads": "SYNCTHREADS",
+    "__syncwarp": "SYNCWARP",
+    "atomicAdd": "ATOMICADD",
+    "atomicSub": "ATOMICSUB",
+    "atomicMax": "ATOMICMAX",
+    "atomicMin": "ATOMICMIN",
+    "atomicExch": "ATOMICEXCH",
+    "atomicCAS": "ATOMICCAS",
+    "typedef": "TYPEDEF",
+    "struct": "STRUCT",
+    "union": "UNION",
+    "enum": "ENUM",
+    "class": "CLASS",
+    "namespace": "NAMESPACE",
+    "template": "TEMPLATE",
+    "typename": "TYPENAME",
+    "extern": "EXTERN",
+    "static": "STATIC",
+    "inline": "INLINE",
+    "const": "CONST",
+    "volatile": "VOLATILE",
+    "mutable": "MUTABLE",
+    "virtual": "VIRTUAL",
+    "public": "PUBLIC",
+    "private": "PRIVATE",
+    "protected": "PROTECTED",
+    "if": "IF",
+    "else": "ELSE",
+    "for": "FOR",
+    "while": "WHILE",
+    "do": "DO",
+    "switch": "SWITCH",
+    "case": "CASE",
+    "default": "DEFAULT",
+    "break": "BREAK",
+    "continue": "CONTINUE",
+    "return": "RETURN",
+    "goto": "GOTO",
+    "void": "VOID",
+    "bool": "BOOL",
+    "char": "CHAR",
+    "short": "SHORT",
+    "int": "INT",
+    "long": "LONG",
+    "float": "FLOAT",
+    "double": "DOUBLE",
+    "signed": "SIGNED",
+    "unsigned": "UNSIGNED",
+    "size_t": "SIZE_T",
+    "true": "TRUE",
+    "false": "FALSE",
+    "NULL": "NULL",
+    "nullptr": "NULLPTR",
+}
+
+
+class TokenType(Enum):
+    COMMENT_SINGLE = auto()
+    COMMENT_MULTI = auto()
+    KERNEL_LAUNCH_START = auto()
+    KERNEL_LAUNCH_END = auto()
+    GLOBAL = auto()
+    DEVICE = auto()
+    HOST = auto()
+    SHARED = auto()
+    CONSTANT = auto()
+    RESTRICT = auto()
+    MANAGED = auto()
+    NOINLINE = auto()
+    FORCEINLINE = auto()
+    THREADIDX = auto()
+    BLOCKIDX = auto()
+    GRIDDIM = auto()
+    BLOCKDIM = auto()
+    WARPSIZE = auto()
+    SYNCTHREADS = auto()
+    SYNCWARP = auto()
+    ATOMICADD = auto()
+    ATOMICSUB = auto()
+    ATOMICMAX = auto()
+    ATOMICMIN = auto()
+    ATOMICEXCH = auto()
+    ATOMICCAS = auto()
+    TYPEDEF = auto()
+    STRUCT = auto()
+    UNION = auto()
+    ENUM = auto()
+    CLASS = auto()
+    NAMESPACE = auto()
+    TEMPLATE = auto()
+    TYPENAME = auto()
+    EXTERN = auto()
+    STATIC = auto()
+    INLINE = auto()
+    CONST = auto()
+    VOLATILE = auto()
+    MUTABLE = auto()
+    VIRTUAL = auto()
+    PUBLIC = auto()
+    PRIVATE = auto()
+    PROTECTED = auto()
+    IF = auto()
+    ELSE = auto()
+    FOR = auto()
+    WHILE = auto()
+    DO = auto()
+    SWITCH = auto()
+    CASE = auto()
+    DEFAULT = auto()
+    BREAK = auto()
+    CONTINUE = auto()
+    RETURN = auto()
+    GOTO = auto()
+    VOID = auto()
+    BOOL = auto()
+    CHAR = auto()
+    SHORT = auto()
+    INT = auto()
+    LONG = auto()
+    FLOAT = auto()
+    DOUBLE = auto()
+    SIGNED = auto()
+    UNSIGNED = auto()
+    SIZE_T = auto()
+    FLOAT2 = auto()
+    FLOAT3 = auto()
+    FLOAT4 = auto()
+    DOUBLE2 = auto()
+    DOUBLE3 = auto()
+    DOUBLE4 = auto()
+    INT2 = auto()
+    INT3 = auto()
+    INT4 = auto()
+    UINT2 = auto()
+    UINT3 = auto()
+    UINT4 = auto()
+    CHAR2 = auto()
+    CHAR3 = auto()
+    CHAR4 = auto()
+    UCHAR2 = auto()
+    UCHAR3 = auto()
+    UCHAR4 = auto()
+    SHORT2 = auto()
+    SHORT3 = auto()
+    SHORT4 = auto()
+    USHORT2 = auto()
+    USHORT3 = auto()
+    USHORT4 = auto()
+    LONG2 = auto()
+    LONG3 = auto()
+    LONG4 = auto()
+    ULONG2 = auto()
+    ULONG3 = auto()
+    ULONG4 = auto()
+    LONGLONG2 = auto()
+    LONGLONG3 = auto()
+    LONGLONG4 = auto()
+    ULONGLONG2 = auto()
+    ULONGLONG3 = auto()
+    ULONGLONG4 = auto()
+    TEXTURE = auto()
+    SURFACE = auto()
+    CUDAARRAY = auto()
+    CUDAARRAYT = auto()
+    TRUE = auto()
+    FALSE = auto()
+    NULL = auto()
+    NULLPTR = auto()
+    IDENTIFIER = auto()
+    NUMBER = auto()
+    STRING = auto()
+    CHAR_LIT = auto()
+    PREPROCESSOR = auto()
+    SCOPE = auto()
+    SHIFT_LEFT = auto()
+    SHIFT_RIGHT = auto()
+    PLUS_EQUALS = auto()
+    MINUS_EQUALS = auto()
+    MULTIPLY_EQUALS = auto()
+    DIVIDE_EQUALS = auto()
+    MODULO_EQUALS = auto()
+    AND_EQUALS = auto()
+    OR_EQUALS = auto()
+    XOR_EQUALS = auto()
+    SHIFT_LEFT_EQUALS = auto()
+    SHIFT_RIGHT_EQUALS = auto()
+    LOGICAL_AND = auto()
+    LOGICAL_OR = auto()
+    EQUAL = auto()
+    NOT_EQUAL = auto()
+    LESS_EQUAL = auto()
+    GREATER_EQUAL = auto()
+    INCREMENT = auto()
+    DECREMENT = auto()
+    ARROW = auto()
+    DOT = auto()
+    PLUS = auto()
+    MINUS = auto()
+    MULTIPLY = auto()
+    DIVIDE = auto()
+    MODULO = auto()
+    ASSIGN = auto()
+    LESS_THAN = auto()
+    GREATER_THAN = auto()
+    LOGICAL_NOT = auto()
+    BITWISE_AND = auto()
+    BITWISE_OR = auto()
+    BITWISE_XOR = auto()
+    BITWISE_NOT = auto()
+    QUESTION = auto()
+    COLON = auto()
+    LBRACE = auto()
+    RBRACE = auto()
+    LPAREN = auto()
+    RPAREN = auto()
+    LBRACKET = auto()
+    RBRACKET = auto()
+    SEMICOLON = auto()
+    COMMA = auto()
+    WHITESPACE = auto()
+
+
+class CudaLexer:
+    def __init__(self, code: str):
+        self._token_patterns = [(name, re.compile(pattern)) for name, pattern in TOKENS]
+        self.code = code
+        self._length = len(code)
+        self.reserved_keywords = KEYWORDS
+
+    def tokenize(self) -> List[Tuple[str, str]]:
+        """Tokenize the input code and return list of tokens"""
+        return list(self.token_generator())
+
+    def token_generator(self) -> Iterator[Tuple[str, str]]:
+        """Generator function that yields tokens one at a time"""
+        pos = 0
+        while pos < self._length:
+            token = self._next_token(pos)
+            if token is None:
+                raise SyntaxError(
+                    f"Illegal character '{self.code[pos]}' at position {pos}"
+                )
+            new_pos, token_type, text = token
+
+            # Check if identifier is a reserved keyword
+            if token_type == "IDENTIFIER" and text in self.reserved_keywords:
+                token_type = self.reserved_keywords[text]
+
+            # Skip whitespace and comments unless needed for debugging
+            if token_type not in SKIP_TOKENS:
+                yield (token_type, text)
+
+            pos = new_pos
+
+        yield ("EOF", "")
+
+    def _next_token(self, pos: int) -> Tuple[int, str, str]:
+        """Find the next token starting at the given position"""
+        for token_type, pattern in self._token_patterns:
+            match = pattern.match(self.code, pos)
+            if match:
+                return match.end(0), token_type, match.group(0)
+        return None
+
+    @classmethod
+    def from_file(cls, filepath: str) -> "CudaLexer":
+        """Create a lexer instance from a file"""
+        with open(filepath, "r") as f:
+            return cls(f.read())
+
+
+# Compatibility wrapper
+class Lexer:
+    """Compatibility wrapper around CudaLexer"""
+
+    def __init__(self, input_str):
+        self.lexer = CudaLexer(input_str)
+        self.tokens = self.lexer.tokenize()
+        self.current_pos = 0
+
+    def next(self):
+        if self.current_pos < len(self.tokens):
+            token = self.tokens[self.current_pos]
+            self.current_pos += 1
+            return token
+        return ("EOF", "")
+
+    def peek(self):
+        if self.current_pos < len(self.tokens):
+            return self.tokens[self.current_pos]
+        return ("EOF", "")
diff --git a/crosstl/backend/CUDA/CudaParser.py b/crosstl/backend/CUDA/CudaParser.py
new file mode 100644
index 0000000..51b8b5a
--- /dev/null
+++ b/crosstl/backend/CUDA/CudaParser.py
@@ -0,0 +1,876 @@
+"""CUDA Parser Implementation"""
+
+from .CudaAst import (
+    ArrayAccessNode,
+    AssignmentNode,
+    AtomicOperationNode,
+    BinaryOpNode,
+    BreakNode,
+    CaseNode,
+    CastNode,
+    ConstantMemoryNode,
+    ContinueNode,
+    CudaBuiltinNode,
+    DoWhileNode,
+    ForNode,
+    FunctionCallNode,
+    FunctionNode,
+    IfNode,
+    KernelLaunchNode,
+    KernelNode,
+    MemberAccessNode,
+    PreprocessorNode,
+    ReturnNode,
+    ShaderNode,
+    SharedMemoryNode,
+    StructNode,
+    SwitchNode,
+    SyncNode,
+    TernaryOpNode,
+    UnaryOpNode,
+    VariableNode,
+    WhileNode,
+)
+from .CudaLexer import CudaLexer
+
+
+class CudaParser:
+    def __init__(self, tokens):
+        self.tokens = tokens
+        self.current_index = 0
+        self.current_token = tokens[0] if tokens else None
+
+    def parse(self):
+        """Parse the entire CUDA program"""
+        includes = []
+        functions = []
+        structs = []
+        global_variables = []
+        kernels = []
+
+        while self.current_token[0] != "EOF":
+            if self.current_token[0] == "PREPROCESSOR":
+                includes.append(self.parse_preprocessor())
+            elif self.current_token[0] == "STRUCT":
+                structs.append(self.parse_struct())
+            elif (
+                self.current_token[0] in ["GLOBAL", "DEVICE", "HOST"]
+                or self.peek_function()
+            ):
+                func = self.parse_function()
+                if isinstance(func, KernelNode):
+                    kernels.append(func)
+                else:
+                    functions.append(func)
+            elif (
+                self.current_token[0] in ["CONSTANT", "SHARED"] or self.peek_variable()
+            ):
+                global_variables.append(self.parse_global_variable())
+            else:
+                # Skip unexpected tokens
+                self.eat(self.current_token[0])
+
+        return ShaderNode(includes, functions, structs, global_variables, kernels)
+
+    def peek_function(self):
+        """Check if the next tokens form a function declaration"""
+        # Look ahead for function pattern: [qualifiers] type name (
+        saved_index = self.current_index
+
+        # Skip qualifiers and type
+        while saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+            "CONST",
+            "STATIC",
+            "INLINE",
+            "EXTERN",
+        ]:
+            saved_index += 1
+
+        # Skip type (could be multiple tokens)
+        if saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+            "VOID",
+            "INT",
+            "FLOAT",
+            "DOUBLE",
+            "CHAR",
+            "SHORT",
+            "LONG",
+            "BOOL",
+            "IDENTIFIER",
+            "UNSIGNED",
+            "SIGNED",
+        ]:
+            saved_index += 1
+
+        # Check for function name followed by parenthesis
+        if (
+            saved_index < len(self.tokens) - 1
+            and self.tokens[saved_index][0] == "IDENTIFIER"
+            and self.tokens[saved_index + 1][0] == "LPAREN"
+        ):
+            return True
+
+        return False
+
+    def peek_variable(self):
+        """Check if the next tokens form a variable declaration"""
+        # Look for type followed by identifier and semicolon or assignment
+        saved_index = self.current_index
+
+        # Skip qualifiers
+        while saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+            "CONST",
+            "STATIC",
+            "EXTERN",
+        ]:
+            saved_index += 1
+
+        # Check for type
+        if saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+            "INT",
+            "FLOAT",
+            "DOUBLE",
+            "CHAR",
+            "SHORT",
+            "LONG",
+            "BOOL",
+            "IDENTIFIER",
+            "UNSIGNED",
+            "SIGNED",
+            "FLOAT2",
+            "FLOAT3",
+            "FLOAT4",
+        ]:
+            saved_index += 1
+            # Check for identifier
+            if (
+                saved_index < len(self.tokens)
+                and self.tokens[saved_index][0] == "IDENTIFIER"
+            ):
+                saved_index += 1
+                # Check for semicolon or assignment
+                if saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+                    "SEMICOLON",
+                    "ASSIGN",
+                ]:
+                    return True
+
+        return False
+
+    def eat(self, expected_type):
+        """Consume a token of the expected type"""
+        if self.current_token[0] == expected_type:
+            token = self.current_token
+            self.current_index += 1
+            if self.current_index < len(self.tokens):
+                self.current_token = self.tokens[self.current_index]
+            else:
+                self.current_token = ("EOF", "")
+            return token
+        else:
+            raise SyntaxError(f"Expected {expected_type}, got {self.current_token[0]}")
+
+    def parse_preprocessor(self):
+        """Parse preprocessor directives"""
+        directive_token = self.eat("PREPROCESSOR")
+        directive_text = directive_token[1].strip()
+
+        if directive_text.startswith("#include"):
+            content = directive_text[8:].strip()
+            return PreprocessorNode("include", content)
+        elif directive_text.startswith("#define"):
+            content = directive_text[7:].strip()
+            return PreprocessorNode("define", content)
+        else:
+            return PreprocessorNode("other", directive_text)
+
+    def parse_struct(self):
+        """Parse struct declaration"""
+        self.eat("STRUCT")
+        name = self.eat("IDENTIFIER")[1]
+        self.eat("LBRACE")
+
+        members = []
+        while self.current_token[0] != "RBRACE":
+            member = self.parse_variable_declaration()
+            members.append(member)
+            self.eat("SEMICOLON")
+
+        self.eat("RBRACE")
+        self.eat("SEMICOLON")
+
+        return StructNode(name, members)
+
+    def parse_function(self):
+        """Parse function declaration including kernels"""
+        qualifiers = []
+
+        # Parse CUDA qualifiers
+        while self.current_token[0] in ["GLOBAL", "DEVICE", "HOST", "INLINE", "STATIC"]:
+            qualifiers.append(self.current_token[1])
+            self.eat(self.current_token[0])
+
+        # Parse return type
+        return_type = self.parse_type()
+
+        # Parse function name
+        name = self.eat("IDENTIFIER")[1]
+
+        # Parse parameters
+        params = self.parse_parameters()
+
+        # Parse function body
+        body = self.parse_block()
+
+        # Check if it's a kernel function
+        if "__global__" in qualifiers:
+            return KernelNode(return_type, name, params, body)
+        else:
+            return FunctionNode(return_type, name, params, body, qualifiers)
+
+    def parse_parameters(self):
+        """Parse function parameters"""
+        self.eat("LPAREN")
+        params = []
+
+        if self.current_token[0] != "RPAREN":
+            params.append(self.parse_parameter())
+
+            while self.current_token[0] == "COMMA":
+                self.eat("COMMA")
+                params.append(self.parse_parameter())
+
+        self.eat("RPAREN")
+        return params
+
+    def parse_parameter(self):
+        """Parse a single parameter"""
+        param_type = self.parse_type()
+        param_name = self.eat("IDENTIFIER")[1]
+        return VariableNode(param_type, param_name)
+
+    def parse_type(self):
+        """Parse type specification"""
+        type_parts = []
+
+        # Handle qualifiers
+        while self.current_token[0] in ["CONST", "VOLATILE", "UNSIGNED", "SIGNED"]:
+            type_parts.append(self.current_token[1])
+            self.eat(self.current_token[0])
+
+        # Handle basic types
+        if self.current_token[0] in [
+            "VOID",
+            "CHAR",
+            "SHORT",
+            "INT",
+            "LONG",
+            "FLOAT",
+            "DOUBLE",
+            "BOOL",
+            "FLOAT2",
+            "FLOAT3",
+            "FLOAT4",
+            "INT2",
+            "INT3",
+            "INT4",
+            "DOUBLE2",
+            "DOUBLE3",
+            "DOUBLE4",
+            "SIZE_T",
+        ]:
+            type_parts.append(self.current_token[1])
+            self.eat(self.current_token[0])
+        elif self.current_token[0] == "IDENTIFIER":
+            type_parts.append(self.current_token[1])
+            self.eat("IDENTIFIER")
+
+        # Handle pointer/reference
+        while self.current_token[0] == "MULTIPLY":
+            type_parts.append("*")
+            self.eat("MULTIPLY")
+
+        # Handle array brackets
+        if self.current_token[0] == "LBRACKET":
+            self.eat("LBRACKET")
+            if self.current_token[0] == "NUMBER":
+                type_parts.append(f"[{self.current_token[1]}]")
+                self.eat("NUMBER")
+            else:
+                type_parts.append("[]")
+            self.eat("RBRACKET")
+
+        return " ".join(type_parts)
+
+    def parse_global_variable(self):
+        """Parse global variable declaration"""
+        qualifiers = []
+
+        # Parse CUDA memory qualifiers
+        while self.current_token[0] in ["CONSTANT", "SHARED", "DEVICE", "MANAGED"]:
+            qualifiers.append(self.current_token[1])
+            self.eat(self.current_token[0])
+
+        var = self.parse_variable_declaration()
+        var.qualifiers = qualifiers
+        self.eat("SEMICOLON")
+
+        # Return specialized nodes for CUDA memory types
+        if "__constant__" in qualifiers:
+            return ConstantMemoryNode(var.vtype, var.name, var.value)
+        elif "__shared__" in qualifiers:
+            return SharedMemoryNode(var.vtype, var.name)
+        else:
+            return var
+
+    def parse_variable_declaration(self):
+        """Parse variable declaration"""
+        qualifiers = []
+
+        # Parse CUDA memory qualifiers
+        while self.current_token[0] in ["SHARED", "CONSTANT"]:
+            qualifiers.append(self.current_token[1])
+            self.eat(self.current_token[0])
+
+        vtype = self.parse_type()
+        name = self.eat("IDENTIFIER")[1]
+
+        # Handle array declarations
+        if self.current_token[0] == "LBRACKET":
+            self.eat("LBRACKET")
+            if self.current_token[0] == "NUMBER":
+                size = self.current_token[1]
+                vtype += f"[{size}]"
+                self.eat("NUMBER")
+            else:
+                vtype += "[]"
+            self.eat("RBRACKET")
+
+        value = None
+        if self.current_token[0] == "ASSIGN":
+            self.eat("ASSIGN")
+            value = self.parse_expression()
+
+        var = VariableNode(vtype, name, value, qualifiers)
+
+        # Return specialized nodes for CUDA memory types
+        if "__shared__" in qualifiers:
+            return SharedMemoryNode(vtype, name)
+        elif "__constant__" in qualifiers:
+            return ConstantMemoryNode(vtype, name, value)
+        else:
+            return var
+
+    def parse_block(self):
+        """Parse a block of statements"""
+        self.eat("LBRACE")
+        statements = []
+
+        while self.current_token[0] != "RBRACE" and self.current_token[0] != "EOF":
+            stmt = self.parse_statement()
+            if stmt:
+                statements.append(stmt)
+
+        self.eat("RBRACE")
+        return statements
+
+    def parse_statement(self):
+        """Parse a single statement"""
+        if self.current_token[0] == "IF":
+            return self.parse_if_statement()
+        elif self.current_token[0] == "FOR":
+            return self.parse_for_statement()
+        elif self.current_token[0] == "WHILE":
+            return self.parse_while_statement()
+        elif self.current_token[0] == "DO":
+            return self.parse_do_while_statement()
+        elif self.current_token[0] == "SWITCH":
+            return self.parse_switch_statement()
+        elif self.current_token[0] == "RETURN":
+            return self.parse_return_statement()
+        elif self.current_token[0] == "BREAK":
+            self.eat("BREAK")
+            self.eat("SEMICOLON")
+            return BreakNode()
+        elif self.current_token[0] == "CONTINUE":
+            self.eat("CONTINUE")
+            self.eat("SEMICOLON")
+            return ContinueNode()
+        elif self.current_token[0] in ["SYNCTHREADS", "SYNCWARP"]:
+            return self.parse_sync_statement()
+        elif self.current_token[0] == "LBRACE":
+            return self.parse_block()
+        elif self.is_variable_declaration():
+            var = self.parse_variable_declaration()
+            self.eat("SEMICOLON")
+            return var
+        else:
+            # Expression statement or assignment
+            expr = self.parse_assignment_expression()
+            self.eat("SEMICOLON")
+            return expr
+
+    def is_variable_declaration(self):
+        """Check if current position is a variable declaration"""
+        # Simple heuristic: type followed by identifier
+        if self.current_token[0] in [
+            "INT",
+            "FLOAT",
+            "DOUBLE",
+            "CHAR",
+            "BOOL",
+            "VOID",
+            "FLOAT2",
+            "FLOAT3",
+            "FLOAT4",
+            "IDENTIFIER",
+            "SHARED",
+            "CONSTANT",
+        ]:
+            # Look ahead for identifier
+            saved_index = self.current_index
+
+            # Skip qualifiers
+            while saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+                "SHARED",
+                "CONSTANT",
+            ]:
+                saved_index += 1
+
+            # Check for type
+            if saved_index < len(self.tokens) and self.tokens[saved_index][0] in [
+                "INT",
+                "FLOAT",
+                "DOUBLE",
+                "CHAR",
+                "BOOL",
+                "VOID",
+                "FLOAT2",
+                "FLOAT3",
+                "FLOAT4",
+                "IDENTIFIER",
+            ]:
+                saved_index += 1
+
+                # Check for identifier
+                if (
+                    saved_index < len(self.tokens)
+                    and self.tokens[saved_index][0] == "IDENTIFIER"
+                ):
+                    return True
+        return False
+
+    def parse_if_statement(self):
+        """Parse if statement"""
+        self.eat("IF")
+        self.eat("LPAREN")
+        condition = self.parse_expression()
+        self.eat("RPAREN")
+
+        if_body = self.parse_statement()
+
+        else_body = None
+        if self.current_token[0] == "ELSE":
+            self.eat("ELSE")
+            else_body = self.parse_statement()
+
+        return IfNode(condition, if_body, else_body)
+
+    def parse_for_statement(self):
+        """Parse for loop"""
+        self.eat("FOR")
+        self.eat("LPAREN")
+
+        # Parse initialization
+        init = None
+        if self.current_token[0] != "SEMICOLON":
+            if self.is_variable_declaration():
+                init = self.parse_variable_declaration()
+            else:
+                init = self.parse_expression()
+        self.eat("SEMICOLON")
+
+        # Parse condition
+        condition = None
+        if self.current_token[0] != "SEMICOLON":
+            condition = self.parse_expression()
+        self.eat("SEMICOLON")
+
+        # Parse update
+        update = None
+        if self.current_token[0] != "RPAREN":
+            update = self.parse_expression()
+        self.eat("RPAREN")
+
+        # Parse body
+        body = self.parse_statement()
+
+        return ForNode(init, condition, update, body)
+
+    def parse_while_statement(self):
+        """Parse while loop"""
+        self.eat("WHILE")
+        self.eat("LPAREN")
+        condition = self.parse_expression()
+        self.eat("RPAREN")
+        body = self.parse_statement()
+
+        return WhileNode(condition, body)
+
+    def parse_do_while_statement(self):
+        """Parse do-while loop"""
+        self.eat("DO")
+        body = self.parse_statement()
+        self.eat("WHILE")
+        self.eat("LPAREN")
+        condition = self.parse_expression()
+        self.eat("RPAREN")
+        self.eat("SEMICOLON")
+
+        return DoWhileNode(body, condition)
+
+    def parse_switch_statement(self):
+        """Parse switch statement"""
+        self.eat("SWITCH")
+        self.eat("LPAREN")
+        expression = self.parse_expression()
+        self.eat("RPAREN")
+        self.eat("LBRACE")
+
+        cases = []
+        default_case = None
+
+        while self.current_token[0] != "RBRACE":
+            if self.current_token[0] == "CASE":
+                self.eat("CASE")
+                value = self.parse_expression()
+                self.eat("COLON")
+                body = []
+                while self.current_token[0] not in ["CASE", "DEFAULT", "RBRACE"]:
+                    body.append(self.parse_statement())
+                cases.append(CaseNode(value, body))
+            elif self.current_token[0] == "DEFAULT":
+                self.eat("DEFAULT")
+                self.eat("COLON")
+                default_case = []
+                while self.current_token[0] not in ["CASE", "RBRACE"]:
+                    default_case.append(self.parse_statement())
+            else:
+                self.eat(self.current_token[0])  # Skip unexpected tokens
+
+        self.eat("RBRACE")
+        return SwitchNode(expression, cases, default_case)
+
+    def parse_return_statement(self):
+        """Parse return statement"""
+        self.eat("RETURN")
+
+        value = None
+        if self.current_token[0] != "SEMICOLON":
+            value = self.parse_expression()
+
+        self.eat("SEMICOLON")
+        return ReturnNode(value)
+
+    def parse_sync_statement(self):
+        """Parse CUDA synchronization statements"""
+        sync_type = self.current_token[1]
+        self.eat(self.current_token[0])
+
+        self.eat("LPAREN")
+        args = []
+        if self.current_token[0] != "RPAREN":
+            args.append(self.parse_expression())
+            while self.current_token[0] == "COMMA":
+                self.eat("COMMA")
+                args.append(self.parse_expression())
+        self.eat("RPAREN")
+        self.eat("SEMICOLON")
+
+        return SyncNode(sync_type, args)
+
+    def parse_expression(self):
+        """Parse expression with precedence"""
+        return self.parse_ternary_expression()
+
+    def parse_ternary_expression(self):
+        """Parse ternary conditional operator"""
+        expr = self.parse_logical_or_expression()
+
+        if self.current_token[0] == "QUESTION":
+            self.eat("QUESTION")
+            true_expr = self.parse_expression()
+            self.eat("COLON")
+            false_expr = self.parse_expression()
+            return TernaryOpNode(expr, true_expr, false_expr)
+
+        return expr
+
+    def parse_logical_or_expression(self):
+        """Parse logical OR expression"""
+        left = self.parse_logical_and_expression()
+
+        while self.current_token[0] == "LOGICAL_OR":
+            op = self.current_token[1]
+            self.eat("LOGICAL_OR")
+            right = self.parse_logical_and_expression()
+            left = BinaryOpNode(left, op, right)
+
+        return left
+
+    def parse_logical_and_expression(self):
+        """Parse logical AND expression"""
+        left = self.parse_equality_expression()
+
+        while self.current_token[0] == "LOGICAL_AND":
+            op = self.current_token[1]
+            self.eat("LOGICAL_AND")
+            right = self.parse_equality_expression()
+            left = BinaryOpNode(left, op, right)
+
+        return left
+
+    def parse_equality_expression(self):
+        """Parse equality expression"""
+        left = self.parse_relational_expression()
+
+        while self.current_token[0] in ["EQUAL", "NOT_EQUAL"]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            right = self.parse_relational_expression()
+            left = BinaryOpNode(left, op, right)
+
+        return left
+
+    def parse_relational_expression(self):
+        """Parse relational expression"""
+        left = self.parse_additive_expression()
+
+        while self.current_token[0] in [
+            "LESS_THAN",
+            "GREATER_THAN",
+            "LESS_EQUAL",
+            "GREATER_EQUAL",
+        ]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            right = self.parse_additive_expression()
+            left = BinaryOpNode(left, op, right)
+
+        return left
+
+    def parse_additive_expression(self):
+        """Parse additive expression"""
+        left = self.parse_multiplicative_expression()
+
+        while self.current_token[0] in ["PLUS", "MINUS"]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            right = self.parse_multiplicative_expression()
+            left = BinaryOpNode(left, op, right)
+
+        return left
+
+    def parse_multiplicative_expression(self):
+        """Parse multiplicative expression"""
+        left = self.parse_unary_expression()
+
+        while self.current_token[0] in ["MULTIPLY", "DIVIDE", "MODULO"]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            right = self.parse_unary_expression()
+            left = BinaryOpNode(left, op, right)
+
+        return left
+
+    def parse_unary_expression(self):
+        """Parse unary expression"""
+        if self.current_token[0] in ["PLUS", "MINUS", "LOGICAL_NOT", "BITWISE_NOT"]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            operand = self.parse_unary_expression()
+            return UnaryOpNode(op, operand)
+        elif self.current_token[0] in ["INCREMENT", "DECREMENT"]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            operand = self.parse_postfix_expression()
+            return UnaryOpNode(op, operand)
+        else:
+            return self.parse_postfix_expression()
+
+    def parse_postfix_expression(self):
+        """Parse postfix expression"""
+        left = self.parse_primary_expression()
+
+        while True:
+            if self.current_token[0] == "DOT":
+                self.eat("DOT")
+                member = self.eat("IDENTIFIER")[1]
+                left = MemberAccessNode(left, member, False)
+            elif self.current_token[0] == "ARROW":
+                self.eat("ARROW")
+                member = self.eat("IDENTIFIER")[1]
+                left = MemberAccessNode(left, member, True)
+            elif self.current_token[0] == "LBRACKET":
+                self.eat("LBRACKET")
+                index = self.parse_expression()
+                self.eat("RBRACKET")
+                left = ArrayAccessNode(left, index)
+            elif self.current_token[0] == "LPAREN":
+                # Function call
+                self.eat("LPAREN")
+                args = []
+                if self.current_token[0] != "RPAREN":
+                    args.append(self.parse_expression())
+                    while self.current_token[0] == "COMMA":
+                        self.eat("COMMA")
+                        args.append(self.parse_expression())
+                self.eat("RPAREN")
+
+                # Check for atomic operations
+                if isinstance(left, str) and left.startswith("atomic"):
+                    left = AtomicOperationNode(left, args)
+                else:
+                    left = FunctionCallNode(left, args)
+            elif self.current_token[0] == "KERNEL_LAUNCH_START":
+                # Kernel launch: kernel<<<blocks, threads>>>(args)
+                return self.parse_kernel_launch(left)
+            elif self.current_token[0] in ["INCREMENT", "DECREMENT"]:
+                op = self.current_token[1]
+                self.eat(self.current_token[0])
+                left = UnaryOpNode(f"post{op}", left)
+            else:
+                break
+
+        return left
+
+    def parse_kernel_launch(self, kernel_name):
+        """Parse CUDA kernel launch syntax"""
+        self.eat("KERNEL_LAUNCH_START")
+
+        # Parse blocks
+        blocks = self.parse_expression()
+        self.eat("COMMA")
+
+        # Parse threads
+        threads = self.parse_expression()
+
+        # Parse optional shared memory and stream
+        shared_mem = None
+        stream = None
+
+        if self.current_token[0] == "COMMA":
+            self.eat("COMMA")
+            shared_mem = self.parse_expression()
+
+            if self.current_token[0] == "COMMA":
+                self.eat("COMMA")
+                stream = self.parse_expression()
+
+        self.eat("KERNEL_LAUNCH_END")
+
+        # Parse arguments
+        self.eat("LPAREN")
+        args = []
+        if self.current_token[0] != "RPAREN":
+            args.append(self.parse_expression())
+            while self.current_token[0] == "COMMA":
+                self.eat("COMMA")
+                args.append(self.parse_expression())
+        self.eat("RPAREN")
+
+        return KernelLaunchNode(kernel_name, blocks, threads, shared_mem, stream, args)
+
+    def parse_primary_expression(self):
+        """Parse primary expression"""
+        if self.current_token[0] == "NUMBER":
+            value = self.current_token[1]
+            self.eat("NUMBER")
+            return value
+        elif self.current_token[0] == "STRING":
+            value = self.current_token[1]
+            self.eat("STRING")
+            return value
+        elif self.current_token[0] == "CHAR_LIT":
+            value = self.current_token[1]
+            self.eat("CHAR_LIT")
+            return value
+        elif self.current_token[0] in ["TRUE", "FALSE"]:
+            value = self.current_token[1]
+            self.eat(self.current_token[0])
+            return value
+        elif self.current_token[0] in ["NULL", "NULLPTR"]:
+            value = self.current_token[1]
+            self.eat(self.current_token[0])
+            return value
+        elif self.current_token[0] == "IDENTIFIER":
+            name = self.current_token[1]
+            self.eat("IDENTIFIER")
+            return name
+        elif self.current_token[0] in [
+            "ATOMICADD",
+            "ATOMICSUB",
+            "ATOMICMAX",
+            "ATOMICMIN",
+            "ATOMICEXCH",
+            "ATOMICCAS",
+        ]:
+            name = self.current_token[1]
+            self.eat(self.current_token[0])
+            return name
+        elif self.current_token[0] in [
+            "THREADIDX",
+            "BLOCKIDX",
+            "GRIDDIM",
+            "BLOCKDIM",
+            "WARPSIZE",
+        ]:
+            builtin_name = self.current_token[1]
+            self.eat(self.current_token[0])
+
+            # Check for component access (.x, .y, .z)
+            if self.current_token[0] == "DOT":
+                self.eat("DOT")
+                component = self.eat("IDENTIFIER")[1]
+                return CudaBuiltinNode(builtin_name, component)
+            else:
+                return CudaBuiltinNode(builtin_name)
+        elif self.current_token[0] == "LPAREN":
+            self.eat("LPAREN")
+            if self.current_token[0] in ["INT", "FLOAT", "DOUBLE", "CHAR"]:
+                # Type cast
+                target_type = self.parse_type()
+                self.eat("RPAREN")
+                expr = self.parse_unary_expression()
+                return CastNode(target_type, expr)
+            else:
+                # Parenthesized expression
+                expr = self.parse_expression()
+                self.eat("RPAREN")
+                return expr
+        else:
+            raise SyntaxError(
+                f"Unexpected token in primary expression: {self.current_token}"
+            )
+
+    def parse_assignment_expression(self):
+        """Parse assignment expression"""
+        left = self.parse_expression()
+
+        if self.current_token[0] in [
+            "ASSIGN",
+            "PLUS_EQUALS",
+            "MINUS_EQUALS",
+            "MULTIPLY_EQUALS",
+            "DIVIDE_EQUALS",
+            "MODULO_EQUALS",
+            "AND_EQUALS",
+            "OR_EQUALS",
+            "XOR_EQUALS",
+        ]:
+            op = self.current_token[1]
+            self.eat(self.current_token[0])
+            right = self.parse_expression()
+            return AssignmentNode(left, right, op)
+
+        return left
diff --git a/crosstl/backend/CUDA/__init__.py b/crosstl/backend/CUDA/__init__.py
new file mode 100644
index 0000000..be2a910
--- /dev/null
+++ b/crosstl/backend/CUDA/__init__.py
@@ -0,0 +1,3 @@
+from .CudaLexer import CudaLexer
+from .CudaParser import CudaParser
+from .CudaCrossGLCodeGen import CudaToCrossGLConverter
diff --git a/crosstl/formatter.py b/crosstl/formatter.py
index dafb233..246d2c4 100644
--- a/crosstl/formatter.py
+++ b/crosstl/formatter.py
@@ -18,6 +18,7 @@ class ShaderLanguage(Enum):
     SPIRV = "spirv"
     MOJO = "mojo"
     RUST = "rust"
+    CUDA = "cuda"
     UNKNOWN = "unknown"
 
 
@@ -64,6 +65,8 @@ class CodeFormatter:
             return ShaderLanguage.SPIRV
         elif ext in [".rs", ".rust"]:
             return ShaderLanguage.RUST
+        elif ext in [".cu", ".cuh", ".cuda"]:
+            return ShaderLanguage.CUDA
         else:
             return ShaderLanguage.UNKNOWN
 
@@ -95,6 +98,7 @@ class CodeFormatter:
             ShaderLanguage.GLSL,
             ShaderLanguage.METAL,
             ShaderLanguage.RUST,
+            ShaderLanguage.CUDA,
         ]:
             return self._format_with_clang(code, language)
         elif language == ShaderLanguage.SPIRV:
@@ -115,6 +119,7 @@ class CodeFormatter:
             ShaderLanguage.GLSL: "Google",
             ShaderLanguage.METAL: "LLVM",
             ShaderLanguage.RUST: "LLVM",
+            ShaderLanguage.CUDA: "Google",
         }
         style = style_map.get(language, "LLVM")
 
@@ -125,6 +130,7 @@ class CodeFormatter:
                 ShaderLanguage.GLSL: ".glsl",
                 ShaderLanguage.METAL: ".metal",
                 ShaderLanguage.RUST: ".rs",
+                ShaderLanguage.CUDA: ".cu",
             }
             ext = ext_map.get(language, ".txt")
 
@@ -323,6 +329,7 @@ def format_shader_code(code, backend, output_path=None):
         "vulkan": ShaderLanguage.SPIRV,
         "mojo": ShaderLanguage.MOJO,
         "rust": ShaderLanguage.RUST,
+        "cuda": ShaderLanguage.CUDA,
     }
 
     language = language_map.get(backend.lower())
diff --git a/crosstl/translator/codegen/__init__.py b/crosstl/translator/codegen/__init__.py
index e69de29..10048f4 100644
--- a/crosstl/translator/codegen/__init__.py
+++ b/crosstl/translator/codegen/__init__.py
@@ -0,0 +1,7 @@
+from . import GLSL_codegen
+from . import directx_codegen
+from . import metal_codegen
+from . import SPIRV_codegen
+from . import mojo_codegen
+from . import rust_codegen
+from . import cuda_codegen
diff --git a/crosstl/translator/codegen/cuda_codegen.py b/crosstl/translator/codegen/cuda_codegen.py
new file mode 100644
index 0000000..0068254
--- /dev/null
+++ b/crosstl/translator/codegen/cuda_codegen.py
@@ -0,0 +1,427 @@
+"""CUDA Code Generator"""
+
+from ..ast import (
+    AssignmentNode,
+    BinaryOpNode,
+    ForNode,
+    FunctionCallNode,
+    IfNode,
+    MemberAccessNode,
+    ReturnNode,
+    StructNode,
+    TernaryOpNode,
+    UnaryOpNode,
+    VariableNode,
+    ArrayAccessNode,
+    ArrayNode,
+    ShaderNode,
+    FunctionNode,
+)
+
+
+class CudaCodeGen:
+    """Generates CUDA code from CrossGL AST"""
+
+    def __init__(self):
+        self.indent_level = 0
+        self.output = []
+
+    def generate(self, ast_node):
+        """Generate CUDA code from CrossGL AST"""
+        self.output = []
+        self.indent_level = 0
+        self.visit(ast_node)
+        return "\n".join(self.output)
+
+    def visit(self, node):
+        """Visit an AST node and generate code"""
+        method_name = f"visit_{type(node).__name__}"
+        visitor = getattr(self, method_name, self.generic_visit)
+        return visitor(node)
+
+    def generic_visit(self, node):
+        """Generic visitor for unknown nodes"""
+        if isinstance(node, str):
+            return node
+        elif isinstance(node, list):
+            return [self.visit(item) for item in node]
+        else:
+            return str(node)
+
+    def emit(self, code):
+        """Emit code with proper indentation"""
+        if code.strip():
+            self.output.append("    " * self.indent_level + code)
+        else:
+            self.output.append("")
+
+    def visit_ShaderNode(self, node):
+        """Visit the main shader node"""
+        # Add CUDA headers
+        self.emit("#include <cuda_runtime.h>")
+        self.emit("#include <device_launch_parameters.h>")
+        self.emit("")
+
+        # Process structs
+        if hasattr(node, "structs") and node.structs:
+            for struct in node.structs:
+                self.visit(struct)
+                self.emit("")
+
+        # Process cbuffers as constant memory
+        if hasattr(node, "cbuffers") and node.cbuffers:
+            for cbuffer in node.cbuffers:
+                self.visit_cbuffer(cbuffer)
+                self.emit("")
+
+        # Process global variables
+        if hasattr(node, "global_variables") and node.global_variables:
+            for var in node.global_variables:
+                self.visit(var)
+                self.emit("")
+
+        # Process functions
+        if hasattr(node, "functions") and node.functions:
+            for func in node.functions:
+                self.visit(func)
+                self.emit("")
+
+    def visit_StructNode(self, node):
+        """Visit struct declaration"""
+        self.emit(f"struct {node.name} {{")
+        self.indent_level += 1
+
+        for member in node.members:
+            member_type = self.convert_crossgl_type_to_cuda(member.vtype)
+            self.emit(f"{member_type} {member.name};")
+
+        self.indent_level -= 1
+        self.emit("};")
+
+    def visit_FunctionNode(self, node):
+        """Visit function declaration"""
+        # Determine function qualifiers based on CrossGL qualifier
+        qualifiers = []
+
+        # Check CrossGL function qualifier
+        if hasattr(node, "qualifier") and node.qualifier:
+            if node.qualifier == "compute":
+                qualifiers.append("__global__")
+            elif node.qualifier == "vertex" or node.qualifier == "fragment":
+                # Vertex and fragment shaders become device functions in CUDA
+                qualifiers.append("__device__")
+            else:
+                qualifiers.append("__device__")
+        else:
+            # Default to device function
+            qualifiers.append("__device__")
+
+        return_type = self.convert_crossgl_type_to_cuda(node.return_type)
+        qualifier_str = " ".join(qualifiers)
+
+        # Generate parameters
+        params = []
+        for param in node.params:
+            param_type = self.convert_crossgl_type_to_cuda(param.vtype)
+            params.append(f"{param_type} {param.name}")
+
+        param_str = ", ".join(params)
+        self.emit(f"{qualifier_str} {return_type} {node.name}({param_str}) {{")
+
+        self.indent_level += 1
+
+        # Add built-in variable mappings for kernels
+        if "__global__" in qualifiers:
+            self.emit("// CUDA built-in variables")
+            self.emit("int3 threadIdx = {threadIdx.x, threadIdx.y, threadIdx.z};")
+            self.emit("int3 blockIdx = {blockIdx.x, blockIdx.y, blockIdx.z};")
+            self.emit("int3 blockDim = {blockDim.x, blockDim.y, blockDim.z};")
+            self.emit("int3 gridDim = {gridDim.x, gridDim.y, gridDim.z};")
+            self.emit("")
+
+        # Process function body
+        for stmt in node.body:
+            self.visit(stmt)
+
+        self.indent_level -= 1
+        self.emit("}")
+
+    def visit_VariableNode(self, node):
+        """Visit variable declaration"""
+        # Handle both declaration and usage cases
+        if hasattr(node, "vtype") and node.vtype:
+            var_type = self.convert_crossgl_type_to_cuda(node.vtype)
+
+            # Check for special memory qualifiers
+            qualifiers = []
+            if hasattr(node, "qualifiers"):
+                if "workgroup" in str(node.qualifiers):
+                    qualifiers.append("__shared__")
+                elif "uniform" in str(node.qualifiers):
+                    qualifiers.append("__constant__")
+
+            qualifier_str = " ".join(qualifiers)
+            if qualifier_str:
+                qualifier_str += " "
+
+            # CrossGL VariableNode doesn't have a value attribute by default
+            # Variables are typically declared and then assigned separately
+            self.emit(f"{qualifier_str}{var_type} {node.name};")
+        else:
+            # This might be a variable reference, just return the name
+            return node.name
+
+    def visit_AssignmentNode(self, node):
+        """Visit assignment statement"""
+        left = self.visit(node.left)
+        right = self.visit(node.right)
+        operator = getattr(node, "operator", "=")
+        self.emit(f"{left} {operator} {right};")
+
+    def visit_BinaryOpNode(self, node):
+        """Visit binary operation"""
+        left = self.visit(node.left)
+        right = self.visit(node.right)
+        return f"({left} {node.op} {right})"
+
+    def visit_UnaryOpNode(self, node):
+        """Visit unary operation"""
+        operand = self.visit(node.operand)
+        return f"({node.op}{operand})"
+
+    def visit_FunctionCallNode(self, node):
+        """Visit function call"""
+        args = [self.visit(arg) for arg in node.args]
+        args_str = ", ".join(args)
+
+        # Convert built-in functions
+        func_name = self.convert_builtin_function(node.name)
+        return f"{func_name}({args_str})"
+
+    def visit_MemberAccessNode(self, node):
+        """Visit member access"""
+        obj = self.visit(node.object)
+        return f"{obj}.{node.member}"
+
+    def visit_ArrayAccessNode(self, node):
+        """Visit array access"""
+        array = self.visit(node.array)
+        index = self.visit(node.index)
+        return f"{array}[{index}]"
+
+    def visit_IfNode(self, node):
+        """Visit if statement"""
+        condition = self.visit(node.if_condition)
+        self.emit(f"if ({condition}) {{")
+
+        self.indent_level += 1
+        for stmt in node.if_body:
+            self.visit(stmt)
+        self.indent_level -= 1
+
+        # Handle else-if conditions
+        if hasattr(node, "else_if_conditions") and node.else_if_conditions:
+            for else_if_condition, else_if_body in zip(
+                node.else_if_conditions, node.else_if_bodies
+            ):
+                condition_str = self.visit(else_if_condition)
+                self.emit(f"}} else if ({condition_str}) {{")
+                self.indent_level += 1
+                for stmt in else_if_body:
+                    self.visit(stmt)
+                self.indent_level -= 1
+
+        if node.else_body:
+            self.emit("} else {")
+            self.indent_level += 1
+            for stmt in node.else_body:
+                self.visit(stmt)
+            self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_ForNode(self, node):
+        """Visit for loop"""
+        init = self.visit(node.init) if node.init else ""
+        condition = self.visit(node.condition) if node.condition else ""
+        update = self.visit(node.update) if node.update else ""
+
+        self.emit(f"for ({init}; {condition}; {update}) {{")
+
+        self.indent_level += 1
+        for stmt in node.body:
+            self.visit(stmt)
+        self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_WhileNode(self, node):
+        """Visit while loop"""
+        condition = self.visit(node.condition)
+        self.emit(f"while ({condition}) {{")
+
+        self.indent_level += 1
+        for stmt in node.body:
+            self.visit(stmt)
+        self.indent_level -= 1
+
+        self.emit("}")
+
+    def visit_ReturnNode(self, node):
+        """Visit return statement"""
+        if node.value:
+            value = self.visit(node.value)
+            self.emit(f"return {value};")
+        else:
+            self.emit("return;")
+
+    def convert_crossgl_type_to_cuda(self, crossgl_type):
+        """Convert CrossGL types to CUDA equivalents"""
+        type_mapping = {
+            # Basic types
+            "void": "void",
+            "bool": "bool",
+            "i8": "char",
+            "u8": "unsigned char",
+            "i16": "short",
+            "u16": "unsigned short",
+            "i32": "int",
+            "u32": "unsigned int",
+            "i64": "long long",
+            "u64": "unsigned long long",
+            "f32": "float",
+            "f64": "double",
+            "int": "int",
+            "float": "float",
+            "double": "double",
+            # Vector types (with generics)
+            "vec2<f32>": "float2",
+            "vec3<f32>": "float3",
+            "vec4<f32>": "float4",
+            "vec2<f64>": "double2",
+            "vec3<f64>": "double3",
+            "vec4<f64>": "double4",
+            "vec2<i32>": "int2",
+            "vec3<i32>": "int3",
+            "vec4<i32>": "int4",
+            "vec2<u32>": "uint2",
+            "vec3<u32>": "uint3",
+            "vec4<u32>": "uint4",
+            # Vector types (without generics - for compatibility)
+            "vec2": "float2",
+            "vec3": "float3",
+            "vec4": "float4",
+            "ivec2": "int2",
+            "ivec3": "int3",
+            "ivec4": "int4",
+            "uvec2": "uint2",
+            "uvec3": "uint3",
+            "uvec4": "uint4",
+            # Matrix types
+            "mat2": "float2x2",
+            "mat3": "float3x3",
+            "mat4": "float4x4",
+        }
+
+        # Handle arrays
+        if crossgl_type.startswith("array<") and crossgl_type.endswith(">"):
+            # Extract element type and size
+            inner = crossgl_type[6:-1]  # Remove "array<" and ">"
+            if "," in inner:
+                parts = inner.split(",")
+                element_type = parts[0].strip()
+                size = parts[1].strip()
+                cuda_element_type = type_mapping.get(element_type, element_type)
+                return f"{cuda_element_type}[{size}]"
+            else:
+                cuda_element_type = type_mapping.get(inner, inner)
+                return f"{cuda_element_type}*"
+
+        # Handle pointers
+        if crossgl_type.startswith("ptr<") and crossgl_type.endswith(">"):
+            element_type = crossgl_type[4:-1]  # Remove "ptr<" and ">"
+            cuda_element_type = type_mapping.get(element_type, element_type)
+            return f"{cuda_element_type}*"
+
+        return type_mapping.get(crossgl_type, crossgl_type)
+
+    def convert_builtin_function(self, func_name):
+        """Convert CrossGL built-in functions to CUDA equivalents"""
+        function_mapping = {
+            # Math functions
+            "sqrt": "sqrtf",
+            "pow": "powf",
+            "sin": "sinf",
+            "cos": "cosf",
+            "tan": "tanf",
+            "log": "logf",
+            "exp": "expf",
+            "abs": "fabsf",
+            "min": "fminf",
+            "max": "fmaxf",
+            "floor": "floorf",
+            "ceil": "ceilf",
+            # Vector constructors
+            "vec2<f32>": "make_float2",
+            "vec3<f32>": "make_float3",
+            "vec4<f32>": "make_float4",
+            "vec2<i32>": "make_int2",
+            "vec3<i32>": "make_int3",
+            "vec4<i32>": "make_int4",
+            # Atomic operations
+            "atomicAdd": "atomicAdd",
+            "atomicSub": "atomicSub",
+            "atomicMax": "atomicMax",
+            "atomicMin": "atomicMin",
+            "atomicExchange": "atomicExch",
+            "atomicCompareExchange": "atomicCAS",
+            # Synchronization
+            "workgroupBarrier": "__syncthreads",
+        }
+
+        return function_mapping.get(func_name, func_name)
+
+    def visit_cbuffer(self, cbuffer):
+        """Visit constant buffer (convert to CUDA constant memory)"""
+        self.emit(f"// Constant buffer: {cbuffer.name}")
+        for member in cbuffer.members:
+            member_type = self.convert_crossgl_type_to_cuda(member.vtype)
+            self.emit(f"__constant__ {member_type} {member.name};")
+
+    def visit_ArrayNode(self, node):
+        """Visit array declaration"""
+        element_type = self.convert_crossgl_type_to_cuda(node.element_type)
+
+        if node.size:
+            self.emit(f"{element_type} {node.name}[{node.size}];")
+        else:
+            # Dynamic array - use pointer in CUDA
+            self.emit(f"{element_type}* {node.name};")
+
+    def visit_TernaryOpNode(self, node):
+        """Visit ternary conditional operator"""
+        condition = self.visit(node.condition)
+        true_expr = self.visit(node.true_expr)
+        false_expr = self.visit(node.false_expr)
+        return f"({condition} ? {true_expr} : {false_expr})"
+
+    def visit_list(self, node_list):
+        """Visit a list of nodes"""
+        results = []
+        for node in node_list:
+            result = self.visit(node)
+            if result:
+                results.append(result)
+        return results
+
+    def visit_str(self, node):
+        """Visit string literals"""
+        return node
+
+    def visit_int(self, node):
+        """Visit integer literals"""
+        return str(node)
+
+    def visit_float(self, node):
+        """Visit float literals"""
+        return str(node)
diff --git a/examples/output/ComplexShader.cu b/examples/output/ComplexShader.cu
new file mode 100644
index 0000000..0a556cb
--- /dev/null
+++ b/examples/output/ComplexShader.cu
@@ -0,0 +1,150 @@
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+
+struct Material {
+  float3 albedo;
+  float roughness;
+  float metallic;
+  float3 emissive;
+  float opacity;
+  bool hasNormalMap;
+  sampler2D albedoMap;
+  sampler2D normalMap;
+  sampler2D metallicRoughnessMap;
+};
+
+struct Light {
+  float3 position;
+  float3 color;
+  float intensity;
+  float radius;
+  bool castShadows;
+  float4x4 viewProjection;
+};
+
+struct Scene {
+  Material[4] materials;
+  Light[8] lights;
+  float3 ambientLight;
+  float time;
+  float elapsedTime;
+  int activeLightCount;
+  float4x4 viewMatrix;
+  float4x4 projectionMatrix;
+};
+
+struct VertexInput {
+  float3 position;
+  float3 normal;
+  float3 tangent;
+  float3 bitangent;
+  float2 texCoord0;
+  float2 texCoord1;
+  float4 color;
+  int materialIndex;
+};
+
+struct VertexOutput {
+  float3 worldPosition;
+  float3 worldNormal;
+  float3 worldTangent;
+  float3 worldBitangent;
+  float2 texCoord0;
+  float2 texCoord1;
+  float4 color;
+  float3x3 TBN;
+  int materialIndex;
+  float4 clipPosition;
+};
+
+struct FragmentOutput {
+  float4 color;
+  float4 normalBuffer;
+  float4 positionBuffer;
+  float depth;
+};
+
+struct GlobalUniforms {
+  Scene scene;
+  float3 cameraPosition;
+  float globalRoughness;
+  float2 screenSize;
+  float nearPlane;
+  float farPlane;
+  int frameCount;
+  float[] noiseValues;
+};
+
+float PI;
+
+float EPSILON;
+
+int MAX_ITERATIONS;
+
+float a2;
+
+float num;
+
+float denom;
+
+__device__ float geometrySchlickGGX(float NdotV, float roughness) {
+  float r;
+  float k;
+  float num;
+  float denom;
+  return ['(num DIVIDE fmaxf(denom, EPSILON))'];
+}
+
+__device__ float geometrySmith(float3 N, float3 V, float3 L, float roughness) {
+  float NdotV;
+  float ggx2;
+  return ['(ggx1 MULTIPLY ggx2)'];
+}
+
+__device__ float3 fresnelSchlick(float cosTheta, float3 F0) {
+  return
+      ['(F0 PLUS ((1.0 MINUS F0) MULTIPLY powf(fmaxf((1.0 MINUS cosTheta), 0.0), 5.0)))'];
+}
+
+__device__ float noise3D(float3 p) {
+  float3 i;
+  float3 u;
+  float n000;
+  float n010;
+  float n100;
+  float n110;
+  float n00;
+  float n10;
+  float n0;
+  return ['mix(n0, n1, u.x)'];
+}
+
+__device__ float fbm(float3 p, int octaves, float lacunarity, float gain) {
+  float sum;
+  float amplitude;
+  float frequency;
+  int i;
+  for ((None = 0); (i < octaves); i++) {
+    if ((i >= MAX_ITERATIONS)) {
+    }
+  }
+  return ['sum'];
+}
+
+__device__ float4 samplePlanarProjection(sampler2D tex, float3 worldPos,
+                                         float3 normal) {
+  float3 absNormal;
+  bool useY;
+  float2 uv;
+  if (useX) {
+    if ((normal.x < 0.0)) {
+    }
+  } else if (useY) {
+    if ((normal.y < 0.0)) {
+    }
+  } else {
+    if ((normal.z < 0.0)) {
+    }
+  }
+  return ['texture(tex, uv)'];
+}
diff --git a/examples/output/PerlinNoise.cu b/examples/output/PerlinNoise.cu
new file mode 100644
index 0000000..de860b7
--- /dev/null
+++ b/examples/output/PerlinNoise.cu
@@ -0,0 +1,31 @@
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+
+struct VertexInput {
+  float3 position;
+};
+
+struct VertexOutput {
+  float2 uv;
+  float4 position;
+};
+
+struct FragmentInput {
+  float2 uv;
+};
+
+struct FragmentOutput {
+  float4 color;
+};
+
+__device__ VertexOutput main(VertexInput input) {
+  VertexOutput output;
+  return ['output'];
+}
+
+__device__ FragmentOutput main(FragmentInput input) {
+  FragmentOutput output;
+  float noise;
+  float3 color;
+  return ['output'];
+}
diff --git a/examples/output/SimpleShader.cu b/examples/output/SimpleShader.cu
new file mode 100644
index 0000000..420615c
--- /dev/null
+++ b/examples/output/SimpleShader.cu
@@ -0,0 +1,33 @@
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+
+struct VertexInput {
+  float3 position;
+  float2 texCoord;
+};
+
+struct VertexOutput {
+  float2 uv;
+  float4 position;
+};
+
+struct FragmentInput {
+  float2 uv;
+};
+
+struct FragmentOutput {
+  float4 color;
+};
+
+__device__ VertexOutput main(VertexInput input) {
+  VertexOutput output;
+  return ['output'];
+}
+
+__device__ FragmentOutput main(FragmentInput input) {
+  FragmentOutput output;
+  float r;
+  float g;
+  float b;
+  return ['output'];
+}
diff --git a/setup.py b/setup.py
index e046047..419f1e1 100644
--- a/setup.py
+++ b/setup.py
@@ -13,6 +13,7 @@ setup(
         "crosstl/backend/SPIRV",
         "crosstl/backend/Mojo",
         "crosstl/backend/Rust",
+        "crosstl/backend/CUDA",
     ],
     version="0.0.1.3",
     author="CrossGL team",


Test patch:
diff --git a/.github/workflows/backend-tests.yml b/.github/workflows/backend-tests.yml
index 4d340a9..281c889 100644
--- a/.github/workflows/backend-tests.yml
+++ b/.github/workflows/backend-tests.yml
@@ -13,7 +13,8 @@ jobs:
     runs-on: ${{ matrix.OS }}
     strategy:
       matrix:
-        backend: [directx, metal, GLSL, "slang", "SPIRV", "mojo", "rust"]
+        backend:
+          [directx, metal, GLSL, "slang", "SPIRV", "mojo", "rust", "CUDA"]
         python-version: ["3.8", "3.9", "3.10", "3.11", "3.12", "3.13"]
         OS: [ubuntu-latest, windows-latest, macOS-latest]
       fail-fast: false
diff --git a/.github/workflows/examples-test.yml b/.github/workflows/examples-test.yml
index fae2fde..db8ebec 100644
--- a/.github/workflows/examples-test.yml
+++ b/.github/workflows/examples-test.yml
@@ -46,6 +46,7 @@ jobs:
           python -c "import crosstl; crosstl.translate('ArrayTest.cgl', backend='vulkan', save_shader='output/ArrayTest.spirv')"
           python -c "import crosstl; crosstl.translate('ArrayTest.cgl', backend='mojo', save_shader='output/ArrayTest.mojo')"
           python -c "import crosstl; crosstl.translate('ArrayTest.cgl', backend='rust', save_shader='output/ArrayTest.rs')"
+          python -c "import crosstl; crosstl.translate('ArrayTest.cgl', backend='cuda', save_shader='output/ArrayTest.cu')"
 
           # Simple Shader
           python -c "import crosstl; crosstl.translate('SimpleShader.cgl', backend='metal', save_shader='output/SimpleShader.metal')"
@@ -54,6 +55,7 @@ jobs:
           python -c "import crosstl; crosstl.translate('SimpleShader.cgl', backend='vulkan', save_shader='output/SimpleShader.spirv')"
           python -c "import crosstl; crosstl.translate('SimpleShader.cgl', backend='mojo', save_shader='output/SimpleShader.mojo')"
           python -c "import crosstl; crosstl.translate('SimpleShader.cgl', backend='rust', save_shader='output/SimpleShader.rs')"
+          python -c "import crosstl; crosstl.translate('SimpleShader.cgl', backend='cuda', save_shader='output/SimpleShader.cu')"
 
           # Perlin Noise
           python -c "import crosstl; crosstl.translate('PerlinNoise.cgl', backend='metal', save_shader='output/PerlinNoise.metal')"
@@ -62,6 +64,7 @@ jobs:
           python -c "import crosstl; crosstl.translate('PerlinNoise.cgl', backend='vulkan', save_shader='output/PerlinNoise.spirv')"
           python -c "import crosstl; crosstl.translate('PerlinNoise.cgl', backend='mojo', save_shader='output/PerlinNoise.mojo')"
           python -c "import crosstl; crosstl.translate('PerlinNoise.cgl', backend='rust', save_shader='output/PerlinNoise.rs')"
+          python -c "import crosstl; crosstl.translate('PerlinNoise.cgl', backend='cuda', save_shader='output/PerlinNoise.cu')"
 
           # Complex Shader
           python -c "import crosstl; crosstl.translate('ComplexShader.cgl', backend='metal', save_shader='output/ComplexShader.metal')"
@@ -70,6 +73,7 @@ jobs:
           python -c "import crosstl; crosstl.translate('ComplexShader.cgl', backend='vulkan', save_shader='output/ComplexShader.spirv')"
           python -c "import crosstl; crosstl.translate('ComplexShader.cgl', backend='mojo', save_shader='output/ComplexShader.mojo')"
           python -c "import crosstl; crosstl.translate('ComplexShader.cgl', backend='rust', save_shader='output/ComplexShader.rs')"
+          python -c "import crosstl; crosstl.translate('ComplexShader.cgl', backend='cuda', save_shader='output/ComplexShader.cu')"
         shell: bash
         continue-on-error: true
 
@@ -117,6 +121,12 @@ jobs:
           - { backend: "rust", example: "PerlinNoise", extension: ".rs" }
           - { backend: "rust", example: "ComplexShader", extension: ".rs" }
 
+          # CUDA backend
+          - { backend: "cuda", example: "ArrayTest", extension: ".cu" }
+          - { backend: "cuda", example: "SimpleShader", extension: ".cu" }
+          - { backend: "cuda", example: "PerlinNoise", extension: ".cu" }
+          - { backend: "cuda", example: "ComplexShader", extension: ".cu" }
+
           # Slang backend is excluded as it's still experimental
           # and not fully implemented
 
diff --git a/.github/workflows/translator-tests.yml b/.github/workflows/translator-tests.yml
index 8617a50..d310647 100644
--- a/.github/workflows/translator-tests.yml
+++ b/.github/workflows/translator-tests.yml
@@ -22,6 +22,7 @@ jobs:
             "SPIRV",
             "mojo",
             "rust",
+            "CUDA",
             "general",
           ]
         python-version: ["3.8", "3.9", "3.10", "3.11", "3.12", "3.13"]
diff --git a/examples/output/ArrayTest.cu b/examples/output/ArrayTest.cu
new file mode 100644
index 0000000..bc48b8f
--- /dev/null
+++ b/examples/output/ArrayTest.cu
@@ -0,0 +1,39 @@
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+
+struct VertexInput {
+  float3 position;
+  float2 texCoord;
+};
+
+struct VertexOutput {
+  float2 uv;
+  float4 position;
+};
+
+struct FragmentInput {
+  float2 uv;
+};
+
+struct FragmentOutput {
+  float4 color;
+};
+
+// Constant buffer: TestBuffer
+__constant__ float[4] values;
+__constant__ vec3[2] colors;
+
+__device__ VertexOutput main(VertexInput input) {
+  VertexOutput output;
+  float scale;
+  float3 position;
+  return ['output'];
+}
+
+__device__ FragmentOutput main(FragmentInput input) {
+  FragmentOutput output;
+  float3 color;
+  if ((input.uv.x > 0.5)) {
+  }
+  return ['output'];
+}
diff --git a/examples/test.py b/examples/test.py
index 0d2ab99..b77b721 100644
--- a/examples/test.py
+++ b/examples/test.py
@@ -34,6 +34,12 @@ def main():
                 save_shader="output/ComplexShader.rs",
             )
 
+            cuda_code = crosstl.translate(
+                "ComplexShader.cgl",
+                backend="cuda",
+                save_shader="output/ComplexShader.cu",
+            )
+
             print("Complex shader translation successful!")
         except Exception as e:
             print(f"Warning: Complex shader translation failed: {e}")
@@ -64,6 +70,12 @@ def main():
                 save_shader="output/PerlinNoise.rs",
             )
 
+            cuda_code = crosstl.translate(
+                "PerlinNoise.cgl",
+                backend="cuda",
+                save_shader="output/PerlinNoise.cu",
+            )
+
             print("Perlin noise shader translation successful!")
         except Exception as e:
             print(f"Warning: Perlin noise shader translation failed: {e}")
@@ -90,6 +102,10 @@ def main():
                 "ArrayTest.cgl", backend="rust", save_shader="output/ArrayTest.rs"
             )
 
+            cuda_code = crosstl.translate(
+                "ArrayTest.cgl", backend="cuda", save_shader="output/ArrayTest.cu"
+            )
+
             print("Array test shader translation successful!")
         except Exception as e:
             print(f"Warning: Array test shader translation failed: {e}")
@@ -115,6 +131,12 @@ def main():
             save_shader="output/SimpleShader.rs",
         )
 
+        cuda_code = crosstl.translate(
+            "SimpleShader.cgl",
+            backend="cuda",
+            save_shader="output/SimpleShader.cu",
+        )
+
         print("Simple shader translation successful!")
         return 0
     except Exception as e:
diff --git a/tests/test_backend/test_CUDA/__init__.py b/tests/test_backend/test_CUDA/__init__.py
new file mode 100644
index 0000000..297801e
--- /dev/null
+++ b/tests/test_backend/test_CUDA/__init__.py
@@ -0,0 +1,1 @@
+# CUDA Backend Tests
diff --git a/tests/test_backend/test_CUDA/test_codegen.py b/tests/test_backend/test_CUDA/test_codegen.py
new file mode 100644
index 0000000..32968ad
--- /dev/null
+++ b/tests/test_backend/test_CUDA/test_codegen.py
@@ -0,0 +1,91 @@
+"""Test CUDA Code Generation"""
+
+import pytest
+from crosstl.backend.CUDA.CudaLexer import CudaLexer
+from crosstl.backend.CUDA.CudaParser import CudaParser
+from crosstl.backend.CUDA.CudaCrossGLCodeGen import CudaToCrossGLConverter
+
+
+class TestCudaCodeGen:
+    def test_basic_kernel_conversion(self):
+        """Test basic kernel to CrossGL conversion"""
+        code = """
+        __global__ void simple_kernel(float* data) {
+            int idx = threadIdx.x;
+            data[idx] = idx;
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        codegen = CudaToCrossGLConverter()
+        result = codegen.generate(ast)
+
+        assert "// CUDA to CrossGL conversion" in result
+        assert "// Kernel: simple_kernel" in result
+
+    def test_device_function_conversion(self):
+        """Test device function conversion"""
+        code = """
+        __device__ float add(float a, float b) {
+            return a + b;
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        codegen = CudaToCrossGLConverter()
+        result = codegen.generate(ast)
+
+        assert "// CUDA to CrossGL conversion" in result
+        assert "// Function: add" in result
+
+    def test_multiple_kernels_conversion(self):
+        """Test multiple kernels conversion"""
+        code = """
+        __global__ void kernel1(float* data) {
+            data[threadIdx.x] = 1.0f;
+        }
+        
+        __global__ void kernel2(int* data) {
+            data[threadIdx.x] = 2;
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        codegen = CudaToCrossGLConverter()
+        result = codegen.generate(ast)
+
+        assert "// Kernel: kernel1" in result
+        assert "// Kernel: kernel2" in result
+
+    def test_type_conversion(self):
+        """Test CUDA type to CrossGL type conversion"""
+        codegen = CudaToCrossGLConverter()
+
+        # Test basic types
+        assert codegen.convert_cuda_type_to_crossgl("int") == "i32"
+        assert codegen.convert_cuda_type_to_crossgl("float") == "f32"
+        assert codegen.convert_cuda_type_to_crossgl("double") == "f64"
+        assert codegen.convert_cuda_type_to_crossgl("bool") == "bool"
+        assert codegen.convert_cuda_type_to_crossgl("void") == "void"
+
+    def test_empty_program(self):
+        """Test empty program conversion"""
+        code = ""
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        codegen = CudaToCrossGLConverter()
+        result = codegen.generate(ast)
+
+        assert "// CUDA to CrossGL conversion" in result
diff --git a/tests/test_backend/test_CUDA/test_lexer.py b/tests/test_backend/test_CUDA/test_lexer.py
new file mode 100644
index 0000000..8606430
--- /dev/null
+++ b/tests/test_backend/test_CUDA/test_lexer.py
@@ -0,0 +1,104 @@
+"""Test CUDA Lexer"""
+
+import pytest
+from crosstl.backend.CUDA.CudaLexer import CudaLexer
+
+
+class TestCudaLexer:
+    def test_cuda_keywords(self):
+        """Test CUDA-specific keyword tokenization"""
+        code = "__global__ __device__ __shared__ __constant__"
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+
+        expected_tokens = [
+            ("GLOBAL", "__global__"),
+            ("DEVICE", "__device__"),
+            ("SHARED", "__shared__"),
+            ("CONSTANT", "__constant__"),
+            ("EOF", ""),
+        ]
+
+        assert tokens == expected_tokens
+
+    def test_cuda_builtin_variables(self):
+        """Test CUDA built-in variable tokenization"""
+        code = "threadIdx blockIdx gridDim blockDim"
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+
+        expected_tokens = [
+            ("THREADIDX", "threadIdx"),
+            ("BLOCKIDX", "blockIdx"),
+            ("GRIDDIM", "gridDim"),
+            ("BLOCKDIM", "blockDim"),
+            ("EOF", ""),
+        ]
+
+        assert tokens == expected_tokens
+
+    def test_kernel_launch_syntax(self):
+        """Test kernel launch syntax tokenization"""
+        code = "kernel<<<blocks, threads>>>"
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+
+        expected_tokens = [
+            ("IDENTIFIER", "kernel"),
+            ("KERNEL_LAUNCH_START", "<<<"),
+            ("IDENTIFIER", "blocks"),
+            ("COMMA", ","),
+            ("IDENTIFIER", "threads"),
+            ("KERNEL_LAUNCH_END", ">>>"),
+            ("EOF", ""),
+        ]
+
+        assert tokens == expected_tokens
+
+    def test_cuda_vector_types(self):
+        """Test CUDA vector type tokenization"""
+        code = "float2 float3 float4 int2 int3 int4"
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+
+        expected_tokens = [
+            ("FLOAT2", "float2"),
+            ("FLOAT3", "float3"),
+            ("FLOAT4", "float4"),
+            ("INT2", "int2"),
+            ("INT3", "int3"),
+            ("INT4", "int4"),
+            ("EOF", ""),
+        ]
+
+        assert tokens == expected_tokens
+
+    def test_atomic_operations(self):
+        """Test atomic operation tokenization"""
+        code = "atomicAdd atomicSub atomicMax atomicMin"
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+
+        expected_tokens = [
+            ("ATOMICADD", "atomicAdd"),
+            ("ATOMICSUB", "atomicSub"),
+            ("ATOMICMAX", "atomicMax"),
+            ("ATOMICMIN", "atomicMin"),
+            ("EOF", ""),
+        ]
+
+        assert tokens == expected_tokens
+
+    def test_sync_functions(self):
+        """Test synchronization function tokenization"""
+        code = "__syncthreads __syncwarp"
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+
+        expected_tokens = [
+            ("SYNCTHREADS", "__syncthreads"),
+            ("SYNCWARP", "__syncwarp"),
+            ("EOF", ""),
+        ]
+
+        assert tokens == expected_tokens
diff --git a/tests/test_backend/test_CUDA/test_parser.py b/tests/test_backend/test_CUDA/test_parser.py
new file mode 100644
index 0000000..9ed2373
--- /dev/null
+++ b/tests/test_backend/test_CUDA/test_parser.py
@@ -0,0 +1,118 @@
+"""Test CUDA Parser"""
+
+import pytest
+from crosstl.backend.CUDA.CudaLexer import CudaLexer
+from crosstl.backend.CUDA.CudaParser import CudaParser
+from crosstl.backend.CUDA.CudaAst import ShaderNode
+
+
+class TestCudaParser:
+    def test_simple_kernel_parsing(self):
+        """Test parsing a simple CUDA kernel"""
+        code = """
+        __global__ void simple_kernel(float* data) {
+            int idx = threadIdx.x;
+            data[idx] = idx;
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.kernels) == 1
+        assert ast.kernels[0].name == "simple_kernel"
+
+    def test_device_function_parsing(self):
+        """Test parsing a device function"""
+        code = """
+        __device__ float add(float a, float b) {
+            return a + b;
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.functions) == 1
+        assert ast.functions[0].name == "add"
+        assert "__device__" in ast.functions[0].qualifiers
+
+    def test_shared_memory_parsing(self):
+        """Test parsing shared memory declaration"""
+        code = """
+        __global__ void kernel() {
+            __shared__ float shared_data[256];
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.kernels) == 1
+
+    def test_builtin_variables_parsing(self):
+        """Test parsing CUDA built-in variables"""
+        code = """
+        __global__ void kernel() {
+            int x = threadIdx.x;
+            int y = blockIdx.y;
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.kernels) == 1
+
+    def test_vector_types_parsing(self):
+        """Test parsing CUDA vector types"""
+        code = """
+        __global__ void kernel(float2* data) {
+            float4 vec = make_float4(1.0f, 2.0f, 3.0f, 4.0f);
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.kernels) == 1
+
+    def test_atomic_operations_parsing(self):
+        """Test parsing atomic operations"""
+        code = """
+        __global__ void kernel(int* counter) {
+            atomicAdd(counter, 1);
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.kernels) == 1
+
+    def test_sync_parsing(self):
+        """Test parsing synchronization functions"""
+        code = """
+        __global__ void kernel() {
+            __syncthreads();
+        }
+        """
+        lexer = CudaLexer(code)
+        tokens = lexer.tokenize()
+        parser = CudaParser(tokens)
+        ast = parser.parse()
+
+        assert isinstance(ast, ShaderNode)
+        assert len(ast.kernels) == 1
diff --git a/tests/test_translator/test_codegen/test_CUDA_codegen.py b/tests/test_translator/test_codegen/test_CUDA_codegen.py
new file mode 100644
index 0000000..99b637c
--- /dev/null
+++ b/tests/test_translator/test_codegen/test_CUDA_codegen.py
@@ -0,0 +1,151 @@
+"""Test CUDA Code Generation from CrossGL"""
+
+import pytest
+from crosstl.translator.lexer import Lexer
+from crosstl.translator.parser import Parser
+from crosstl.translator.codegen.cuda_codegen import CudaCodeGen
+
+
+class TestCudaCodeGen:
+    def test_simple_function_generation(self):
+        """Test generating a simple CUDA function from CrossGL"""
+        source_code = """
+        shader TestShader {
+            vertex {
+                void main() {
+                    int x;
+                }
+            }
+        }
+        """
+
+        lexer = Lexer(source_code)
+        parser = Parser(lexer.tokens)
+        ast = parser.parse()
+
+        codegen = CudaCodeGen()
+        cuda_code = codegen.generate(ast)
+
+        assert "#include <cuda_runtime.h>" in cuda_code
+        assert "#include <device_launch_parameters.h>" in cuda_code
+        assert "__device__ void main()" in cuda_code
+
+    def test_compute_shader_to_kernel(self):
+        """Test converting a compute shader to CUDA kernel"""
+        source_code = """
+        shader TestShader {
+            compute {
+                void main() {
+                    int id;
+                }
+            }
+        }
+        """
+
+        lexer = Lexer(source_code)
+        parser = Parser(lexer.tokens)
+        ast = parser.parse()
+
+        codegen = CudaCodeGen()
+        cuda_code = codegen.generate(ast)
+
+        assert "__global__ void main()" in cuda_code
+
+    def test_type_conversion(self):
+        """Test CrossGL to CUDA type conversion"""
+        codegen = CudaCodeGen()
+
+        # Test basic types
+        assert codegen.convert_crossgl_type_to_cuda("i32") == "int"
+        assert codegen.convert_crossgl_type_to_cuda("f32") == "float"
+        assert codegen.convert_crossgl_type_to_cuda("f64") == "double"
+        assert codegen.convert_crossgl_type_to_cuda("bool") == "bool"
+        assert codegen.convert_crossgl_type_to_cuda("void") == "void"
+
+        # Test vector types
+        assert codegen.convert_crossgl_type_to_cuda("vec2<f32>") == "float2"
+        assert codegen.convert_crossgl_type_to_cuda("vec3<f32>") == "float3"
+        assert codegen.convert_crossgl_type_to_cuda("vec4<f32>") == "float4"
+        assert codegen.convert_crossgl_type_to_cuda("vec2<i32>") == "int2"
+
+    def test_function_conversion(self):
+        """Test CrossGL to CUDA function conversion"""
+        codegen = CudaCodeGen()
+
+        # Test math functions
+        assert codegen.convert_builtin_function("sqrt") == "sqrtf"
+        assert codegen.convert_builtin_function("pow") == "powf"
+        assert codegen.convert_builtin_function("sin") == "sinf"
+        assert codegen.convert_builtin_function("cos") == "cosf"
+
+        # Test vector constructors
+        assert codegen.convert_builtin_function("vec2<f32>") == "make_float2"
+        assert codegen.convert_builtin_function("vec3<f32>") == "make_float3"
+        assert codegen.convert_builtin_function("vec4<f32>") == "make_float4"
+
+        # Test atomic operations
+        assert codegen.convert_builtin_function("atomicAdd") == "atomicAdd"
+        assert codegen.convert_builtin_function("atomicExchange") == "atomicExch"
+
+        # Test synchronization
+        assert codegen.convert_builtin_function("workgroupBarrier") == "__syncthreads"
+
+    def test_struct_generation(self):
+        """Test struct generation"""
+        source_code = """
+        shader TestShader {
+            struct Vertex {
+                vec3 position;
+                vec3 normal;
+            }
+        }
+        """
+
+        lexer = Lexer(source_code)
+        parser = Parser(lexer.tokens)
+        ast = parser.parse()
+
+        codegen = CudaCodeGen()
+        cuda_code = codegen.generate(ast)
+
+        assert "struct Vertex {" in cuda_code
+        assert "float3 position;" in cuda_code
+        assert "float3 normal;" in cuda_code
+
+    def test_variable_with_qualifiers(self):
+        """Test variable generation with memory qualifiers"""
+        source_code = """
+        shader TestShader {
+            compute {
+                void main() {
+                    float shared_data;
+                    float constants;
+                }
+            }
+        }
+        """
+
+        lexer = Lexer(source_code)
+        parser = Parser(lexer.tokens)
+        ast = parser.parse()
+
+        codegen = CudaCodeGen()
+        cuda_code = codegen.generate(ast)
+
+        # For now, just check that basic variables are generated
+        assert "float shared_data;" in cuda_code
+        assert "float constants;" in cuda_code
+
+    def test_empty_shader(self):
+        """Test empty shader generation"""
+        source_code = ""
+
+        lexer = Lexer(source_code)
+        parser = Parser(lexer.tokens)
+        ast = parser.parse()
+
+        codegen = CudaCodeGen()
+        cuda_code = codegen.generate(ast)
+
+        assert "#include <cuda_runtime.h>" in cuda_code
+        assert "#include <device_launch_parameters.h>" in cuda_code