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>

__global__ void threshold_scale_negate_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    float threshold,
    float scale,
    int size
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        float x = input[idx];

        if (x > threshold) {
            output[idx] = -(x * scale);
        } else {
            output[idx] = 0.0f;
        }
    }
}

torch::Tensor threshold_scale_negate_cuda(torch::Tensor input, float threshold, float scale) {
    auto output = torch::empty_like(input);
    int size = input.numel();

    const int block_size = 256;
    int num_blocks = (size + block_size - 1) / block_size;

    threshold_scale_negate_kernel<<<num_blocks, block_size>>>(
        input.data_ptr<float>(),
        output.data_ptr<float>(),
        threshold,
        scale,
        size
    );

    return output;
}
"""

cpp_source = """
torch::Tensor threshold_scale_negate_cuda(torch::Tensor input, float threshold, float scale);
"""

module = load_inline(
    name="threshold_scale_negate",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["threshold_scale_negate_cuda"],
    verbose=True
)


class ModelNew(nn.Module):
    def __init__(self, threshold, scale):
        super(ModelNew, self).__init__()
        self.threshold = threshold
        self.scale = scale
        self.module = module

    def forward(self, x):
        return self.module.threshold_scale_negate_cuda(x, self.threshold, self.scale)