add some activation kernels

.gitattributes

@@ -34,9 +34,3 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
 *.zst filter=lfs diff=lfs merge=lfs -text
 *tfevents* filter=lfs diff=lfs merge=lfs -text
 *.so filter=lfs diff=lfs merge=lfs -text
-build/torch210-cu128-x86_64-windows/activation/_activation_e1b4b08.pyd filter=lfs diff=lfs merge=lfs -text
-media/benches.gif filter=lfs diff=lfs merge=lfs -text
-media/benches.mp4 filter=lfs diff=lfs merge=lfs -text
-build/torch211-cu128-x86_64-windows/_activation_cuda_47d6ce5.pyd filter=lfs diff=lfs merge=lfs -text
-build/torch211-cu128-x86_64-windows/_activation_cuda_edabc93.pyd filter=lfs diff=lfs merge=lfs -text
-build/torch211-cu128-x86_64-windows/_activation_cuda_ce29f4e.pyd filter=lfs diff=lfs merge=lfs -text

README.md

@@ -1,26 +1,10 @@
 ---
 tags:
-  - kernels
+- kernel
 ---
 
+![Status](https://hubwebhook.dholtz.com/shield?repo=kernels-community/activation)
 
 ## Activation
 
-Activation kernels from [vLLM](https://github.com/vllm-project/vllm/blob/main/csrc/activation_kernels.cu).
-
-Kernel source: https://github.com/huggingface/kernels-community/tree/main/activation
-
-### Performance 
-
-
-<img class="dark:hidden border border-gray-200 dark:border-gray-700 rounded-lg" src="media/benches_light_animation.svg" />
-<img class="hidden dark:block border border-gray-200 dark:border-gray-700 rounded-lg" src="media/benches_dark_animation.svg" />
-
-<img class="dark:hidden border border-gray-200 dark:border-gray-700 rounded-lg" src="media/benches_light_latency.svg" />
-<img class="hidden dark:block border border-gray-200 dark:border-gray-700 rounded-lg" src="media/benches_dark_latency.svg" />
-
-<img class="dark:hidden border border-gray-200 dark:border-gray-700 rounded-lg" src="media/benches_light_throughput.svg" />
-<img class="hidden dark:block border border-gray-200 dark:border-gray-700 rounded-lg" src="media/benches_dark_throughput.svg" />
-
-
-
+Activation kernels from [vLLM](https://github.com/vllm-project/vllm/blob/main/csrc/activation_kernels.cu).

activation/activation_kernels.cu

@@ -0,0 +1,225 @@
+#include <ATen/cuda/CUDAContext.h>
+#include <torch/all.h>
+#include <c10/cuda/CUDAGuard.h>
+
+#include <cmath>
+
+#include "cuda_compat.h"
+#include "dispatch_utils.h"
+
+namespace vllm {
+
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&),
+          bool act_first>
+__device__ __forceinline__ scalar_t compute(const scalar_t& x,
+                                            const scalar_t& y) {
+  return act_first ? ACT_FN(x) * y : x * ACT_FN(y);
+}
+// Activation and gating kernel template.
+
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&),
+          bool act_first>
+__global__ void act_and_mul_kernel(
+    scalar_t* __restrict__ out,          // [..., d]
+    const scalar_t* __restrict__ input,  // [..., 2, d]
+    const int d) {
+  const int64_t token_idx = blockIdx.x;
+  for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
+    const scalar_t x = VLLM_LDG(&input[token_idx * 2 * d + idx]);
+    const scalar_t y = VLLM_LDG(&input[token_idx * 2 * d + d + idx]);
+    out[token_idx * d + idx] = compute<scalar_t, ACT_FN, act_first>(x, y);
+  }
+}
+
+template <typename T>
+__device__ __forceinline__ T silu_kernel(const T& x) {
+  // x * sigmoid(x)
+  return (T)(((float)x) / (1.0f + expf((float)-x)));
+}
+
+template <typename T>
+__device__ __forceinline__ T gelu_kernel(const T& x) {
+  // Equivalent to PyTorch GELU with 'none' approximation.
+  // Refer to:
+  // https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L36-L38
+  const float f = (float)x;
+  constexpr float ALPHA = M_SQRT1_2;
+  return (T)(f * 0.5f * (1.0f + ::erf(f * ALPHA)));
+}
+
+template <typename T>
+__device__ __forceinline__ T gelu_tanh_kernel(const T& x) {
+  // Equivalent to PyTorch GELU with 'tanh' approximation.
+  // Refer to:
+  // https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L25-L30
+  const float f = (float)x;
+  constexpr float BETA = M_SQRT2 * M_2_SQRTPI * 0.5f;
+  constexpr float KAPPA = 0.044715;
+  float x_cube = f * f * f;
+  float inner = BETA * (f + KAPPA * x_cube);
+  return (T)(0.5f * f * (1.0f + ::tanhf(inner)));
+}
+
+}  // namespace vllm
+
+// Launch activation and gating kernel.
+// Use ACT_FIRST (bool) indicating whether to apply the activation function
+// first.
+#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, ACT_FIRST)                 \
+  int d = input.size(-1) / 2;                                            \
+  int64_t num_tokens = input.numel() / input.size(-1);                   \
+  dim3 grid(num_tokens);                                                 \
+  dim3 block(std::min(d, 1024));                                         \
+  if (num_tokens == 0) {                                                 \
+    return;                                                              \
+  }                                                                      \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));      \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();          \
+  VLLM_DISPATCH_FLOATING_TYPES(                                          \
+      input.scalar_type(), "act_and_mul_kernel", [&] {                   \
+        vllm::act_and_mul_kernel<scalar_t, KERNEL<scalar_t>, ACT_FIRST>  \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),       \
+                                         input.data_ptr<scalar_t>(), d); \
+      });
+
+void silu_and_mul(torch::Tensor& out,    // [..., d]
+                  torch::Tensor& input)  // [..., 2 * d]
+{
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, true);
+}
+
+void mul_and_silu(torch::Tensor& out,    // [..., d]
+                  torch::Tensor& input)  // [..., 2 * d]
+{
+  // The difference between mul_and_silu and silu_and_mul is that mul_and_silu
+  // applies the silu to the latter half of the input.
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, false);
+}
+
+void gelu_and_mul(torch::Tensor& out,    // [..., d]
+                  torch::Tensor& input)  // [..., 2 * d]
+{
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_kernel, true);
+}
+
+void gelu_tanh_and_mul(torch::Tensor& out,    // [..., d]
+                       torch::Tensor& input)  // [..., 2 * d]
+{
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_tanh_kernel, true);
+}
+
+namespace vllm {
+
+template <typename T>
+__device__ __forceinline__ T fatrelu_kernel(const T& x, const float threshold) {
+  const float f = (float)x;
+  return (T)(f > threshold ? f : 0.0f);
+}
+
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&, const float)>
+__global__ void act_and_mul_kernel_with_param(
+    scalar_t* __restrict__ out, const scalar_t* __restrict__ input, const int d,
+    const float param) {
+  const int64_t token_idx = blockIdx.x;
+  for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
+    const scalar_t x = VLLM_LDG(&input[token_idx * 2 * d + idx]);
+    const scalar_t y = VLLM_LDG(&input[token_idx * 2 * d + d + idx]);
+    out[token_idx * d + idx] = ACT_FN(x, param) * y;
+  }
+}
+
+}  // namespace vllm
+
+#define LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(KERNEL, PARAM)         \
+  int d = input.size(-1) / 2;                                           \
+  int64_t num_tokens = input.numel() / input.size(-1);                  \
+  dim3 grid(num_tokens);                                                \
+  dim3 block(std::min(d, 1024));                                        \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));     \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();         \
+  VLLM_DISPATCH_FLOATING_TYPES(                                         \
+      input.scalar_type(), "act_and_mul_kernel_with_param", [&] {       \
+        vllm::act_and_mul_kernel_with_param<scalar_t, KERNEL<scalar_t>> \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),      \
+                                         input.data_ptr<scalar_t>(), d, \
+                                         PARAM);                        \
+      });
+
+void fatrelu_and_mul(torch::Tensor& out,    // [..., d],
+                     torch::Tensor& input,  // [..., 2 * d]
+                     double threshold) {
+  LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(vllm::fatrelu_kernel, threshold);
+}
+namespace vllm {
+
+// Element-wise activation kernel template.
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&)>
+__global__ void activation_kernel(
+    scalar_t* __restrict__ out,          // [..., d]
+    const scalar_t* __restrict__ input,  // [..., d]
+    const int d) {
+  const int64_t token_idx = blockIdx.x;
+  for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
+    const scalar_t x = VLLM_LDG(&input[token_idx * d + idx]);
+    out[token_idx * d + idx] = ACT_FN(x);
+  }
+}
+
+}  // namespace vllm
+
+// Launch element-wise activation kernel.
+#define LAUNCH_ACTIVATION_KERNEL(KERNEL)                                       \
+  int d = input.size(-1);                                                      \
+  int64_t num_tokens = input.numel() / d;                                      \
+  dim3 grid(num_tokens);                                                       \
+  dim3 block(std::min(d, 1024));                                               \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));            \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();                \
+  VLLM_DISPATCH_FLOATING_TYPES(input.scalar_type(), "activation_kernel", [&] { \
+    vllm::activation_kernel<scalar_t, KERNEL<scalar_t>>                        \
+        <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),                 \
+                                     input.data_ptr<scalar_t>(), d);           \
+  });
+
+namespace vllm {
+
+template <typename T>
+__device__ __forceinline__ T gelu_new_kernel(const T& x) {
+  const float x3 = (float)(x * x * x);
+  const T t = (T)tanhf((T)(0.79788456f * (float)(x + (T)(0.044715f * x3))));
+  return ((T)0.5) * x * (((T)1.0) + t);
+}
+
+template <typename T>
+__device__ __forceinline__ T gelu_fast_kernel(const T& x) {
+  const float f = (float)x;
+  const T t =
+      (T)tanhf(((T)(f * 0.79788456f)) * (((T)1.0) + (T)(0.044715f * f) * x));
+  return ((T)0.5) * x * (((T)1.0) + t);
+}
+
+template <typename T>
+__device__ __forceinline__ T gelu_quick_kernel(const T& x) {
+  // x * sigmoid(1.702 * x)
+  return (T)(((float)x) / (1.0f + expf(-1.702f * (float)x)));
+}
+
+}  // namespace vllm
+
+void gelu_new(torch::Tensor& out,    // [..., d]
+              torch::Tensor& input)  // [..., d]
+{
+  LAUNCH_ACTIVATION_KERNEL(vllm::gelu_new_kernel);
+}
+
+void gelu_fast(torch::Tensor& out,    // [..., d]
+               torch::Tensor& input)  // [..., d]
+{
+  LAUNCH_ACTIVATION_KERNEL(vllm::gelu_fast_kernel);
+}
+
+void gelu_quick(torch::Tensor& out,    // [..., d]
+                torch::Tensor& input)  // [..., d]
+{
+  LAUNCH_ACTIVATION_KERNEL(vllm::gelu_quick_kernel);
+}

activation/cuda_compat.h

@@ -0,0 +1,49 @@
+#pragma once
+
+#ifdef USE_ROCM
+  #include <hip/hip_runtime.h>
+#endif
+
+#if defined(USE_ROCM) && defined(__GFX9__)
+  #define WARP_SIZE 64
+#else
+  #define WARP_SIZE 32
+#endif
+
+#ifndef USE_ROCM
+  #define VLLM_LDG(arg) __ldg(arg)
+#else
+  #define VLLM_LDG(arg) *(arg)
+#endif
+
+#ifndef USE_ROCM
+  #define VLLM_SHFL_XOR_SYNC(var, lane_mask) \
+    __shfl_xor_sync(uint32_t(-1), var, lane_mask)
+  #define VLLM_SHFL_XOR_SYNC_WIDTH(var, lane_mask, width) \
+    __shfl_xor_sync(uint32_t(-1), var, lane_mask, width)
+#else
+  #define VLLM_SHFL_XOR_SYNC(var, lane_mask) __shfl_xor(var, lane_mask)
+  #define VLLM_SHFL_XOR_SYNC_WIDTH(var, lane_mask, width) \
+    __shfl_xor(var, lane_mask, width)
+#endif
+
+#ifndef USE_ROCM
+  #define VLLM_SHFL_SYNC(var, src_lane) __shfl_sync(uint32_t(-1), var, src_lane)
+#else
+  #define VLLM_SHFL_SYNC(var, src_lane) __shfl(var, src_lane)
+#endif
+
+#ifndef USE_ROCM
+  #define VLLM_SHFL_DOWN_SYNC(var, lane_delta) \
+    __shfl_down_sync(uint32_t(-1), var, lane_delta)
+#else
+  #define VLLM_SHFL_DOWN_SYNC(var, lane_delta) __shfl_down(var, lane_delta)
+#endif
+
+#ifndef USE_ROCM
+  #define VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize(FUNC, VAL) \
+    cudaFuncSetAttribute(FUNC, cudaFuncAttributeMaxDynamicSharedMemorySize, VAL)
+#else
+  #define VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize(FUNC, VAL) \
+    hipFuncSetAttribute(FUNC, hipFuncAttributeMaxDynamicSharedMemorySize, VAL)
+#endif

activation/dispatch_utils.h

@@ -0,0 +1,83 @@
+/*
+ * Adapted from
+ * https://github.com/pytorch/pytorch/blob/v2.0.1/aten/src/ATen/Dispatch.h
+ */
+#pragma once
+
+#include <torch/all.h>
+
+// Need a special dispatch case macro since we will nest the FP8 dispatch.
+// Instead of the usual 'scalar_t', this names the dispatched type 'fp8_t'.
+#define AT_DISPATCH_FP8_CASE(enum_type, ...) \
+  AT_PRIVATE_CASE_TYPE_USING_HINT(enum_type, fp8_t, __VA_ARGS__)
+
+#define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
+  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
+  AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)  \
+  AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__)
+
+#define VLLM_DISPATCH_FLOATING_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__))
+
+// ROCm devices might use either fn or fnuz, so set up dispatch table for both.
+// A host-based check at runtime will create a preferred FP8 type for ROCm
+// such that the correct kernel is dispatched.
+#ifdef USE_ROCM
+  #define VLLM_DISPATCH_CASE_FP8_TYPES(...)                          \
+    AT_DISPATCH_FP8_CASE(at::ScalarType::Float8_e4m3fn, __VA_ARGS__) \
+    AT_DISPATCH_FP8_CASE(at::ScalarType::Float8_e4m3fnuz, __VA_ARGS__)
+
+  #define VLLM_DISPATCH_CASE_QUANT_TYPES(...)                      \
+    AT_DISPATCH_CASE(at::ScalarType::Float8_e4m3fn, __VA_ARGS__)   \
+    AT_DISPATCH_CASE(at::ScalarType::Float8_e4m3fnuz, __VA_ARGS__) \
+    AT_DISPATCH_CASE(at::ScalarType::Char, __VA_ARGS__)
+#else
+  #define VLLM_DISPATCH_CASE_FP8_TYPES(...) \
+    AT_DISPATCH_FP8_CASE(at::ScalarType::Float8_e4m3fn, __VA_ARGS__)
+
+  #define VLLM_DISPATCH_CASE_QUANT_TYPES(...)                    \
+    AT_DISPATCH_CASE(at::ScalarType::Float8_e4m3fn, __VA_ARGS__) \
+    AT_DISPATCH_CASE(at::ScalarType::Char, __VA_ARGS__)
+#endif
+
+// When using this dispatch macro, the type is 'fp8_t' not 'scalar_t'.
+// See AT_DISPATCH_FP8_CASE above.
+#define VLLM_DISPATCH_FP8_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_FP8_TYPES(__VA_ARGS__))
+
+#define VLLM_DISPATCH_QUANT_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_QUANT_TYPES(__VA_ARGS__))
+
+#define VLLM_DISPATCH_CASE_FLOATING_AND_BYTE_TYPES(...)   \
+  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__)    \
+  AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)     \
+  AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__) \
+  AT_DISPATCH_CASE(at::ScalarType::Byte, __VA_ARGS__)
+
+#define VLLM_DISPATCH_FLOATING_AND_BYTE_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(TYPE, NAME,                               \
+                     VLLM_DISPATCH_CASE_FLOATING_AND_BYTE_TYPES(__VA_ARGS__))
+
+#define VLLM_DISPATCH_CASE_INTEGRAL_TYPES(...)         \
+  AT_DISPATCH_CASE(at::ScalarType::Byte, __VA_ARGS__)  \
+  AT_DISPATCH_CASE(at::ScalarType::Char, __VA_ARGS__)  \
+  AT_DISPATCH_CASE(at::ScalarType::Short, __VA_ARGS__) \
+  AT_DISPATCH_CASE(at::ScalarType::Int, __VA_ARGS__)   \
+  AT_DISPATCH_CASE(at::ScalarType::Long, __VA_ARGS__)
+
+#define VLLM_DISPATCH_CASE_INTEGRAL_AND_UNSIGNED_TYPES(...) \
+  AT_DISPATCH_CASE(at::ScalarType::Byte, __VA_ARGS__)       \
+  AT_DISPATCH_CASE(at::ScalarType::Char, __VA_ARGS__)       \
+  AT_DISPATCH_CASE(at::ScalarType::Short, __VA_ARGS__)      \
+  AT_DISPATCH_CASE(at::ScalarType::Int, __VA_ARGS__)        \
+  AT_DISPATCH_CASE(at::ScalarType::Long, __VA_ARGS__)       \
+  AT_DISPATCH_CASE(at::ScalarType::UInt16, __VA_ARGS__)     \
+  AT_DISPATCH_CASE(at::ScalarType::UInt32, __VA_ARGS__)     \
+  AT_DISPATCH_CASE(at::ScalarType::UInt64, __VA_ARGS__)
+
+#define VLLM_DISPATCH_INTEGRAL_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_INTEGRAL_TYPES(__VA_ARGS__))
+
+#define VLLM_DISPATCH_INTEGRAL_AND_UNSIGNED_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(                                              \
+      TYPE, NAME, VLLM_DISPATCH_CASE_INTEGRAL_AND_UNSIGNED_TYPES(__VA_ARGS__))

benchmarks/benchmark.py

@@ -1,5 +0,0 @@
-from kernels.benchmarks import SiluAndMulBenchmark
-
-
-class SiluWorkloads(SiluAndMulBenchmark):
-    pass

build.toml

@@ -0,0 +1,18 @@
+[general]
+name = "activation"
+universal = false
+
+[torch]
+src = [
+    "torch-ext/torch_binding.cpp",
+    "torch-ext/torch_binding.h",
+]
+
+[kernel.activation]
+backend = "cuda"
+depends = ["torch"]
+src = [
+    "activation/activation_kernels.cu",
+    "activation/cuda_compat.h",
+    "activation/dispatch_utils.h",
+]

build/torch210-cu128-x86_64-windows/activation/__init__.py

@@ -1,75 +0,0 @@
-import torch
-
-from ._ops import ops
-
-from . import layers
-
-
-def silu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu_and_mul(out, x)
-    return out
-
-
-def mul_and_silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.mul_and_silu(out, x)
-    return out
-
-
-def gelu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_and_mul(out, x)
-    return out
-
-
-def gelu_tanh_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh_and_mul(out, x)
-    return out
-
-
-def fatrelu_and_mul(out: torch.Tensor, x: torch.Tensor, threshold: float = 0.0) -> None:
-    ops.fatrelu_and_mul(out, x, threshold)
-    return out
-
-
-def gelu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu(out, x)
-    return out
-
-def silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu(out, x)
-    return out
-
-
-def gelu_tanh(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh(out, x)
-    return out
-
-
-def gelu_fast(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_fast(out, x)
-    return out
-
-
-def gelu_new(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_new(out, x)
-    return out
-
-
-def gelu_quick(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_quick(out, x)
-    return out
-
-
-__all__ = [
-    "silu_and_mul",
-    "mul_and_silu",
-    "gelu_and_mul",
-    "gelu_tanh_and_mul",
-    "fatrelu_and_mul",
-    "gelu_fast",
-    "gelu_new",
-    "gelu_quick",
-    "gelu_tanh",
-    "silu",
-    "gelu",
-    "layers",
-]

build/torch210-cu128-x86_64-windows/activation/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _activation_e1b4b08
-ops = torch.ops._activation_e1b4b08
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_activation_e1b4b08::{op_name}"

build/torch210-cu128-x86_64-windows/activation/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cu128-x86_64-windows/metadata.json

@@ -1,4 +0,0 @@
-{
-  "version": 1,
-  "python-depends": []
-}

build/torch210-cxx11-cu126-aarch64-linux/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-cxx11-cu126-aarch64-linux/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cxx11-cu126-aarch64-linux/metadata.json

@@ -1,18 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "7.0",
-      "7.2",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0+PTX"
-    ]
-  }
-}

build/torch210-cxx11-cu126-x86_64-linux/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-cxx11-cu126-x86_64-linux/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cxx11-cu126-x86_64-linux/metadata.json

@@ -1,18 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "7.0",
-      "7.2",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0+PTX"
-    ]
-  }
-}

build/torch210-cxx11-cu128-aarch64-linux/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-cxx11-cu128-aarch64-linux/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cxx11-cu128-aarch64-linux/metadata.json

@@ -1,21 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "10.1",
-      "12.0+PTX",
-      "7.0",
-      "7.2",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch210-cxx11-cu128-x86_64-linux/_activation_cuda_ccf9ce9.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:c2946ede3bdc0bffe6e0c1f68dbcdbbdcfcd9c8f2d53723f17159005aac6b247
-size 4406632

build/torch210-cxx11-cu128-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _activation_cuda_ccf9ce9
-ops = torch.ops._activation_cuda_ccf9ce9
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_activation_cuda_ccf9ce9::{op_name}"

build/torch210-cxx11-cu128-x86_64-linux/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-cxx11-cu128-x86_64-linux/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cxx11-cu128-x86_64-linux/metadata.json

@@ -1,21 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "10.1",
-      "12.0+PTX",
-      "7.0",
-      "7.2",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch210-cxx11-cu130-aarch64-linux/__init__.py

@@ -1,75 +0,0 @@
-import torch
-
-from ._ops import ops
-
-from . import layers
-
-
-def silu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu_and_mul(out, x)
-    return out
-
-
-def mul_and_silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.mul_and_silu(out, x)
-    return out
-
-
-def gelu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_and_mul(out, x)
-    return out
-
-
-def gelu_tanh_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh_and_mul(out, x)
-    return out
-
-
-def fatrelu_and_mul(out: torch.Tensor, x: torch.Tensor, threshold: float = 0.0) -> None:
-    ops.fatrelu_and_mul(out, x, threshold)
-    return out
-
-
-def gelu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu(out, x)
-    return out
-
-def silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu(out, x)
-    return out
-
-
-def gelu_tanh(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh(out, x)
-    return out
-
-
-def gelu_fast(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_fast(out, x)
-    return out
-
-
-def gelu_new(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_new(out, x)
-    return out
-
-
-def gelu_quick(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_quick(out, x)
-    return out
-
-
-__all__ = [
-    "silu_and_mul",
-    "mul_and_silu",
-    "gelu_and_mul",
-    "gelu_tanh_and_mul",
-    "fatrelu_and_mul",
-    "gelu_fast",
-    "gelu_new",
-    "gelu_quick",
-    "gelu_tanh",
-    "silu",
-    "gelu",
-    "layers",
-]

build/torch210-cxx11-cu130-aarch64-linux/_activation_cuda_ccf9ce9.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:6c836df90ec0e7853341065857441a9a48891ae86a15d5fd0d044fb03b6320ea
-size 4294136

build/torch210-cxx11-cu130-aarch64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _activation_cuda_ccf9ce9
-ops = torch.ops._activation_cuda_ccf9ce9
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_activation_cuda_ccf9ce9::{op_name}"

build/torch210-cxx11-cu130-aarch64-linux/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-cxx11-cu130-aarch64-linux/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cxx11-cu130-aarch64-linux/metadata.json

@@ -1,19 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "11.0",
-      "12.0+PTX",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch210-cxx11-cu130-x86_64-linux/__init__.py

@@ -1,75 +0,0 @@
-import torch
-
-from ._ops import ops
-
-from . import layers
-
-
-def silu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu_and_mul(out, x)
-    return out
-
-
-def mul_and_silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.mul_and_silu(out, x)
-    return out
-
-
-def gelu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_and_mul(out, x)
-    return out
-
-
-def gelu_tanh_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh_and_mul(out, x)
-    return out
-
-
-def fatrelu_and_mul(out: torch.Tensor, x: torch.Tensor, threshold: float = 0.0) -> None:
-    ops.fatrelu_and_mul(out, x, threshold)
-    return out
-
-
-def gelu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu(out, x)
-    return out
-
-def silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu(out, x)
-    return out
-
-
-def gelu_tanh(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh(out, x)
-    return out
-
-
-def gelu_fast(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_fast(out, x)
-    return out
-
-
-def gelu_new(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_new(out, x)
-    return out
-
-
-def gelu_quick(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_quick(out, x)
-    return out
-
-
-__all__ = [
-    "silu_and_mul",
-    "mul_and_silu",
-    "gelu_and_mul",
-    "gelu_tanh_and_mul",
-    "fatrelu_and_mul",
-    "gelu_fast",
-    "gelu_new",
-    "gelu_quick",
-    "gelu_tanh",
-    "silu",
-    "gelu",
-    "layers",
-]

build/torch210-cxx11-cu130-x86_64-linux/_activation_cuda_ccf9ce9.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:57ee9bcb4b40ac6b5fa006806f222b3961ea49513e425b4b3caaf60c760d604d
-size 4186584

build/torch210-cxx11-cu130-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _activation_cuda_ccf9ce9
-ops = torch.ops._activation_cuda_ccf9ce9
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_activation_cuda_ccf9ce9::{op_name}"

build/torch210-cxx11-cu130-x86_64-linux/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-cxx11-cu130-x86_64-linux/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-cxx11-cu130-x86_64-linux/metadata.json

@@ -1,19 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "11.0",
-      "12.0+PTX",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch210-metal-aarch64-darwin/__init__.py

@@ -1,75 +0,0 @@
-import torch
-
-from ._ops import ops
-
-from . import layers
-
-
-def silu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu_and_mul(out, x)
-    return out
-
-
-def mul_and_silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.mul_and_silu(out, x)
-    return out
-
-
-def gelu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_and_mul(out, x)
-    return out
-
-
-def gelu_tanh_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh_and_mul(out, x)
-    return out
-
-
-def fatrelu_and_mul(out: torch.Tensor, x: torch.Tensor, threshold: float = 0.0) -> None:
-    ops.fatrelu_and_mul(out, x, threshold)
-    return out
-
-
-def gelu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu(out, x)
-    return out
-
-def silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu(out, x)
-    return out
-
-
-def gelu_tanh(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh(out, x)
-    return out
-
-
-def gelu_fast(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_fast(out, x)
-    return out
-
-
-def gelu_new(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_new(out, x)
-    return out
-
-
-def gelu_quick(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_quick(out, x)
-    return out
-
-
-__all__ = [
-    "silu_and_mul",
-    "mul_and_silu",
-    "gelu_and_mul",
-    "gelu_tanh_and_mul",
-    "fatrelu_and_mul",
-    "gelu_fast",
-    "gelu_new",
-    "gelu_quick",
-    "gelu_tanh",
-    "silu",
-    "gelu",
-    "layers",
-]

build/torch210-metal-aarch64-darwin/_activation_metal_ccf9ce9.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:83cc62ee4c6546a6b77f4ad26285861bd36894ddacb4836e38db664e75c317f3
-size 205056

build/torch210-metal-aarch64-darwin/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _activation_metal_ccf9ce9
-ops = torch.ops._activation_metal_ccf9ce9
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_activation_metal_ccf9ce9::{op_name}"

build/torch210-metal-aarch64-darwin/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch210-metal-aarch64-darwin/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch210-metal-aarch64-darwin/metadata.json

@@ -1,8 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "metal"
-  }
-}

build/torch211-cu128-x86_64-windows/__init__.py

@@ -1,75 +0,0 @@
-import torch
-
-from ._ops import ops
-
-from . import layers
-
-
-def silu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu_and_mul(out, x)
-    return out
-
-
-def mul_and_silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.mul_and_silu(out, x)
-    return out
-
-
-def gelu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_and_mul(out, x)
-    return out
-
-
-def gelu_tanh_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh_and_mul(out, x)
-    return out
-
-
-def fatrelu_and_mul(out: torch.Tensor, x: torch.Tensor, threshold: float = 0.0) -> None:
-    ops.fatrelu_and_mul(out, x, threshold)
-    return out
-
-
-def gelu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu(out, x)
-    return out
-
-def silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu(out, x)
-    return out
-
-
-def gelu_tanh(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh(out, x)
-    return out
-
-
-def gelu_fast(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_fast(out, x)
-    return out
-
-
-def gelu_new(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_new(out, x)
-    return out
-
-
-def gelu_quick(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_quick(out, x)
-    return out
-
-
-__all__ = [
-    "silu_and_mul",
-    "mul_and_silu",
-    "gelu_and_mul",
-    "gelu_tanh_and_mul",
-    "fatrelu_and_mul",
-    "gelu_fast",
-    "gelu_new",
-    "gelu_quick",
-    "gelu_tanh",
-    "silu",
-    "gelu",
-    "layers",
-]

build/torch211-cu128-x86_64-windows/_activation_cuda_ce29f4e.pyd

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:b3b882f56a5eaffdb32e94f924b611a4df6ee97400021b2dc34be9a062b941a7
-size 2464256

build/torch211-cu128-x86_64-windows/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _activation_cuda_ce29f4e
-ops = torch.ops._activation_cuda_ce29f4e
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_activation_cuda_ce29f4e::{op_name}"

build/torch211-cu128-x86_64-windows/activation/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch211-cu128-x86_64-windows/layers.py

@@ -1,201 +0,0 @@
-import torch
-import torch.nn as nn
-
-from ._ops import ops
-
-
-class SiluAndMul(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> silu(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.silu_and_mul(out, x)
-        return out
-
-class Silu(nn.Module):
-    """An activation function for SiLU.
-
-    The function computes x -> silu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.silu(out, x)
-        return out
-
-class Gelu(nn.Module):
-    """An activation function for GELU.
-
-    The function computes x -> gelu(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu(out, x)
-        return out
-
-class GeluTanh(nn.Module):
-    """An activation function for GELU with `tanh` approximation.
-
-    The function computes x -> gelu_tanh(x).
-
-    Shapes:
-        x: (num_tokens, d) or (batch_size, seq_len, d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_tanh(out, x)
-        return out
-
-
-class MulAndSilu(nn.Module):
-    """An activation function for SwiGLU.
-
-    The function computes x -> x[:d] * silu(x[d:]) where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.mul_and_silu(out, x)
-        return out
-
-
-class GeluAndMul(nn.Module):
-    """An activation function for GeGLU.
-
-    The function computes x -> GELU(x[:d]) * x[d:] where d = x.shape[-1] // 2.
-
-    Shapes:
-        x: (batch_size, seq_len, 2 * d) or (num_tokens, 2 * d)
-        return: (batch_size, seq_len, d) or (num_tokens, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_and_mul(out, x)
-        return out
-
-
-class GeluTanhAndMul(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.gelu_tanh_and_mul(out, x)
-        return out
-
-
-class FatreluAndMul(nn.Module):
-    """An activation function for FATReLU.
-
-    The function computes x -> FATReLU(x[:d]) * x[d:] where
-    d = x.shape[-1] // 2.
-    This is used in openbmb/MiniCPM-S-1B-sft.
-
-    Shapes:
-        x: (num_tokens, 2 * d) or (batch_size, seq_len, 2 * d)
-        return: (num_tokens, d) or (batch_size, seq_len, d)
-    """
-
-    can_torch_compile: bool = True
-
-    def __init__(self, threshold: float = 0.0):
-        super().__init__()
-        self.threshold = threshold
-
-    def forward(self, x: torch.Tensor):
-        if not x.is_contiguous():
-            x = x.contiguous()
-        d = x.shape[-1] // 2
-        output_shape = x.shape[:-1] + (d,)
-        out = torch.empty(output_shape, dtype=x.dtype, device=x.device)
-        ops.fatrelu_and_mul(out, x, self.threshold)
-        return out
-
-
-class FastGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_fast(out, x)
-        return out
-
-
-class NewGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_new(out, x)
-        return out
-
-
-class QuickGELU(nn.Module):
-    can_torch_compile: bool = True
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        if not x.is_contiguous():
-            x = x.contiguous()
-        out = torch.empty_like(x)
-        ops.gelu_quick(out, x)
-        return out

build/torch211-cu128-x86_64-windows/metadata.json

@@ -1,21 +0,0 @@
-{
-  "version": 1,
-  "license": "Apache-2.0",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "10.1",
-      "12.0+PTX",
-      "7.0",
-      "7.2",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.7",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch211-cxx11-cu126-aarch64-linux/__init__.py

@@ -1,75 +0,0 @@
-import torch
-
-from ._ops import ops
-
-from . import layers
-
-
-def silu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu_and_mul(out, x)
-    return out
-
-
-def mul_and_silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.mul_and_silu(out, x)
-    return out
-
-
-def gelu_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_and_mul(out, x)
-    return out
-
-
-def gelu_tanh_and_mul(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh_and_mul(out, x)
-    return out
-
-
-def fatrelu_and_mul(out: torch.Tensor, x: torch.Tensor, threshold: float = 0.0) -> None:
-    ops.fatrelu_and_mul(out, x, threshold)
-    return out
-
-
-def gelu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu(out, x)
-    return out
-
-def silu(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.silu(out, x)
-    return out
-
-
-def gelu_tanh(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_tanh(out, x)
-    return out
-
-
-def gelu_fast(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_fast(out, x)
-    return out
-
-
-def gelu_new(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_new(out, x)
-    return out
-
-
-def gelu_quick(out: torch.Tensor, x: torch.Tensor) -> None:
-    ops.gelu_quick(out, x)
-    return out
-
-
-__all__ = [
-    "silu_and_mul",
-    "mul_and_silu",
-    "gelu_and_mul",
-    "gelu_tanh_and_mul",
-    "fatrelu_and_mul",
-    "gelu_fast",
-    "gelu_new",
-    "gelu_quick",
-    "gelu_tanh",
-    "silu",
-    "gelu",
-    "layers",
-]