feat:add minkowski_hingeloss #70

S1/wut0n_#70/minkowski_hinge_cudacode.py

@@ -0,0 +1,125 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+minkowski_hinge_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+#include <math.h>
+
+// 融合内核：计算Hinge Loss，其中距离为欧几里得距离
+__global__ void minkowski_hinge_p2_kernel(
+    const float* __restrict__ anchor,
+    const float* __restrict__ positive,
+    const float* __restrict__ negative,
+    float* __restrict__ loss,
+    int batch_size,
+    int feature_dim,
+    float margin
+) {
+    int sample_idx = blockIdx.x;
+    if (sample_idx >= batch_size) return;
+    
+    int tid = threadIdx.x;
+    int base = sample_idx * feature_dim;
+    
+    // --- 融合计算 ---
+    // 同时计算 d(anchor, positive)^2 和 d(anchor, negative)^2
+    float sum_sq_pos = 0.0f;
+    float sum_sq_neg = 0.0f;
+    
+    for (int dim = tid; dim < feature_dim; dim += blockDim.x) {
+        float diff_pos = anchor[base + dim] - positive[base + dim];
+        float diff_neg = anchor[base + dim] - negative[base + dim];
+        sum_sq_pos += diff_pos * diff_pos;
+        sum_sq_neg += diff_neg * diff_neg;
+    }
+
+    // 使用共享内存进行归约
+    extern __shared__ float s_mem[];
+    float* s_pos = s_mem;
+    float* s_neg = &s_mem[blockDim.x];
+
+    s_pos[tid] = sum_sq_pos;
+    s_neg[tid] = sum_sq_neg;
+    __syncthreads();
+
+    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            s_pos[tid] += s_pos[tid + s];
+            s_neg[tid] += s_neg[tid + s];
+        }
+        __syncthreads();
+    }
+
+    if (tid == 0) {
+        float d_pos = sqrtf(s_pos[0]);
+        float d_neg = sqrtf(s_neg[0]);
+        
+        // 直接计算最终loss，无需中间结果
+        float hinge_loss = margin + d_pos - d_neg;
+        if (hinge_loss < 0) {
+            hinge_loss = 0.0f; // ReLU
+        }
+        loss[sample_idx] = hinge_loss;
+    }
+}
+
+torch::Tensor minkowski_hinge_p2_cuda(
+    torch::Tensor anchor,
+    torch::Tensor positive,
+    torch::Tensor negative,
+    float margin
+) {
+    TORCH_CHECK(anchor.scalar_type() == torch::kFloat32, "Anchor must be float32");
+    TORCH_CHECK(positive.sizes() == anchor.sizes(), "Positive must have same shape as Anchor");
+    TORCH_CHECK(negative.sizes() == anchor.sizes(), "Negative must have same shape as Anchor");
+    
+    int batch_size = anchor.size(0);
+    int feature_dim = anchor.size(1);
+    
+    auto loss = torch::zeros({batch_size}, anchor.options());
+    
+    const int block_size = 256;
+    size_t shared_mem = 2 * block_size * sizeof(float); // For both sums
+    
+    minkowski_hinge_p2_kernel<<<batch_size, block_size, shared_mem>>>(
+        anchor.data_ptr<float>(),
+        positive.data_ptr<float>(),
+        negative.data_ptr<float>(),
+        loss.data_ptr<float>(),
+        batch_size,
+        feature_dim,
+        margin
+    );
+    
+    return loss;
+}
+"""
+
+minkowski_hinge_cpp_source = """
+torch::Tensor minkowski_hinge_p2_cuda(torch::Tensor anchor, torch::Tensor positive, torch::Tensor negative, float margin);
+"""
+
+minkowski_hinge = load_inline(
+    name="minkowski_hinge_p2",
+    cpp_sources=minkowski_hinge_cpp_source,
+    cuda_sources=minkowski_hinge_source,
+    functions=["minkowski_hinge_p2_cuda"],
+    verbose=True
+)
+
+# --- 修正部分 ---
+class ModelNew(torch.nn.Module):
+    # 1. 修改 __init__ 的签名，移除 p 参数，与 Model 保持一致
+    def __init__(self, margin=1.0):
+        super(ModelNew, self).__init__()
+        # 2. 移除不必要的 p 检查
+        self.margin = margin
+        self.minkowski_hinge = minkowski_hinge
+
+    def forward(self, anchor, positive, negative):
+        return self.minkowski_hinge.minkowski_hinge_p2_cuda(anchor, positive, negative, self.margin)
+
+# 3. 修改 get_init_inputs，只返回 margin，与 torchcode 保持一致
+def get_init_inputs():
+    return [1.0] # margin

S1/wut0n_#70/minkowski_hinge_torchcode.py

@@ -0,0 +1,38 @@
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+    """
+    Hinge Loss implementation using PyTorch's optimized built-in functions.
+    This serves as a strong, fair baseline for our custom CUDA kernel.
+    """
+    def __init__(self, margin=1.0):
+        super(Model, self).__init__()
+        self.margin = margin
+    
+    def forward(self, anchor: torch.Tensor, positive: torch.Tensor, negative: torch.Tensor) -> torch.Tensor:
+        """
+        Compute the Hinge Loss using the most efficient combination of PyTorch's primitives.
+        """
+        # Step 1: Compute Euclidean distances using the highly optimized `torch.pairwise_distance`
+        # This is the standard, performant way to do it in PyTorch.
+        d_pos = torch.pairwise_distance(anchor, positive, p=2)
+        d_neg = torch.pairwise_distance(anchor, negative, p=2)
+        
+        # Step 2: Compute the Hinge Loss formula
+        # Loss = max(margin + d_pos - d_neg, 0)
+        loss = torch.nn.functional.relu(self.margin + d_pos - d_neg)
+        
+        return loss
+
+batch_size = 256
+feature_dim = 512
+
+def get_inputs():
+    a = torch.randn(batch_size, feature_dim)
+    p = torch.randn(batch_size, feature_dim)
+    n = torch.randn(batch_size, feature_dim)
+    return [a, p, n]
+
+def get_init_inputs():
+    return [1.0] # margin

S1/wut0n_#70/prompt.txt

@@ -0,0 +1,134 @@
+You are an expert CUDA programmer tasked with accelerating a PyTorch model by replacing its operators with a highly optimized, custom CUDA kernel. You should consider operator fusion and algorithmic optimizations to achieve maximum speedup.
+
+Here's an example to show you the syntax of inline embedding custom CUDA operators in PyTorch. The example given architecture is a simple ReLU:
+
+python
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+class Model(nn.Module):
+def init(self) -> None:
+super().init()
+
+def forward(self, x):
+    return torch.relu(x)
+def get_inputs():
+x = torch.randn(1, 128).cuda()
+return [x]
+
+def get_init_inputs():
+return []
+
+
+
+The example new architecture with a custom CUDA kernel looks like this:
+
+python
+import torch
+from torch.utils.cpp_extension import load_inline
+
+relu_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+global void relu_kernel(const float* x, float* y, int size) {
+int idx = blockIdx.x * blockDim.x + threadIdx.x;
+if (idx < size) {
+y[idx] = fmaxf(x[idx], 0.f);
+}
+}
+
+torch::Tensor relu_cuda(torch::Tensor x) {
+auto size = x.numel();
+auto y = torch::empty_like(x);
+const int block_size = 256;
+int num_blocks = (size + block_size - 1) / block_size;
+relu_kernel<<<num_blocks, block_size>>>(x.data_ptr<float>(), y.data_ptr<float>(), size);
+return y;
+}
+"""
+
+relu_cpp_source = """
+torch::Tensor relu_cuda(torch::Tensor x);
+"""
+
+Compile the inline CUDA code
+relu = load_inline(
+name=“relu”,
+cpp_sources=relu_cpp_source,
+cuda_sources=relu_source,
+functions=[“relu_cuda”],
+verbose=True
+)
+
+class ModelNew(torch.nn.Module):
+def init(self):
+super(ModelNew, self).init()
+self.relu = relu # The module containing the kernel
+
+def forward(self, x):
+    return self.relu.relu_cuda(x)
+def get_inputs():
+x = torch.randn(1, 128).cuda()
+return [x]
+
+def get_init_inputs():
+return []
+
+
+
+---
+
+Now, you are given the following PyTorch architecture to accelerate. The model computes a Hinge Loss based on distances between anchor, positive, and negative samples. This baseline implementation uses PyTorch's optimized `pairwise_distance` function.
+
+python
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+“”"
+Hinge Loss implementation using PyTorch’s optimized built-in functions.
+This serves as a strong, fair baseline for our custom CUDA kernel.
+“”"
+def init(self, margin=1.0):
+super(Model, self).init()
+self.margin = margin
+
+def forward(self, anchor: torch.Tensor, positive: torch.Tensor, negative: torch.Tensor) -> torch.Tensor:
+    """
+    Compute the Hinge Loss using the most efficient combination of PyTorch's primitives.
+    """
+    # Step 1: Compute Euclidean distances using the highly optimized `torch.pairwise_distance`
+    # This is the standard, performant way to do it in PyTorch.
+    d_pos = torch.pairwise_distance(anchor, positive, p=2)
+    d_neg = torch.pairwise_distance(anchor, negative, p=2)
+    
+    # Step 2: Compute the Hinge Loss formula
+    # Loss = max(margin + d_pos - d_neg, 0)
+    loss = torch.nn.functional.relu(self.margin + d_pos - d_neg)
+    
+    return loss
+batch_size = 256
+feature_dim = 512
+
+def get_inputs():
+a = torch.randn(batch_size, feature_dim)
+p = torch.randn(batch_size, feature_dim)
+n = torch.randn(batch_size, feature_dim)
+return [a, p, n]
+
+def get_init_inputs():
+return [1.0] # margin
+
+Your task is to generate the `ModelNew` architecture with a custom CUDA kernel that fuses the two distance calculations and the final Hinge Loss computation into a single kernel launch, thereby eliminating the intermediate distance tensors.
+
+**CRITICAL REQUIREMENTS:**
+
+1.  **Operator Fusion:** The entire logic—computing both Euclidean distances and then using them to compute the final loss—must be performed inside a **single CUDA kernel**. No intermediate distance tensors should be written to global memory.
+2.  **Algorithmic Specialization:** The implementation must be specialized for Euclidean distance (p=2) for maximum performance. The `__init__` of the new model should only accept `margin`.
+3.  **Kernel Logic:** Each thread block should be responsible for computing the loss for a single sample in the batch. The kernel should compute both squared distances in parallel within the block, using a shared memory reduction. Then, thread 0 of the block should compute the final loss value.
+4.  **Performance Optimization:** Use a multi-threaded reduction within a thread block. Use `extern __shared__` for the reduction and a standard tree-based reduction pattern for numerical stability. The kernel should compute both distances simultaneously to improve data locality.
+5.  **Data Type:** The entire computation must be performed using `float32`. Do not use `double` for any part of the calculation.
+6.  **Code Structure:** Follow the exact structure of the provided example, including `load_inline`, the CUDA source string, the C++ wrapper source string, and the `ModelNew` class. The `get_init_inputs` function must return `[1.0]` to match the baseline.
+7.  **No Fast Math:** Do not use `--use_fast_math` in the compilation flags to ensure numerical accuracy.

S1/wut0n_#70/run_code.py

@@ -0,0 +1,74 @@
+###########################################################
+# 性能和精度验证程序
+###########################################################
+import torch
+import torch.nn as nn
+import time
+from minkowski_hinge_torchcode import Model, get_inputs, get_init_inputs
+from minkowski_hinge_cudacode 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 = 1000
+    
+    # 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 minkowski_hinge 平均执行时间: {torch_time:.6f} 秒")
+    print(f"自定义 CUDA minkowski_hinge 平均执行时间: {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()