diff --git a/.claude/settings.local.json b/.claude/settings.local.json
new file mode 100644
index 0000000..5f9d937
--- /dev/null
+++ b/.claude/settings.local.json
@@ -0,0 +1,11 @@
+{
+  "permissions": {
+    "allow": [
+      "Bash(python:*)",
+      "Bash(pip install:*)",
+      "Bash(mv:*)",
+      "Bash(git clone:*)"
+    ],
+    "deny": []
+  }
+}
\ No newline at end of file
diff --git a/CMakeLists.txt b/CMakeLists.txt
index efe03e0..9782a13 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -22,26 +22,101 @@
 
 cmake_minimum_required(VERSION 3.20)
 
-project(DiffRast LANGUAGES CUDA CXX)
+# Detect platform and available languages
+if(APPLE)
+    project(DiffRast LANGUAGES CXX OBJCXX)
+    set(USE_CUDA OFF)
+    set(USE_METAL ON)
+else()
+    find_package(CUDA QUIET)
+    if(CUDA_FOUND)
+        project(DiffRast LANGUAGES CUDA CXX)
+        set(USE_CUDA ON)
+        set(USE_METAL OFF)
+    else()
+        project(DiffRast LANGUAGES CXX)
+        set(USE_CUDA OFF)
+        set(USE_METAL OFF)
+    endif()
+endif()
 
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_EXTENSIONS OFF)
-set(CMAKE_CUDA_STANDARD 17)
 
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
+if(USE_CUDA)
+    set(CMAKE_CUDA_STANDARD 17)
+endif()
 
-add_library(CudaRasterizer
-	cuda_rasterizer/backward.h
-	cuda_rasterizer/backward.cu
-	cuda_rasterizer/forward.h
-	cuda_rasterizer/forward.cu
-	cuda_rasterizer/auxiliary.h
-	cuda_rasterizer/rasterizer_impl.cu
-	cuda_rasterizer/rasterizer_impl.h
-	cuda_rasterizer/rasterizer.h
-)
+# Find OpenMP for CPU fallback
+find_package(OpenMP)
+if(OpenMP_CXX_FOUND)
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}")
+endif()
 
-set_target_properties(CudaRasterizer PROPERTIES CUDA_ARCHITECTURES "70;75;86")
+# Build appropriate backend
+if(USE_CUDA)
+    # CUDA backend (original)
+    add_library(CudaRasterizer
+        cuda_rasterizer/backward.h
+        cuda_rasterizer/backward.cu
+        cuda_rasterizer/forward.h
+        cuda_rasterizer/forward.cu
+        cuda_rasterizer/auxiliary.h
+        cuda_rasterizer/rasterizer_impl.cu
+        cuda_rasterizer/rasterizer_impl.h
+        cuda_rasterizer/rasterizer.h
+    )
+    
+    set_target_properties(CudaRasterizer PROPERTIES CUDA_ARCHITECTURES "70;75;86")
+    target_include_directories(CudaRasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/cuda_rasterizer)
+    target_include_directories(CudaRasterizer PRIVATE third_party/glm ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
+    
+    set(RASTERIZER_LIB CudaRasterizer)
+    
+elseif(USE_METAL AND APPLE)
+    # Metal backend for Apple Silicon
+    add_library(MetalRasterizer
+        metal_rasterizer/rasterizer_metal.h
+        metal_rasterizer/rasterizer_metal.mm
+    )
+    
+    find_library(METAL_FRAMEWORK Metal)
+    find_library(METALKIT_FRAMEWORK MetalKit)
+    find_library(MPS_FRAMEWORK MetalPerformanceShaders)
+    find_library(FOUNDATION_FRAMEWORK Foundation)
+    
+    target_link_libraries(MetalRasterizer 
+        ${METAL_FRAMEWORK} 
+        ${METALKIT_FRAMEWORK} 
+        ${MPS_FRAMEWORK}
+        ${FOUNDATION_FRAMEWORK}
+    )
+    
+    target_include_directories(MetalRasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/metal_rasterizer)
+    target_include_directories(MetalRasterizer PRIVATE third_party/glm)
+    target_compile_definitions(MetalRasterizer PRIVATE USE_METAL=1)
+    
+    set(RASTERIZER_LIB MetalRasterizer)
+    
+else()
+    # CPU fallback
+    add_library(CPURasterizer
+        cpu_rasterizer/rasterizer_cpu.h
+        cpu_rasterizer/rasterizer_cpu.cpp
+    )
+    
+    if(OpenMP_CXX_FOUND)
+        target_link_libraries(CPURasterizer OpenMP::OpenMP_CXX)
+        target_compile_definitions(CPURasterizer PRIVATE USE_OPENMP=1)
+    endif()
+    
+    target_include_directories(CPURasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/cpu_rasterizer)
+    target_include_directories(CPURasterizer PRIVATE third_party/glm)
+    target_compile_definitions(CPURasterizer PRIVATE USE_CPU=1)
+    
+    set(RASTERIZER_LIB CPURasterizer)
+    
+endif()
 
-target_include_directories(CudaRasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/cuda_rasterizer)
-target_include_directories(CudaRasterizer PRIVATE third_party/glm ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
+# Export the selected backend
+set(DIFF_RAST_BACKEND ${RASTERIZER_LIB} PARENT_SCOPE)
diff --git a/cpu_rasterizer/rasterizer_cpu.cpp b/cpu_rasterizer/rasterizer_cpu.cpp
new file mode 100644
index 0000000..01fb820
--- /dev/null
+++ b/cpu_rasterizer/rasterizer_cpu.cpp
@@ -0,0 +1,286 @@
+/*
+ * CPU fallback implementation for triangle rasterization
+ * Copyright (C) 2024 - Mac M1 Port
+ */
+
+#include "rasterizer_cpu.h"
+#include <algorithm>
+#include <cmath>
+#include <cstring>
+#include <iostream>
+
+namespace CPURasterizer {
+
+// Spherical harmonics constants (from original CUDA code)
+constexpr float SH_C0 = 0.28209479177387814f;
+constexpr float SH_C1 = 0.4886025119029199f;
+constexpr float SH_C2[] = {
+    1.0925484305920792f,
+    -1.0925484305920792f,
+    0.31539156525252005f,
+    -1.0925484305920792f,
+    0.5462742152960396f
+};
+
+CPURasterizer::CPURasterizer() {
+    // Initialize temporary buffers
+    temp_buffer.reserve(1024 * 1024);
+    tile_ranges.reserve(1024);
+    point_list.reserve(1024 * 1024);
+}
+
+CPURasterizer::~CPURasterizer() {
+    // Cleanup handled by destructors
+}
+
+glm::vec3 CPURasterizer::computeColorFromSH(int idx, int deg, int max_coeffs, 
+                                           const glm::vec3& means, const glm::vec3& campos, 
+                                           const float* shs, bool* clamped) {
+    glm::vec3 pos = means;
+    glm::vec3 dir = glm::normalize(pos - campos);
+    
+    const glm::vec3* sh = reinterpret_cast<const glm::vec3*>(shs) + idx * max_coeffs;
+    glm::vec3 result = SH_C0 * sh[0];
+    
+    if (deg > 0) {
+        float x = dir.x;
+        float y = dir.y;
+        float z = dir.z;
+        result = result - SH_C1 * y * sh[1] + SH_C1 * z * sh[2] - SH_C1 * x * sh[3];
+        
+        if (deg > 1) {
+            float xx = x * x, yy = y * y, zz = z * z;
+            float xy = x * y, yz = y * z, xz = x * z;
+            result = result +
+                SH_C2[0] * xy * sh[4] +
+                SH_C2[1] * yz * sh[5] +
+                SH_C2[2] * (2.0f * zz - xx - yy) * sh[6] +
+                SH_C2[3] * xz * sh[7] +
+                SH_C2[4] * (xx - yy) * sh[8];
+        }
+    }
+    
+    result += 0.5f;
+    
+    // Clamp colors
+    if (clamped) {
+        *clamped = (result.x < 0 || result.y < 0 || result.z < 0 ||
+                   result.x > 1 || result.y > 1 || result.z > 1);
+    }
+    
+    return glm::clamp(result, 0.0f, 1.0f);
+}
+
+glm::mat3 CPURasterizer::computeCov2D(const glm::vec3& mean, float focal_x, float focal_y, 
+                                     float tan_fovx, float tan_fovy, const float* cov3D, 
+                                     const float* viewmatrix) {
+    // Transform point to camera space
+    glm::mat4 W = glm::mat4(
+        viewmatrix[0], viewmatrix[4], viewmatrix[8], viewmatrix[12],
+        viewmatrix[1], viewmatrix[5], viewmatrix[9], viewmatrix[13],
+        viewmatrix[2], viewmatrix[6], viewmatrix[10], viewmatrix[14],
+        viewmatrix[3], viewmatrix[7], viewmatrix[11], viewmatrix[15]
+    );
+    
+    glm::vec4 p_hom = W * glm::vec4(mean, 1.0f);
+    glm::vec3 p_view = glm::vec3(p_hom) / p_hom.w;
+    
+    // Compute Jacobian of perspective projection
+    float t = p_view.z;
+    float limx = 1.3f * tan_fovx;
+    float limy = 1.3f * tan_fovy;
+    float txtz = t * t;
+    float x = p_view.x;
+    float y = p_view.y;
+    
+    glm::mat3 J = glm::mat3(
+        focal_x / t, 0.0f, -(focal_x * x) / txtz,
+        0.0f, focal_y / t, -(focal_y * y) / txtz,
+        0.0f, 0.0f, 0.0f
+    );
+    
+    // Transform covariance to camera space
+    glm::mat3 W3 = glm::mat3(W);
+    glm::mat3 Vrk = glm::mat3(
+        cov3D[0], cov3D[1], cov3D[2],
+        cov3D[3], cov3D[4], cov3D[5],
+        cov3D[6], cov3D[7], cov3D[8]
+    );
+    
+    glm::mat3 T = W3 * Vrk * glm::transpose(W3);
+    glm::mat3 cov2D = J * T * glm::transpose(J);
+    
+    return cov2D;
+}
+
+void CPURasterizer::computeCov3D(const glm::vec3& scale, float mod, const glm::vec4& rot, 
+                                float* cov3D) {
+    // Build rotation matrix from quaternion
+    float r = rot.x;
+    float x = rot.y;
+    float y = rot.z;
+    float z = rot.w;
+    
+    glm::mat3 R = glm::mat3(
+        1.f - 2.f * (y * y + z * z), 2.f * (x * y - r * z), 2.f * (x * z + r * y),
+        2.f * (x * y + r * z), 1.f - 2.f * (x * x + z * z), 2.f * (y * z - r * x),
+        2.f * (x * z - r * y), 2.f * (y * z + r * x), 1.f - 2.f * (x * x + y * y)
+    );
+    
+    glm::mat3 S = glm::mat3(
+        mod * scale.x, 0, 0,
+        0, mod * scale.y, 0,
+        0, 0, mod * scale.z
+    );
+    
+    glm::mat3 M = S * R;
+    glm::mat3 Sigma = glm::transpose(M) * M;
+    
+    cov3D[0] = Sigma[0][0];
+    cov3D[1] = Sigma[0][1];
+    cov3D[2] = Sigma[0][2];
+    cov3D[3] = Sigma[1][0];
+    cov3D[4] = Sigma[1][1];
+    cov3D[5] = Sigma[1][2];
+    cov3D[6] = Sigma[2][0];
+    cov3D[7] = Sigma[2][1];
+    cov3D[8] = Sigma[2][2];
+}
+
+int CPURasterizer::rasterize_triangles(
+    const RasterizeArgs& args,
+    float* out_color,
+    float* out_depth,
+    int* out_alpha,
+    float* out_cum_alpha,
+    int* radii,
+    float* geomBuffer,
+    float* binningBuffer,
+    float* imgBuffer
+) {
+    const int P = args.P;
+    const int H = args.H;
+    const int W = args.W;
+    
+    // Initialize output
+    std::fill_n(out_color, H * W * 3, 0.0f);
+    if (out_depth) std::fill_n(out_depth, H * W, 0.0f);
+    
+    // Copy background
+    if (args.background) {
+        for (int i = 0; i < H * W; i++) {
+            out_color[i * 3 + 0] = args.background[0];
+            out_color[i * 3 + 1] = args.background[1];
+            out_color[i * 3 + 2] = args.background[2];
+        }
+    }
+    
+    // Simple CPU rasterization - process each point
+    glm::vec3 campos(args.cam_pos[0], args.cam_pos[1], args.cam_pos[2]);
+    
+    #ifdef _OPENMP
+    #pragma omp parallel for
+    #endif
+    for (int idx = 0; idx < P; idx++) {
+        // Get point data
+        glm::vec3 p_world(args.means3D[idx * 3], args.means3D[idx * 3 + 1], args.means3D[idx * 3 + 2]);
+        
+        // Transform to screen space
+        glm::vec4 p_hom = glm::vec4(p_world, 1.0f);
+        
+        // Simple orthographic projection for now
+        float x_screen = (p_world.x + 1.0f) * 0.5f * W;
+        float y_screen = (p_world.y + 1.0f) * 0.5f * H;
+        
+        int px = static_cast<int>(x_screen);
+        int py = static_cast<int>(y_screen);
+        
+        if (px >= 0 && px < W && py >= 0 && py < H) {
+            int pixel_idx = py * W + px;
+            
+            // Simple alpha blending
+            float alpha = args.opacity ? args.opacity[idx] : 1.0f;
+            
+            if (args.colors) {
+                out_color[pixel_idx * 3 + 0] = args.colors[idx * args.D + 0] * alpha;
+                out_color[pixel_idx * 3 + 1] = args.colors[idx * args.D + 1] * alpha;
+                out_color[pixel_idx * 3 + 2] = args.colors[idx * args.D + 2] * alpha;
+            }
+            
+            if (radii) radii[idx] = 1;
+        } else {
+            if (radii) radii[idx] = 0;
+        }
+    }
+    
+    return 0;
+}
+
+int CPURasterizer::rasterize_triangles_backward(
+    const RasterizeArgs& args,
+    const float* grad_out_color,
+    const float* grad_out_depth,
+    const int* radii,
+    const float* geomBuffer,
+    const float* binningBuffer,
+    const float* imgBuffer,
+    float* grad_means3D,
+    float* grad_colors,
+    float* grad_opacity,
+    float* grad_scales,
+    float* grad_rotations,
+    float* grad_cov3D_precomp
+) {
+    const int P = args.P;
+    const int H = args.H;
+    const int W = args.W;
+    
+    // Initialize gradients
+    if (grad_means3D) std::fill_n(grad_means3D, P * 3, 0.0f);
+    if (grad_colors) std::fill_n(grad_colors, P * args.D, 0.0f);
+    if (grad_opacity) std::fill_n(grad_opacity, P, 0.0f);
+    if (grad_scales) std::fill_n(grad_scales, P * 3, 0.0f);
+    if (grad_rotations) std::fill_n(grad_rotations, P * 4, 0.0f);
+    if (grad_cov3D_precomp) std::fill_n(grad_cov3D_precomp, P * 6, 0.0f);
+    
+    // Simple backward pass
+    #ifdef _OPENMP
+    #pragma omp parallel for
+    #endif
+    for (int idx = 0; idx < P; idx++) {
+        if (!radii || radii[idx] == 0) continue;
+        
+        // Simple gradient computation
+        glm::vec3 p_world(args.means3D[idx * 3], args.means3D[idx * 3 + 1], args.means3D[idx * 3 + 2]);
+        
+        float x_screen = (p_world.x + 1.0f) * 0.5f * W;
+        float y_screen = (p_world.y + 1.0f) * 0.5f * H;
+        
+        int px = static_cast<int>(x_screen);
+        int py = static_cast<int>(y_screen);
+        
+        if (px >= 0 && px < W && py >= 0 && py < H) {
+            int pixel_idx = py * W + px;
+            
+            // Backpropagate color gradients
+            if (grad_colors && args.colors) {
+                float alpha = args.opacity ? args.opacity[idx] : 1.0f;
+                grad_colors[idx * args.D + 0] = grad_out_color[pixel_idx * 3 + 0] * alpha;
+                grad_colors[idx * args.D + 1] = grad_out_color[pixel_idx * 3 + 1] * alpha;
+                grad_colors[idx * args.D + 2] = grad_out_color[pixel_idx * 3 + 2] * alpha;
+            }
+            
+            // Backpropagate opacity gradients
+            if (grad_opacity && args.colors) {
+                grad_opacity[idx] = 
+                    grad_out_color[pixel_idx * 3 + 0] * args.colors[idx * args.D + 0] +
+                    grad_out_color[pixel_idx * 3 + 1] * args.colors[idx * args.D + 1] +
+                    grad_out_color[pixel_idx * 3 + 2] * args.colors[idx * args.D + 2];
+            }
+        }
+    }
+    
+    return 0;
+}
+
+} // namespace CPURasterizer
\ No newline at end of file
diff --git a/cpu_rasterizer/rasterizer_cpu.h b/cpu_rasterizer/rasterizer_cpu.h
new file mode 100644
index 0000000..f6fe5dd
--- /dev/null
+++ b/cpu_rasterizer/rasterizer_cpu.h
@@ -0,0 +1,119 @@
+/*
+ * CPU fallback for triangle rasterization
+ * Copyright (C) 2024 - Mac M1 Port
+ */
+
+#pragma once
+
+#include <glm/glm.hpp>
+#include <vector>
+#include <memory>
+
+#ifdef _OPENMP
+#include <omp.h>
+#endif
+
+namespace CPURasterizer {
+
+struct RasterizeArgs {
+    int P;
+    int D; 
+    int M;
+    int R;
+    int H;
+    int W;
+    float focal_x, focal_y;
+    float tan_fovx, tan_fovy;
+    float* background;
+    float* means3D;
+    float* colors;
+    float* opacity;
+    float* scales;
+    float* rotations;
+    float* cov3D_precomp;
+    float* viewmatrix;
+    float* projmatrix;
+    float* cam_pos;
+    bool prefiltered;
+};
+
+class CPURasterizer {
+private:
+    std::vector<float> temp_buffer;
+    std::vector<int> tile_ranges;
+    std::vector<uint64_t> point_list;
+    
+    // Helper functions
+    glm::vec3 computeColorFromSH(int idx, int deg, int max_coeffs, 
+                                const glm::vec3& means, const glm::vec3& campos, 
+                                const float* shs, bool* clamped);
+    
+    glm::mat3 computeCov2D(const glm::vec3& mean, float focal_x, float focal_y, 
+                          float tan_fovx, float tan_fovy, const float* cov3D, 
+                          const float* viewmatrix);
+    
+    void computeCov3D(const glm::vec3& scale, float mod, const glm::vec4& rot, 
+                     float* cov3D);
+    
+    void preprocessCPU(int P, int D, int M, const float* means3D, 
+                      const glm::vec3& campos, const float* colors, 
+                      const float* opacity, const float* scales, 
+                      const glm::vec4* rot, const float* cov3D_precomp,
+                      const float* viewmatrix, const float* projmatrix,
+                      float tan_fovx, float tan_fovy, float focal_x, float focal_y,
+                      int img_height, int img_width, int* radii, 
+                      float2* means2D, float* depths, float* cov3Ds, 
+                      float* rgb, float4* conic_opacity, const dim3& grid, 
+                      uint32_t* tiles_touched, bool prefiltered);
+    
+    void renderCPU(const dim3& grid, const dim3& block, const RasterizeArgs& args,
+                  uint2* ranges, uint32_t* point_list, int W, int H,
+                  float2* means2D, float* colors, float4* conic_opacity,
+                  float* final_T, uint32_t* n_contrib, const float* bg_color,
+                  float* out_color);
+
+public:
+    CPURasterizer();
+    ~CPURasterizer();
+    
+    int rasterize_triangles(
+        const RasterizeArgs& args,
+        float* out_color,
+        float* out_depth,
+        int* out_alpha,
+        float* out_cum_alpha,
+        int* radii,
+        float* geomBuffer,
+        float* binningBuffer,
+        float* imgBuffer
+    );
+    
+    int rasterize_triangles_backward(
+        const RasterizeArgs& args,
+        const float* grad_out_color,
+        const float* grad_out_depth,
+        const int* radii,
+        const float* geomBuffer,
+        const float* binningBuffer,
+        const float* imgBuffer,
+        float* grad_means3D,
+        float* grad_colors,
+        float* grad_opacity,
+        float* grad_scales,
+        float* grad_rotations,
+        float* grad_cov3D_precomp
+    );
+};
+
+// CUDA-style types for compatibility
+struct dim3 {
+    unsigned int x, y, z;
+    dim3(unsigned int _x = 1, unsigned int _y = 1, unsigned int _z = 1) : x(_x), y(_y), z(_z) {}
+};
+
+struct float2 { float x, y; };
+struct float3 { float x, y, z; };
+struct float4 { float x, y, z, w; };
+struct uint2 { unsigned int x, y; };
+
+} // namespace CPURasterizer
\ No newline at end of file
diff --git a/diff_triangle_rasterization/_C.cpython-313-darwin.so b/diff_triangle_rasterization/_C.cpython-313-darwin.so
new file mode 100755
index 0000000..32c77d3
Binary files /dev/null and b/diff_triangle_rasterization/_C.cpython-313-darwin.so differ
diff --git a/diff_triangle_rasterization/__pycache__/__init__.cpython-313.pyc b/diff_triangle_rasterization/__pycache__/__init__.cpython-313.pyc
new file mode 100644
index 0000000..1e1173f
Binary files /dev/null and b/diff_triangle_rasterization/__pycache__/__init__.cpython-313.pyc differ
diff --git a/ext_cpu.cpp b/ext_cpu.cpp
new file mode 100644
index 0000000..3c83748
--- /dev/null
+++ b/ext_cpu.cpp
@@ -0,0 +1,238 @@
+/*
+ * PyTorch extension for CPU backend
+ * Copyright (C) 2024 - Mac M1 Port
+ */
+
+#include <torch/extension.h>
+#include "rasterizer_cpu.h"
+
+using namespace CPURasterizer;
+
+static std::unique_ptr<CPURasterizer> g_cpu_rasterizer;
+
+std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor>
+RasterizetrianglesCUDA(
+    const torch::Tensor& background,
+    const torch::Tensor& means3D,
+    const torch::Tensor& colors,
+    const torch::Tensor& opacity,
+    const torch::Tensor& scales,
+    const torch::Tensor& rotations,
+    const torch::Tensor& cov3D_precomp,
+    const torch::Tensor& viewmatrix,
+    const torch::Tensor& projmatrix,
+    const torch::Tensor& cam_pos,
+    const float tan_fovx,
+    const float tan_fovy,
+    const int image_height,
+    const int image_width,
+    const torch::Tensor& sh,
+    const int degree,
+    const bool prefiltered,
+    const bool debug
+) {
+    if (!g_cpu_rasterizer) {
+        g_cpu_rasterizer = std::make_unique<CPURasterizer>();
+    }
+    
+    const int P = means3D.size(0);
+    const int D = colors.size(1);
+    const int M = 0; // Max SH degree
+    const int R = 0; // Rendered features
+    const int H = image_height;
+    const int W = image_width;
+    const float focal_x = (float)W / (2.0f * tan_fovx);
+    const float focal_y = (float)H / (2.0f * tan_fovy);
+    
+    // Create output tensors
+    auto out_color = torch::zeros({H, W, 3}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto out_depth = torch::zeros({H, W}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto out_alpha = torch::zeros({H, W}, torch::dtype(torch::kInt32).device(torch::kCPU));
+    auto out_cum_alpha = torch::zeros({H, W}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto radii = torch::zeros({P}, torch::dtype(torch::kInt32).device(torch::kCPU));
+    
+    // Prepare arguments
+    RasterizeArgs args;
+    args.P = P;
+    args.D = D;
+    args.M = M;
+    args.R = R;
+    args.H = H;
+    args.W = W;
+    args.focal_x = focal_x;
+    args.focal_y = focal_y;
+    args.tan_fovx = tan_fovx;
+    args.tan_fovy = tan_fovy;
+    args.background = background.data_ptr<float>();
+    args.means3D = means3D.data_ptr<float>();
+    args.colors = colors.data_ptr<float>();
+    args.opacity = opacity.data_ptr<float>();
+    args.scales = scales.data_ptr<float>();
+    args.rotations = rotations.data_ptr<float>();
+    args.cov3D_precomp = cov3D_precomp.numel() > 0 ? cov3D_precomp.data_ptr<float>() : nullptr;
+    args.viewmatrix = viewmatrix.data_ptr<float>();
+    args.projmatrix = projmatrix.data_ptr<float>();
+    args.cam_pos = cam_pos.data_ptr<float>();
+    args.prefiltered = prefiltered;
+    
+    // Temporary buffers
+    auto geomBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto binningBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto imgBuffer = torch::zeros({H * W * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    
+    // Call CPU rasterizer
+    int result = g_cpu_rasterizer->rasterize_triangles(
+        args,
+        out_color.data_ptr<float>(),
+        out_depth.data_ptr<float>(),
+        out_alpha.data_ptr<int>(),
+        out_cum_alpha.data_ptr<float>(),
+        radii.data_ptr<int>(),
+        geomBuffer.data_ptr<float>(),
+        binningBuffer.data_ptr<float>(),
+        imgBuffer.data_ptr<float>()
+    );
+    
+    if (result != 0) {
+        throw std::runtime_error("CPU rasterization failed");
+    }
+    
+    return std::make_tuple(out_color, out_depth, out_alpha, out_cum_alpha, radii);
+}
+
+std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor>
+RasterizetrianglesBackwardCUDA(
+    const torch::Tensor& background,
+    const torch::Tensor& means3D,
+    const torch::Tensor& colors,
+    const torch::Tensor& opacity,
+    const torch::Tensor& scales,
+    const torch::Tensor& rotations,
+    const torch::Tensor& cov3D_precomp,
+    const torch::Tensor& viewmatrix,
+    const torch::Tensor& projmatrix,
+    const torch::Tensor& cam_pos,
+    const float tan_fovx,
+    const float tan_fovy,
+    const torch::Tensor& radii,
+    const torch::Tensor& sh,
+    const int degree,
+    const bool prefiltered,
+    const bool debug,
+    const torch::Tensor& grad_out_color,
+    const torch::Tensor& grad_out_depth
+) {
+    if (!g_cpu_rasterizer) {
+        throw std::runtime_error("CPU rasterizer not initialized");
+    }
+    
+    const int P = means3D.size(0);
+    const int D = colors.size(1);
+    const int M = 0;
+    const int R = 0;
+    const int H = grad_out_color.size(0);
+    const int W = grad_out_color.size(1);
+    const float focal_x = (float)W / (2.0f * tan_fovx);
+    const float focal_y = (float)H / (2.0f * tan_fovy);
+    
+    // Create gradient tensors
+    auto grad_means3D = torch::zeros({P, 3}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_colors = torch::zeros({P, D}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_opacity = torch::zeros({P}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_scales = torch::zeros({P, 3}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_rotations = torch::zeros({P, 4}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_cov3D_precomp = torch::zeros({P, 6}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    
+    // Prepare arguments
+    RasterizeArgs args;
+    args.P = P;
+    args.D = D;
+    args.M = M;
+    args.R = R;
+    args.H = H;
+    args.W = W;
+    args.focal_x = focal_x;
+    args.focal_y = focal_y;
+    args.tan_fovx = tan_fovx;
+    args.tan_fovy = tan_fovy;
+    args.background = background.data_ptr<float>();
+    args.means3D = means3D.data_ptr<float>();
+    args.colors = colors.data_ptr<float>();
+    args.opacity = opacity.data_ptr<float>();
+    args.scales = scales.data_ptr<float>();
+    args.rotations = rotations.data_ptr<float>();
+    args.cov3D_precomp = cov3D_precomp.numel() > 0 ? cov3D_precomp.data_ptr<float>() : nullptr;
+    args.viewmatrix = viewmatrix.data_ptr<float>();
+    args.projmatrix = projmatrix.data_ptr<float>();
+    args.cam_pos = cam_pos.data_ptr<float>();
+    args.prefiltered = prefiltered;
+    
+    // Temporary buffers
+    auto geomBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto binningBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto imgBuffer = torch::zeros({H * W * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    
+    // Call CPU backward pass
+    int result = g_cpu_rasterizer->rasterize_triangles_backward(
+        args,
+        grad_out_color.data_ptr<float>(),
+        grad_out_depth.data_ptr<float>(),
+        radii.data_ptr<int>(),
+        geomBuffer.data_ptr<float>(),
+        binningBuffer.data_ptr<float>(),
+        imgBuffer.data_ptr<float>(),
+        grad_means3D.data_ptr<float>(),
+        grad_colors.data_ptr<float>(),
+        grad_opacity.data_ptr<float>(),
+        grad_scales.data_ptr<float>(),
+        grad_rotations.data_ptr<float>(),
+        grad_cov3D_precomp.data_ptr<float>()
+    );
+    
+    if (result != 0) {
+        throw std::runtime_error("CPU backward pass failed");
+    }
+    
+    return std::make_tuple(grad_means3D, grad_colors, grad_opacity, grad_scales, grad_rotations, grad_cov3D_precomp);
+}
+
+torch::Tensor markVisible(
+    torch::Tensor& means3D,
+    torch::Tensor& viewmatrix,
+    torch::Tensor& projmatrix
+) {
+    // Simple visibility marking for CPU backend
+    const int P = means3D.size(0);
+    auto visible = torch::ones({P}, torch::dtype(torch::kBool).device(torch::kCPU));
+    return visible;
+}
+
+torch::Tensor ComputeRelocationCUDA(
+    torch::Tensor& opacity,
+    torch::Tensor& scales,
+    torch::Tensor& rotations,
+    torch::Tensor& cov3D_precomp,
+    torch::Tensor& means3D,
+    torch::Tensor& viewmatrix,
+    torch::Tensor& projmatrix,
+    float focal_x,
+    float focal_y,
+    float tan_fovx,
+    float tan_fovy,
+    int image_height,
+    int image_width,
+    torch::Tensor& radii,
+    int block_size
+) {
+    // Simple relocation computation for CPU backend
+    const int P = means3D.size(0);
+    auto relocation = torch::zeros({P}, torch::dtype(torch::kInt32).device(torch::kCPU));
+    return relocation;
+}
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    m.def("rasterize_triangles", &RasterizetrianglesCUDA);
+    m.def("rasterize_triangles_backward", &RasterizetrianglesBackwardCUDA);
+    m.def("mark_visible", &markVisible);
+    m.def("compute_relocation", &ComputeRelocationCUDA);
+}
\ No newline at end of file
diff --git a/ext_metal.mm b/ext_metal.mm
new file mode 100644
index 0000000..25f9cda
--- /dev/null
+++ b/ext_metal.mm
@@ -0,0 +1,244 @@
+/*
+ * PyTorch extension for Metal backend
+ * Copyright (C) 2024 - Mac M1 Port
+ */
+
+#include <torch/extension.h>
+#ifdef __APPLE__
+#import <Foundation/Foundation.h>
+#endif
+#include "rasterizer_metal.h"
+
+using namespace MetalRasterizer;
+
+static std::unique_ptr<MetalRasterizer::MetalRasterizerImpl> g_metal_rasterizer;
+
+std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor>
+RasterizetrianglesCUDA(
+    const torch::Tensor& background,
+    const torch::Tensor& means3D,
+    const torch::Tensor& colors,
+    const torch::Tensor& opacity,
+    const torch::Tensor& scales,
+    const torch::Tensor& rotations,
+    const torch::Tensor& cov3D_precomp,
+    const torch::Tensor& viewmatrix,
+    const torch::Tensor& projmatrix,
+    const torch::Tensor& cam_pos,
+    const float tan_fovx,
+    const float tan_fovy,
+    const int image_height,
+    const int image_width,
+    const torch::Tensor& sh,
+    const int degree,
+    const bool prefiltered,
+    const bool debug
+) {
+    if (!g_metal_rasterizer) {
+        g_metal_rasterizer = std::make_unique<MetalRasterizer::MetalRasterizerImpl>();
+        if (!g_metal_rasterizer->initialize()) {
+            throw std::runtime_error("Failed to initialize Metal rasterizer");
+        }
+    }
+    
+    const int P = means3D.size(0);
+    const int D = colors.size(1);
+    const int M = 0; // Max SH degree
+    const int R = 0; // Rendered features
+    const int H = image_height;
+    const int W = image_width;
+    const float focal_x = (float)W / (2.0f * tan_fovx);
+    const float focal_y = (float)H / (2.0f * tan_fovy);
+    
+    // Create output tensors
+    auto out_color = torch::zeros({H, W, 3}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto out_depth = torch::zeros({H, W}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto out_alpha = torch::zeros({H, W}, torch::dtype(torch::kInt32).device(torch::kCPU));
+    auto out_cum_alpha = torch::zeros({H, W}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto radii = torch::zeros({P}, torch::dtype(torch::kInt32).device(torch::kCPU));
+    
+    // Prepare arguments
+    RasterizeArgs args;
+    args.P = P;
+    args.D = D;
+    args.M = M;
+    args.R = R;
+    args.H = H;
+    args.W = W;
+    args.focal_x = focal_x;
+    args.focal_y = focal_y;
+    args.tan_fovx = tan_fovx;
+    args.tan_fovy = tan_fovy;
+    args.background = background.data_ptr<float>();
+    args.means3D = means3D.data_ptr<float>();
+    args.colors = colors.data_ptr<float>();
+    args.opacity = opacity.data_ptr<float>();
+    args.scales = scales.data_ptr<float>();
+    args.rotations = rotations.data_ptr<float>();
+    args.cov3D_precomp = cov3D_precomp.numel() > 0 ? cov3D_precomp.data_ptr<float>() : nullptr;
+    args.viewmatrix = viewmatrix.data_ptr<float>();
+    args.projmatrix = projmatrix.data_ptr<float>();
+    args.cam_pos = cam_pos.data_ptr<float>();
+    args.prefiltered = prefiltered;
+    
+    // Temporary buffers
+    auto geomBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto binningBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto imgBuffer = torch::zeros({H * W * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    
+    // Call Metal rasterizer
+    int result = g_metal_rasterizer->rasterize_triangles(
+        args,
+        out_color.data_ptr<float>(),
+        out_depth.data_ptr<float>(),
+        out_alpha.data_ptr<int>(),
+        out_cum_alpha.data_ptr<float>(),
+        radii.data_ptr<int>(),
+        geomBuffer.data_ptr<float>(),
+        binningBuffer.data_ptr<float>(),
+        imgBuffer.data_ptr<float>()
+    );
+    
+    if (result != 0) {
+        throw std::runtime_error("Metal rasterization failed");
+    }
+    
+    return std::make_tuple(out_color, out_depth, out_alpha, out_cum_alpha, radii);
+}
+
+std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor>
+RasterizetrianglesBackwardCUDA(
+    const torch::Tensor& background,
+    const torch::Tensor& means3D,
+    const torch::Tensor& colors,
+    const torch::Tensor& opacity,
+    const torch::Tensor& scales,
+    const torch::Tensor& rotations,
+    const torch::Tensor& cov3D_precomp,
+    const torch::Tensor& viewmatrix,
+    const torch::Tensor& projmatrix,
+    const torch::Tensor& cam_pos,
+    const float tan_fovx,
+    const float tan_fovy,
+    const torch::Tensor& radii,
+    const torch::Tensor& sh,
+    const int degree,
+    const bool prefiltered,
+    const bool debug,
+    const torch::Tensor& grad_out_color,
+    const torch::Tensor& grad_out_depth
+) {
+    if (!g_metal_rasterizer) {
+        throw std::runtime_error("Metal rasterizer not initialized");
+    }
+    
+    const int P = means3D.size(0);
+    const int D = colors.size(1);
+    const int M = 0;
+    const int R = 0;
+    const int H = grad_out_color.size(0);
+    const int W = grad_out_color.size(1);
+    const float focal_x = (float)W / (2.0f * tan_fovx);
+    const float focal_y = (float)H / (2.0f * tan_fovy);
+    
+    // Create gradient tensors
+    auto grad_means3D = torch::zeros({P, 3}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_colors = torch::zeros({P, D}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_opacity = torch::zeros({P}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_scales = torch::zeros({P, 3}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_rotations = torch::zeros({P, 4}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto grad_cov3D_precomp = torch::zeros({P, 6}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    
+    // Prepare arguments
+    RasterizeArgs args;
+    args.P = P;
+    args.D = D;
+    args.M = M;
+    args.R = R;
+    args.H = H;
+    args.W = W;
+    args.focal_x = focal_x;
+    args.focal_y = focal_y;
+    args.tan_fovx = tan_fovx;
+    args.tan_fovy = tan_fovy;
+    args.background = background.data_ptr<float>();
+    args.means3D = means3D.data_ptr<float>();
+    args.colors = colors.data_ptr<float>();
+    args.opacity = opacity.data_ptr<float>();
+    args.scales = scales.data_ptr<float>();
+    args.rotations = rotations.data_ptr<float>();
+    args.cov3D_precomp = cov3D_precomp.numel() > 0 ? cov3D_precomp.data_ptr<float>() : nullptr;
+    args.viewmatrix = viewmatrix.data_ptr<float>();
+    args.projmatrix = projmatrix.data_ptr<float>();
+    args.cam_pos = cam_pos.data_ptr<float>();
+    args.prefiltered = prefiltered;
+    
+    // Temporary buffers
+    auto geomBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto binningBuffer = torch::zeros({P * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    auto imgBuffer = torch::zeros({H * W * 256}, torch::dtype(torch::kFloat32).device(torch::kCPU));
+    
+    // Call Metal backward pass
+    int result = g_metal_rasterizer->rasterize_triangles_backward(
+        args,
+        grad_out_color.data_ptr<float>(),
+        grad_out_depth.data_ptr<float>(),
+        radii.data_ptr<int>(),
+        geomBuffer.data_ptr<float>(),
+        binningBuffer.data_ptr<float>(),
+        imgBuffer.data_ptr<float>(),
+        grad_means3D.data_ptr<float>(),
+        grad_colors.data_ptr<float>(),
+        grad_opacity.data_ptr<float>(),
+        grad_scales.data_ptr<float>(),
+        grad_rotations.data_ptr<float>(),
+        grad_cov3D_precomp.data_ptr<float>()
+    );
+    
+    if (result != 0) {
+        throw std::runtime_error("Metal backward pass failed");
+    }
+    
+    return std::make_tuple(grad_means3D, grad_colors, grad_opacity, grad_scales, grad_rotations, grad_cov3D_precomp);
+}
+
+torch::Tensor markVisible(
+    torch::Tensor& means3D,
+    torch::Tensor& viewmatrix,
+    torch::Tensor& projmatrix
+) {
+    // Simple visibility marking for Metal backend
+    const int P = means3D.size(0);
+    auto visible = torch::ones({P}, torch::dtype(torch::kBool).device(torch::kCPU));
+    return visible;
+}
+
+torch::Tensor ComputeRelocationCUDA(
+    torch::Tensor& opacity,
+    torch::Tensor& scales,
+    torch::Tensor& rotations,
+    torch::Tensor& cov3D_precomp,
+    torch::Tensor& means3D,
+    torch::Tensor& viewmatrix,
+    torch::Tensor& projmatrix,
+    float focal_x,
+    float focal_y,
+    float tan_fovx,
+    float tan_fovy,
+    int image_height,
+    int image_width,
+    torch::Tensor& radii,
+    int block_size
+) {
+    // Simple relocation computation for Metal backend
+    const int P = means3D.size(0);
+    auto relocation = torch::zeros({P}, torch::dtype(torch::kInt32).device(torch::kCPU));
+    return relocation;
+}
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    m.def("rasterize_triangles", &RasterizetrianglesCUDA);
+    m.def("rasterize_triangles_backward", &RasterizetrianglesBackwardCUDA);
+    m.def("mark_visible", &markVisible);
+    m.def("compute_relocation", &ComputeRelocationCUDA);
+}
\ No newline at end of file
diff --git a/metal_rasterizer/rasterizer_metal.h b/metal_rasterizer/rasterizer_metal.h
new file mode 100644
index 0000000..3505d33
--- /dev/null
+++ b/metal_rasterizer/rasterizer_metal.h
@@ -0,0 +1,85 @@
+/*
+ * Metal backend for triangle rasterization on Apple Silicon
+ * Copyright (C) 2024 - Mac M1 Port
+ */
+
+#pragma once
+
+#ifdef __APPLE__
+#include <Metal/Metal.h>
+#include <MetalPerformanceShaders/MetalPerformanceShaders.h>
+#endif
+
+#include <glm/glm.hpp>
+#include <memory>
+
+namespace MetalRasterizer {
+
+struct RasterizeArgs {
+    int P;
+    int D; 
+    int M;
+    int R;
+    int H;
+    int W;
+    float focal_x, focal_y;
+    float tan_fovx, tan_fovy;
+    float* background;
+    float* means3D;
+    float* colors;
+    float* opacity;
+    float* scales;
+    float* rotations;
+    float* cov3D_precomp;
+    float* viewmatrix;
+    float* projmatrix;
+    float* cam_pos;
+    bool prefiltered;
+};
+
+class MetalRasterizerImpl {
+private:
+#ifdef __APPLE__
+    id<MTLDevice> device;
+    id<MTLCommandQueue> commandQueue;
+    id<MTLLibrary> library;
+    id<MTLComputePipelineState> rasterizePipeline;
+    id<MTLComputePipelineState> backwardPipeline;
+#endif
+    
+public:
+    MetalRasterizerImpl();
+    ~MetalRasterizerImpl();
+    
+    bool initialize();
+    
+    int rasterize_triangles(
+        const RasterizeArgs& args,
+        float* out_color,
+        float* out_depth,
+        int* out_alpha,
+        float* out_cum_alpha,
+        int* radii,
+        float* geomBuffer,
+        float* binningBuffer,
+        float* imgBuffer
+    );
+    
+    int rasterize_triangles_backward(
+        const RasterizeArgs& args,
+        const float* grad_out_color,
+        const float* grad_out_depth,
+        const int* radii,
+        const float* geomBuffer,
+        const float* binningBuffer,
+        const float* imgBuffer,
+        float* grad_means3D,
+        float* grad_colors,
+        float* grad_opacity,
+        float* grad_scales,
+        float* grad_rotations,
+        float* grad_cov3D_precomp
+    );
+};
+
+} // namespace MetalRasterizer
\ No newline at end of file
diff --git a/metal_rasterizer/rasterizer_metal.mm b/metal_rasterizer/rasterizer_metal.mm
new file mode 100644
index 0000000..efea985
--- /dev/null
+++ b/metal_rasterizer/rasterizer_metal.mm
@@ -0,0 +1,284 @@
+/*
+ * Metal backend implementation for triangle rasterization on Apple Silicon
+ * Copyright (C) 2024 - Mac M1 Port
+ */
+
+#include "rasterizer_metal.h"
+#include <iostream>
+
+#ifdef __APPLE__
+#import <Foundation/Foundation.h>
+#import <Metal/Metal.h>
+#import <MetalPerformanceShaders/MetalPerformanceShaders.h>
+#endif
+
+namespace MetalRasterizer {
+
+MetalRasterizerImpl::MetalRasterizerImpl() {
+#ifdef __APPLE__
+    device = nil;
+    commandQueue = nil;
+    library = nil;
+    rasterizePipeline = nil;
+    backwardPipeline = nil;
+#endif
+}
+
+MetalRasterizerImpl::~MetalRasterizerImpl() {
+#ifdef __APPLE__
+    if (rasterizePipeline) {
+        [rasterizePipeline release];
+    }
+    if (backwardPipeline) {
+        [backwardPipeline release];
+    }
+    if (library) {
+        [library release];
+    }
+    if (commandQueue) {
+        [commandQueue release];
+    }
+    if (device) {
+        [device release];
+    }
+#endif
+}
+
+bool MetalRasterizerImpl::initialize() {
+#ifdef __APPLE__
+    // Get default Metal device
+    device = MTLCreateSystemDefaultDevice();
+    if (!device) {
+        std::cerr << "Metal is not supported on this device" << std::endl;
+        return false;
+    }
+    
+    // Create command queue
+    commandQueue = [device newCommandQueue];
+    if (!commandQueue) {
+        std::cerr << "Failed to create Metal command queue" << std::endl;
+        return false;
+    }
+    
+    // Load Metal shaders (would need .metal files)
+    NSError* error = nil;
+    NSString* shaderSource = @R"(
+        #include <metal_stdlib>
+        using namespace metal;
+        
+        struct RasterizeParams {
+            int P, D, M, R, H, W;
+            float focal_x, focal_y;
+            float tan_fovx, tan_fovy;
+        };
+        
+        kernel void rasterize_triangles_kernel(
+            device float* means3D [[buffer(0)]],
+            device float* colors [[buffer(1)]],
+            device float* opacity [[buffer(2)]],
+            device float* out_color [[buffer(3)]],
+            device float* out_depth [[buffer(4)]],
+            constant RasterizeParams& params [[buffer(5)]],
+            uint3 gid [[thread_position_in_grid]]
+        ) {
+            // Basic rasterization kernel - simplified version
+            uint idx = gid.x;
+            if (idx >= params.P) return;
+            
+            // TODO: Implement full rasterization logic
+            // This is a placeholder that copies input to output
+            if (idx < params.P * 3) {
+                out_color[idx] = colors[idx];
+            }
+        }
+        
+        kernel void rasterize_backward_kernel(
+            device float* grad_out_color [[buffer(0)]],
+            device float* grad_means3D [[buffer(1)]],
+            device float* grad_colors [[buffer(2)]],
+            constant RasterizeParams& params [[buffer(3)]],
+            uint3 gid [[thread_position_in_grid]]
+        ) {
+            // Basic backward pass - simplified version
+            uint idx = gid.x;
+            if (idx >= params.P) return;
+            
+            // TODO: Implement full backward logic
+            if (idx < params.P * 3) {
+                grad_colors[idx] = grad_out_color[idx];
+            }
+        }
+    )";
+    
+    library = [device newLibraryWithSource:shaderSource 
+                                   options:nil 
+                                     error:&error];
+    if (!library) {
+        std::cerr << "Failed to create Metal library: " 
+                  << [[error localizedDescription] UTF8String] << std::endl;
+        return false;
+    }
+    
+    // Create compute pipeline states
+    id<MTLFunction> rasterizeFunction = [library newFunctionWithName:@"rasterize_triangles_kernel"];
+    id<MTLFunction> backwardFunction = [library newFunctionWithName:@"rasterize_backward_kernel"];
+    
+    rasterizePipeline = [device newComputePipelineStateWithFunction:rasterizeFunction error:&error];
+    if (!rasterizePipeline) {
+        std::cerr << "Failed to create rasterize pipeline: " 
+                  << [[error localizedDescription] UTF8String] << std::endl;
+        return false;
+    }
+    
+    backwardPipeline = [device newComputePipelineStateWithFunction:backwardFunction error:&error];
+    if (!backwardPipeline) {
+        std::cerr << "Failed to create backward pipeline: " 
+                  << [[error localizedDescription] UTF8String] << std::endl;
+        return false;
+    }
+    
+    return true;
+#else
+    std::cerr << "Metal backend not available on non-Apple platforms" << std::endl;
+    return false;
+#endif
+}
+
+int MetalRasterizerImpl::rasterize_triangles(
+    const RasterizeArgs& args,
+    float* out_color,
+    float* out_depth,
+    int* out_alpha,
+    float* out_cum_alpha,
+    int* radii,
+    float* geomBuffer,
+    float* binningBuffer,
+    float* imgBuffer
+) {
+#ifdef __APPLE__
+    @autoreleasepool {
+        id<MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];
+        id<MTLComputeCommandEncoder> encoder = [commandBuffer computeCommandEncoder];
+        
+        [encoder setComputePipelineState:rasterizePipeline];
+        
+        // Create Metal buffers
+        size_t means3D_size = args.P * 3 * sizeof(float);
+        size_t colors_size = args.P * args.D * sizeof(float);
+        size_t output_size = args.H * args.W * 3 * sizeof(float);
+        
+        id<MTLBuffer> means3DBuffer = [device newBufferWithBytes:args.means3D 
+                                                          length:means3D_size 
+                                                         options:MTLResourceStorageModeShared];
+        id<MTLBuffer> colorsBuffer = [device newBufferWithBytes:args.colors 
+                                                         length:colors_size 
+                                                        options:MTLResourceStorageModeShared];
+        id<MTLBuffer> outputBuffer = [device newBufferWithLength:output_size 
+                                                         options:MTLResourceStorageModeShared];
+        
+        [encoder setBuffer:means3DBuffer offset:0 atIndex:0];
+        [encoder setBuffer:colorsBuffer offset:0 atIndex:1];
+        [encoder setBuffer:outputBuffer offset:0 atIndex:3];
+        
+        // Set compute parameters
+        struct {
+            int P, D, M, R, H, W;
+            float focal_x, focal_y;
+            float tan_fovx, tan_fovy;
+        } params = {
+            args.P, args.D, args.M, args.R, args.H, args.W,
+            args.focal_x, args.focal_y, args.tan_fovx, args.tan_fovy
+        };
+        
+        [encoder setBytes:&params length:sizeof(params) atIndex:5];
+        
+        // Dispatch threads
+        MTLSize gridSize = MTLSizeMake(args.P, 1, 1);
+        MTLSize threadgroupSize = MTLSizeMake(std::min(args.P, 256), 1, 1);
+        
+        [encoder dispatchThreads:gridSize threadsPerThreadgroup:threadgroupSize];
+        [encoder endEncoding];
+        
+        [commandBuffer commit];
+        [commandBuffer waitUntilCompleted];
+        
+        // Copy results back
+        memcpy(out_color, [outputBuffer contents], output_size);
+        
+        return 0;
+    }
+#else
+    std::cerr << "Metal backend not available" << std::endl;
+    return -1;
+#endif
+}
+
+int MetalRasterizerImpl::rasterize_triangles_backward(
+    const RasterizeArgs& args,
+    const float* grad_out_color,
+    const float* grad_out_depth,
+    const int* radii,
+    const float* geomBuffer,
+    const float* binningBuffer,
+    const float* imgBuffer,
+    float* grad_means3D,
+    float* grad_colors,
+    float* grad_opacity,
+    float* grad_scales,
+    float* grad_rotations,
+    float* grad_cov3D_precomp
+) {
+#ifdef __APPLE__
+    @autoreleasepool {
+        id<MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];
+        id<MTLComputeCommandEncoder> encoder = [commandBuffer computeCommandEncoder];
+        
+        [encoder setComputePipelineState:backwardPipeline];
+        
+        // Create Metal buffers for backward pass
+        size_t grad_size = args.H * args.W * 3 * sizeof(float);
+        size_t output_grad_size = args.P * 3 * sizeof(float);
+        
+        id<MTLBuffer> gradInputBuffer = [device newBufferWithBytes:grad_out_color 
+                                                            length:grad_size 
+                                                           options:MTLResourceStorageModeShared];
+        id<MTLBuffer> gradOutputBuffer = [device newBufferWithLength:output_grad_size 
+                                                             options:MTLResourceStorageModeShared];
+        
+        [encoder setBuffer:gradInputBuffer offset:0 atIndex:0];
+        [encoder setBuffer:gradOutputBuffer offset:0 atIndex:1];
+        
+        // Set compute parameters
+        struct {
+            int P, D, M, R, H, W;
+            float focal_x, focal_y;
+            float tan_fovx, tan_fovy;
+        } params = {
+            args.P, args.D, args.M, args.R, args.H, args.W,
+            args.focal_x, args.focal_y, args.tan_fovx, args.tan_fovy
+        };
+        
+        [encoder setBytes:&params length:sizeof(params) atIndex:3];
+        
+        // Dispatch threads
+        MTLSize gridSize = MTLSizeMake(args.P, 1, 1);
+        MTLSize threadgroupSize = MTLSizeMake(std::min(args.P, 256), 1, 1);
+        
+        [encoder dispatchThreads:gridSize threadsPerThreadgroup:threadgroupSize];
+        [encoder endEncoding];
+        
+        [commandBuffer commit];
+        [commandBuffer waitUntilCompleted];
+        
+        // Copy results back
+        memcpy(grad_means3D, [gradOutputBuffer contents], output_grad_size);
+        
+        return 0;
+    }
+#else
+    std::cerr << "Metal backend not available" << std::endl;
+    return -1;
+#endif
+}
+
+} // namespace MetalRasterizer
\ No newline at end of file
diff --git a/setup.py b/setup.py
index 94b274e..71e5686 100644
--- a/setup.py
+++ b/setup.py
@@ -21,25 +21,92 @@
 #
 
 from setuptools import setup
-from torch.utils.cpp_extension import CUDAExtension, BuildExtension
+from torch.utils.cpp_extension import CUDAExtension, CppExtension, BuildExtension
 import os
-os.path.dirname(os.path.abspath(__file__))
+import sys
+import platform
+import torch
+
+def get_extension():
+    """Determine which extension to build based on platform and CUDA availability"""
+    base_dir = os.path.dirname(os.path.abspath(__file__))
+    glm_include = os.path.join(base_dir, "third_party/glm/")
+    
+    # Check if we're on Apple Silicon
+    is_apple_silicon = (platform.system() == "Darwin" and 
+                       (platform.machine() == "arm64" or "arm" in platform.processor().lower()))
+    
+    # Check CUDA availability
+    cuda_available = torch.cuda.is_available() and not is_apple_silicon
+    
+    if cuda_available:
+        # CUDA backend (original)
+        return CUDAExtension(
+            name="diff_triangle_rasterization._C",
+            sources=[
+                "cuda_rasterizer/rasterizer_impl.cu",
+                "cuda_rasterizer/forward.cu", 
+                "cuda_rasterizer/backward.cu",
+                "cuda_rasterizer/utils.cu",
+                "rasterize_points.cu",
+                "ext.cpp"
+            ],
+            extra_compile_args={
+                "nvcc": ["-I" + glm_include, "--use_fast_math"],
+                "cxx": ["-I" + glm_include]
+            }
+        )
+    
+    elif is_apple_silicon:
+        # Metal backend for Apple Silicon
+        return CppExtension(
+            name="diff_triangle_rasterization._C",
+            sources=[
+                "metal_rasterizer/rasterizer_metal.mm",
+                "ext_metal.mm"
+            ],
+            include_dirs=[glm_include, "metal_rasterizer/"],
+            extra_compile_args={
+                "cxx": ["-std=c++17", "-I" + glm_include, "-DUSE_METAL=1", "-ObjC++"]
+            },
+            extra_link_args=[
+                "-framework", "Metal",
+                "-framework", "MetalKit", 
+                "-framework", "MetalPerformanceShaders",
+                "-framework", "Foundation"
+            ]
+        )
+    
+    else:
+        # CPU fallback
+        extra_args = ["-std=c++17", "-I" + glm_include, "-DUSE_CPU=1"]
+        extra_link_args = []
+        
+        # Try to enable OpenMP
+        try:
+            import subprocess
+            result = subprocess.run(["gcc", "--version"], capture_output=True, text=True)
+            if result.returncode == 0:
+                extra_args.extend(["-fopenmp", "-DUSE_OPENMP=1"])
+                extra_link_args.append("-fopenmp")
+        except:
+            pass
+            
+        return CppExtension(
+            name="diff_triangle_rasterization._C",
+            sources=[
+                "cpu_rasterizer/rasterizer_cpu.cpp",
+                "ext_cpu.cpp"
+            ],
+            include_dirs=[glm_include, "cpu_rasterizer/"],
+            extra_compile_args={"cxx": extra_args},
+            extra_link_args=extra_link_args
+        )
 
 setup(
     name="diff_triangle_rasterization",
     packages=['diff_triangle_rasterization'],
-    ext_modules=[
-        CUDAExtension(
-            name="diff_triangle_rasterization._C",
-            sources=[
-            "cuda_rasterizer/rasterizer_impl.cu",
-            "cuda_rasterizer/forward.cu",
-            "cuda_rasterizer/backward.cu",
-            "cuda_rasterizer/utils.cu",
-            "rasterize_points.cu",
-            "ext.cpp"],
-            extra_compile_args={"nvcc": ["-I" + os.path.join(os.path.dirname(os.path.abspath(__file__)), "third_party/glm/"), "--use_fast_math"]})
-        ],
+    ext_modules=[get_extension()],
     cmdclass={
         'build_ext': BuildExtension
     }
diff --git a/third_party/glm b/third_party/glm
index 5c46b9c..2d4c4b4 160000
--- a/third_party/glm
+++ b/third_party/glm
@@ -1 +1 @@
-Subproject commit 5c46b9c07008ae65cb81ab79cd677ecc1934b903
+Subproject commit 2d4c4b4dd31fde06cfffad7915c2b3006402322f