finish complex_exp_lop_power #85

S1/uucoco_#85/complex_exp_log_power_cuda.py

@@ -0,0 +1,103 @@
+import torch
+import torch.nn as nn
+from torch.utils.cpp_extension import load_inline
+
+cuda_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+__device__ void complex_exp(float re, float im, float& out_re, float& out_im) {
+    float exp_re = expf(re);
+    out_re = exp_re * cosf(im);
+    out_im = exp_re * sinf(im);
+}
+
+__device__ void complex_log(float re, float im, float& out_re, float& out_im) {
+    out_re = 0.5f * logf(re * re + im * im);
+    out_im = atan2f(im, re);
+}
+
+__device__ void complex_pow(float z_re, float z_im, float p_re, float p_im, float& out_re, float& out_im) {
+    float log_re, log_im;
+    complex_log(z_re, z_im, log_re, log_im);
+
+    float prod_re = p_re * log_re - p_im * log_im;
+    float prod_im = p_re * log_im + p_im * log_re;
+
+    complex_exp(prod_re, prod_im, out_re, out_im);
+}
+
+__global__ void complex_ops_kernel(
+    const float* __restrict__ z_f,
+    float* __restrict__ output,
+    float p_re,
+    float p_im,
+    int batch_size) {
+
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (idx < batch_size) {
+        float z_re = z_f[idx * 2];
+        float z_im = z_f[idx * 2 + 1];
+
+        float y_re, y_im;
+        complex_exp(z_re, z_im, y_re, y_im);
+
+        float w_re, w_im;
+        complex_log(y_re, y_im, w_re, w_im);
+
+        float out_re, out_im;
+        complex_pow(w_re, w_im, p_re, p_im, out_re, out_im);
+
+        output[idx * 2] = out_re;
+        output[idx * 2 + 1] = out_im;
+    }
+}
+
+torch::Tensor complex_ops_cuda(
+    torch::Tensor z_f,
+    float p_re,
+    float p_im) {
+
+    int batch_size = z_f.size(0);
+    auto output = torch::empty_like(z_f);
+
+    int threads = 256;
+    int blocks = (batch_size + threads - 1) / threads;
+
+    complex_ops_kernel<<<blocks, threads>>>(
+        z_f.data_ptr<float>(),
+        output.data_ptr<float>(),
+        p_re,
+        p_im,
+        batch_size
+    );
+
+    return output;
+}
+"""
+
+cpp_source = """
+torch::Tensor complex_ops_cuda(
+    torch::Tensor z_f,
+    float p_re,
+    float p_im);
+"""
+
+cuda_module = load_inline(
+    name="complex_ops_module",
+    cpp_sources=cpp_source,
+    cuda_sources=cuda_source,
+    functions=["complex_ops_cuda"],
+    verbose=True
+)
+
+
+class ModelNew(nn.Module):
+    def __init__(self, p_re, p_im):
+        super(ModelNew, self).__init__()
+        self.p_re = p_re
+        self.p_im = p_im
+
+    def forward(self, z_f):
+        return cuda_module.complex_ops_cuda(z_f, self.p_re, self.p_im)

S1/uucoco_#85/complex_exp_log_power_torch.py

@@ -0,0 +1,30 @@
+import torch
+import torch.nn as nn
+
+
+class Model(nn.Module):
+    def __init__(self, p_re, p_im):
+        super(Model, self).__init__()
+        self.p_c = torch.complex(torch.tensor(p_re), torch.tensor(p_im))
+
+    def forward(self, z_f):
+        z_c = torch.complex(z_f[..., 0], z_f[..., 1])
+
+        y_c = torch.exp(z_c)
+        w_c = torch.log(y_c)
+        out_c = torch.pow(w_c, self.p_c)
+
+        return torch.stack([out_c.real, out_c.imag], dim=-1)
+
+
+batch_size = 1024
+dim = 2
+
+
+def get_inputs():
+    z_f = torch.randn(batch_size, dim)
+    return [z_f]
+
+
+def get_init_inputs():
+    return [0.5, 0.5]

S1/uucoco_#85/prompt.txt

@@ -0,0 +1,110 @@
+You write custom CUDA kernels to replace the pytorch operators in the given GeGLU architecture to get speedups.
+
+You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining chunk+gelu+elementwise_mul), or algorithmic changes (such as optimized memory access patterns). You are only limited by your imagination.
+
+# Technologies Used in This Code
+
+## Core Libraries
+- **PyTorch**: Deep learning framework
+- **CUDA**: NVIDIA GPU parallel computing
+- **C++**: Kernel implementation
+
+## CUDA Components
+- **CUDA kernel**: `complex_ops_kernel`
+- **CUDA math functions**: `expf()`, `logf()`, `sinf()`, `cosf()`, `atan2f()`
+- **Device functions**: `__device__` helper functions for complex operations
+- **Complex arithmetic**: Three custom complex number operations
+
+## Complex Number Operations
+1. **Complex exponential**: exp(z) = exp(real)·(cos(imag) + i·sin(imag))
+2. **Complex logarithm**: log(z) = log|z| + i·arg(z)
+3. **Complex power**: wᵖ = exp(p·log(w))
+- **Composition**: Computes exp(log(exp(z))ᵖ) = exp(z)ᵖ (mathematically)
+
+## Mathematical Implementation
+- **Complex exponential**: Euler's formula implementation
+- **Complex logarithm**: Polar form using atan2 for angle
+- **Complex power**: Via logarithm and exponential (zᵖ = exp(p·log(z)))
+- **Parameterized power**: User-defined complex exponent (p_re + i·p_im)
+
+## Architecture
+- **Device functions**: Reusable complex operation helpers
+- **Element-wise parallelism**: One thread per complex number
+- **Three-step pipeline**: exp → log → pow composition
+- **Batch processing**: Handles multiple complex numbers
+
+## CUDA Optimizations
+- **Modular design**: Separate device functions for each operation
+- **Mathematical identities**: Leverages exp(log(exp(z))) = exp(z)
+- **Efficient operations**: Optimized complex arithmetic
+- **Single kernel**: Fused three operations
+
+## Performance Features
+- **GPU acceleration**: Parallel complex operations
+- **Reusable functions**: Modular device function design
+- **Numerical precision**: Proper complex number handling
+- **Parameterized**: User-defined complex exponent
+
+## Numerical Considerations
+- **Branch cuts**: Complex logarithm has branch cut on negative real axis
+- **Overflow**: exp() can overflow for large real parts
+- **Domain issues**: log(0) undefined
+- **Multiple values**: Complex power may have multiple values
+
+## Mathematical Properties
+- **Identity relation**: exp(log(exp(z))) = exp(z) exactly
+- **Complex power**: Generalization of real exponentiation
+- **Analytic functions**: exp and log are analytic (except branch cuts)
+- **Composition**: exp ∘ log ∘ exp = exp (mathematically)
+
+## Use Case Applications
+- **Complex analysis**: Advanced complex number manipulations
+- **Signal processing**: Complex exponent operations
+- **Physics**: Quantum mechanics wave functions
+- **Mathematics**: Complex function evaluation
+
+## Implementation Details
+- **Tensor shape**: Expects [batch_size, 2] for complex numbers
+- **Device functions**: `__device__` for GPU-only reusable code
+- **Complex exponent**: User provides p_re and p_im parameters
+- **Output format**: Same interleaved complex format as input
+
+## Unique Aspects
+- **Three-operation chain**: Unique exp-log-pow composition
+- **Parameterized power**: Complex-valued exponent
+- **Mathematical identity**: Should compute exp(z)ᵖ (within numerical error)
+- **Device function library**: Reusable complex arithmetic functions
+
+
+
+Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is:
+import torch
+import torch.nn as nn
+
+
+class Model(nn.Module):
+    def __init__(self, p_re, p_im):
+        super(Model, self).__init__()
+        self.p_c = torch.complex(torch.tensor(p_re), torch.tensor(p_im))
+
+    def forward(self, z_f):
+        z_c = torch.complex(z_f[..., 0], z_f[..., 1])
+
+        y_c = torch.exp(z_c)
+        w_c = torch.log(y_c)
+        out_c = torch.pow(w_c, self.p_c)
+
+        return torch.stack([out_c.real, out_c.imag], dim=-1)
+
+
+batch_size = 1024
+dim = 2
+
+
+def get_inputs():
+    z_f = torch.randn(batch_size, dim)
+    return [z_f]
+
+
+def get_init_inputs():
+    return [0.5, 0.5]

S1/uucoco_#85/run_code.py

@@ -0,0 +1,77 @@
+###########################################################
+# 性能和精度验证程序
+###########################################################
+import torch
+import torch.nn as nn
+import time
+from complex_exp_log_power_torch import Model, get_inputs, get_init_inputs
+from complex_exp_log_power_cuda  import ModelNew
+
+
+def run_benchmark():
+    # 检查 CUDA 是否可用
+    if not torch.cuda.is_available():
+        print("CUDA 不可用，请确保您有可用的 NVIDIA GPU 并已正确安装 PyTorch CUDA 版本。")
+        return
+    else:
+        device = torch.device("cuda")
+
+    # 初始化模型
+    init_inputs = get_init_inputs()
+    init_inputs = [
+        x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in init_inputs
+    ]
+    inputs = get_inputs()
+    inputs = [
+        x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in inputs
+    ]
+
+    torch_model = Model(*init_inputs).cuda()
+    cuda_model = ModelNew(*init_inputs).cuda()
+
+    torch_model.eval()
+    cuda_model.eval()
+
+    print("-------------------- 精度对齐验证 --------------------")
+    with torch.no_grad():
+        output_torch = torch_model(*inputs)
+        output_cuda = cuda_model(*inputs)
+
+    precision_flag = torch.allclose(output_torch, output_cuda, rtol=1e-03)
+    if precision_flag:
+        print("✅ 精度对齐：两个模型的输出结果非常接近。")
+    else:
+        print("❌ 精度不一致！")
+
+    print("\n-------------------- 性能加速比测试 --------------------")
+    num_iterations = 100
+
+    # PyTorch 模型计时
+    torch.cuda.synchronize()
+    start_time = time.time()
+    for _ in range(num_iterations):
+        _ = torch_model(*inputs)
+    torch.cuda.synchronize()
+    torch_time = (time.time() - start_time) / num_iterations
+
+    # 自定义 CUDA 内核计时
+    torch.cuda.synchronize()
+    start_time = time.time()
+    for _ in range(num_iterations):
+        _ = cuda_model(*inputs)
+    torch.cuda.synchronize()
+    cuda_time = (time.time() - start_time) / num_iterations
+
+    print(f"PyTorch torch.relu 平均执行时间: {torch_time:.6f} 秒")
+    print(f"自定义 CUDA 内核 平均执行时间: {cuda_time:.6f} 秒")
+    speedup = 0
+    if cuda_time > 0:
+        speedup = torch_time / cuda_time
+        print(f"加速比 (Speedup): {speedup:.2f}x")
+    else:
+        print("CUDA 内核执行时间为0，无法计算加速比。")
+    return precision_flag, speedup
+
+
+if __name__ == "__main__":
+    precision_flag, speedup = run_benchmark()