diff --git a/.github/workflows/build_kernel.yaml b/.github/workflows/build_kernel.yaml index 953cef32..eae52609 100644 --- a/.github/workflows/build_kernel.yaml +++ b/.github/workflows/build_kernel.yaml @@ -57,7 +57,6 @@ jobs: cpp20-symbols-kernel cutlass-gemm-kernel cutlass-gemm-tvm-ffi-kernel - extra-data relu-kernel relu-torch-stable-abi-kernel relu-tvm-ffi-kernel diff --git a/examples/kernels/extra-data/CARD.md b/examples/kernels/extra-data/CARD.md deleted file mode 100644 index 927ad93f..00000000 --- a/examples/kernels/extra-data/CARD.md +++ /dev/null @@ -1,66 +0,0 @@ ---- -library_name: kernels -{% if license %}license: {{ license }} -{% endif %}--- - -This is the repository card of {{ repo_id }} that has been pushed on the Hub. It was built to be used with the [`kernels` library](https://github.com/huggingface/kernels). This card was automatically generated. - -## How to use -{% if functions %} - -```python -# make sure `kernels` is installed: `pip install -U kernels` -from kernels import get_kernel - -# If the org / user isn't a trusted publisher, pass `trust_remote_code=True` to the -# `get_kernel` call. You can find whether this kernel is from a trusted publisher -# by going to the kernel's Hub page and finding the "Trusted publisher" status at -# the top of the page. -kernel_module = get_kernel("{{ repo_id }}", version={{ version }}) -{{ functions[0] }} = kernel_module.{{ functions[0] }} - -{{ functions[0] }}(...) -``` -{% else %} - -Usage example not available. -{% endif %} - -## Available functions -{% if functions %} -{% for func in functions %} -- `{{ func }}` -{% endfor %} -{% else %} - -Function list not available. -{% endif %} -{% if layers %} - -## Available layers -{% for layer in layers %} -- `{{ layer }}` -{% endfor %} -{% endif %} - -## Benchmarks -{% if has_benchmark %} - -Benchmarking script is available for this kernel. Run `kernels benchmark {{ repo_id }} --version {{ version }}`. -{% else %} - -No benchmark available yet. -{% endif %} -{% if upstream %} - -## Upstream - -The original source code for this kernel comes from {{ upstream }}. -{% endif %} -{% if source %} - -## Source - -The kernel-builder formatted source for this kernel is available at {{ source }}. -{% endif %} - diff --git a/examples/kernels/extra-data/build.toml b/examples/kernels/extra-data/build.toml deleted file mode 100644 index 18883c9b..00000000 --- a/examples/kernels/extra-data/build.toml +++ /dev/null @@ -1,65 +0,0 @@ -[general] -name = "extra-data" -version = 1 -edition = 5 -license = "Apache-2.0" -backends = [ - "cpu", - "cuda", - "metal", - "rocm", - "xpu", -] - -[general.hub] -repo-id = "kernels-test/extra-data" - -[torch] -pyext = [ - "json", - "py", -] -src = [ - "torch-ext/torch_binding.cpp", - "torch-ext/torch_binding.h", -] - -[kernel.relu_rocm] -backend = "rocm" -depends = ["torch"] -rocm-archs = [ - "gfx906", - "gfx908", - "gfx90a", - "gfx940", - "gfx941", - "gfx942", - "gfx1030", - "gfx1100", - "gfx1101", -] -src = ["relu_cuda/relu.cu"] - -[kernel.relu_xpu] -backend = "xpu" -depends = ["torch"] -src = ["relu_xpu/relu.cpp"] - -[kernel.relu_metal] -backend = "metal" -depends = ["torch"] -src = [ - "relu_metal/relu.mm", - "relu_metal/relu.metal", - "relu_metal/common.h", -] - -[kernel.relu] -backend = "cuda" -depends = ["torch"] -src = ["relu_cuda/relu.cu"] - -[kernel.relu_cpu] -backend = "cpu" -depends = ["torch"] -src = ["relu_cpu/relu_cpu.cpp"] diff --git a/examples/kernels/extra-data/flake.nix b/examples/kernels/extra-data/flake.nix deleted file mode 100644 index 41e51eef..00000000 --- a/examples/kernels/extra-data/flake.nix +++ /dev/null @@ -1,17 +0,0 @@ -{ - description = "Flake for ReLU kernel"; - - inputs = { - kernel-builder.url = "path:../../.."; - }; - - outputs = - { - self, - kernel-builder, - }: - kernel-builder.lib.genKernelFlakeOutputs { - inherit self; - path = ./.; - }; -} diff --git a/examples/kernels/extra-data/relu_cpu/relu_cpu.cpp b/examples/kernels/extra-data/relu_cpu/relu_cpu.cpp deleted file mode 100644 index 6197a9f1..00000000 --- a/examples/kernels/extra-data/relu_cpu/relu_cpu.cpp +++ /dev/null @@ -1,56 +0,0 @@ -#include - -#ifdef __SSE__ -#include -#endif - -#ifdef __ARM_NEON -#include -#endif - -#ifdef __SSE__ -void relu_forward_sse(float* out, const float* input, size_t size) { - size_t i = 0; - - for (; i + 4 <= size; i += 4) { - __m128 vec_input = _mm_load_ps(input + i); - __m128 vec_zero = _mm_setzero_ps(); - __m128 vec_output = _mm_max_ps(vec_input, vec_zero); - _mm_store_ps(out + i, vec_output); - } - - for (; i < size; ++i) { - out[i] = input[i] > 0 ? input[i] : 0; - } -} -#endif - -#ifdef __ARM_NEON -void relu_forward_neon(float* out, const float* input, size_t size) { - size_t i = 0; - - for (; i + 4 <= size; i += 4) { - float32x4_t vec_input = vld1q_f32(input + i); - float32x4_t vec_output = vmaxq_f32(vec_input, vdupq_n_f32(0)); - vst1q_f32(out + i, vec_output); - } - - for (; i < size; ++i) { - out[i] = input[i] > 0 ? input[i] : 0; - } -} -#endif - -void relu(torch::Tensor &out, torch::Tensor const &input) { - TORCH_CHECK(out.dtype() == torch::kFloat32, "Output tensor must be of dtype float"); - TORCH_CHECK(input.dtype() == torch::kFloat32, "Input tensor must be of dtype float"); - TORCH_CHECK(out.numel() == input.numel(), "Input and output tensors must have the same number of elements"); - -#if defined(__SSE__) - relu_forward_sse(out.data_ptr(), input.data_ptr(), input.numel()); -#elif defined(__ARM_NEON) - relu_forward_neon(out.data_ptr(), input.data_ptr(), input.numel()); -#else - #error "Unsupported architecture; please use a CPU with SSE or ARM NEON support." -#endif -} diff --git a/examples/kernels/extra-data/relu_cuda/relu.cu b/examples/kernels/extra-data/relu_cuda/relu.cu deleted file mode 100644 index 6bbe3160..00000000 --- a/examples/kernels/extra-data/relu_cuda/relu.cu +++ /dev/null @@ -1,43 +0,0 @@ -#include -#include -#include - -#include - -__global__ void relu_kernel(float *__restrict__ out, - float const *__restrict__ input, const int d) { - const int64_t token_idx = blockIdx.x; - for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) { - auto x = input[token_idx * d + idx]; - out[token_idx * d + idx] = x > 0.0f ? x : 0.0f; - } -} - -void relu(torch::Tensor &out, torch::Tensor const &input) { - TORCH_CHECK(input.device().is_cuda(), "input must be a CUDA tensor"); - TORCH_CHECK(input.is_contiguous(), "input must be contiguous"); - TORCH_CHECK(input.scalar_type() == at::ScalarType::Float && - input.scalar_type() == at::ScalarType::Float, - "relu_kernel only supports float32"); - - TORCH_CHECK(input.sizes() == out.sizes(), - "Tensors must have the same shape. Got input shape: ", - input.sizes(), " and output shape: ", out.sizes()); - - TORCH_CHECK(input.scalar_type() == out.scalar_type(), - "Tensors must have the same data type. Got input dtype: ", - input.scalar_type(), " and output dtype: ", out.scalar_type()); - - TORCH_CHECK(input.device() == out.device(), - "Tensors must be on the same device. Got input device: ", - input.device(), " and output device: ", out.device()); - - 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(); - relu_kernel<<>>(out.data_ptr(), - input.data_ptr(), d); -} diff --git a/examples/kernels/extra-data/relu_metal/common.h b/examples/kernels/extra-data/relu_metal/common.h deleted file mode 100644 index 1b891fad..00000000 --- a/examples/kernels/extra-data/relu_metal/common.h +++ /dev/null @@ -1,10 +0,0 @@ -#ifndef COMMON_H -#define COMMON_H - -#include -using namespace metal; - -// Common constants and utilities for Metal kernels -constant float RELU_THRESHOLD = 0.0f; - -#endif // COMMON_H \ No newline at end of file diff --git a/examples/kernels/extra-data/relu_metal/relu.metal b/examples/kernels/extra-data/relu_metal/relu.metal deleted file mode 100644 index 286b46fe..00000000 --- a/examples/kernels/extra-data/relu_metal/relu.metal +++ /dev/null @@ -1,17 +0,0 @@ -#include -#include "common.h" -using namespace metal; - -kernel void relu_forward_kernel_float(device const float *inA [[buffer(0)]], - device float *outC [[buffer(1)]], - uint index [[thread_position_in_grid]]) { - // Explicitly write to output - outC[index] = max(RELU_THRESHOLD, inA[index]); -} - -kernel void relu_forward_kernel_half(device const half *inA [[buffer(0)]], - device half *outC [[buffer(1)]], - uint index [[thread_position_in_grid]]) { - // Explicitly write to output - outC[index] = max(static_cast(0.0), inA[index]); -} \ No newline at end of file diff --git a/examples/kernels/extra-data/relu_metal/relu.mm b/examples/kernels/extra-data/relu_metal/relu.mm deleted file mode 100644 index 7636737b..00000000 --- a/examples/kernels/extra-data/relu_metal/relu.mm +++ /dev/null @@ -1,105 +0,0 @@ -#include - -#import -#import - -// Include the auto-generated header with embedded metallib -#ifdef EMBEDDED_METALLIB_HEADER -#include EMBEDDED_METALLIB_HEADER -#else -#error "EMBEDDED_METALLIB_HEADER not defined" -#endif - -static inline id getMTLBufferStorage(const torch::Tensor &tensor) { - return __builtin_bit_cast(id, tensor.storage().data()); -} - - -torch::Tensor &dispatchReluKernel(torch::Tensor const &input, - torch::Tensor &output) { - @autoreleasepool { - id device = MTLCreateSystemDefaultDevice(); - - int numThreads = input.numel(); - - // Load the embedded Metal library from memory - NSError *error = nil; - id customKernelLibrary = EMBEDDED_METALLIB_NAMESPACE::createLibrary(device, &error); - TORCH_CHECK(customKernelLibrary, - "Failed to create Metal library from embedded data: ", - error.localizedDescription.UTF8String); - - std::string kernel_name = - std::string("relu_forward_kernel_") + - (input.scalar_type() == torch::kFloat ? "float" : "half"); - id customReluFunction = [customKernelLibrary - newFunctionWithName:[NSString - stringWithUTF8String:kernel_name.c_str()]]; - TORCH_CHECK(customReluFunction, - "Failed to create function state object for ", - kernel_name.c_str()); - - id reluPSO = - [device newComputePipelineStateWithFunction:customReluFunction - error:&error]; - TORCH_CHECK(reluPSO, error.localizedDescription.UTF8String); - - id commandBuffer = torch::mps::get_command_buffer(); - TORCH_CHECK(commandBuffer, "Failed to retrieve command buffer reference"); - - dispatch_queue_t serialQueue = torch::mps::get_dispatch_queue(); - - dispatch_sync(serialQueue, ^() { - id computeEncoder = - [commandBuffer computeCommandEncoder]; - TORCH_CHECK(computeEncoder, "Failed to create compute command encoder"); - - [computeEncoder setComputePipelineState:reluPSO]; - [computeEncoder setBuffer:getMTLBufferStorage(input) - offset:input.storage_offset() * input.element_size() - atIndex:0]; - [computeEncoder setBuffer:getMTLBufferStorage(output) - offset:output.storage_offset() * output.element_size() - atIndex:1]; - - MTLSize gridSize = MTLSizeMake(numThreads, 1, 1); - - NSUInteger threadGroupSize = reluPSO.maxTotalThreadsPerThreadgroup; - if (threadGroupSize > numThreads) { - threadGroupSize = numThreads; - } - MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1); - - [computeEncoder dispatchThreads:gridSize - threadsPerThreadgroup:threadgroupSize]; - - [computeEncoder endEncoding]; - - torch::mps::commit(); - }); - } - - return output; -} - -void relu(torch::Tensor &out, torch::Tensor const &input) { - TORCH_CHECK(input.device().is_mps(), "input must be a MPS tensor"); - TORCH_CHECK(input.is_contiguous(), "input must be contiguous"); - TORCH_CHECK(input.scalar_type() == torch::kFloat || - input.scalar_type() == torch::kHalf, - "Unsupported data type: ", input.scalar_type()); - - TORCH_CHECK(input.sizes() == out.sizes(), - "Tensors must have the same shape. Got input shape: ", - input.sizes(), " and output shape: ", out.sizes()); - - TORCH_CHECK(input.scalar_type() == out.scalar_type(), - "Tensors must have the same data type. Got input dtype: ", - input.scalar_type(), " and output dtype: ", out.scalar_type()); - - TORCH_CHECK(input.device() == out.device(), - "Tensors must be on the same device. Got input device: ", - input.device(), " and output device: ", out.device()); - - dispatchReluKernel(input, out); -} diff --git a/examples/kernels/extra-data/relu_xpu/relu.cpp b/examples/kernels/extra-data/relu_xpu/relu.cpp deleted file mode 100644 index 1809de08..00000000 --- a/examples/kernels/extra-data/relu_xpu/relu.cpp +++ /dev/null @@ -1,40 +0,0 @@ -#include -#include - -using namespace sycl; - -void relu_xpu_impl(torch::Tensor& output, const torch::Tensor& input) { - // Create SYCL queue directly - sycl::queue queue; - - auto input_ptr = input.data_ptr(); - auto output_ptr = output.data_ptr(); - auto numel = input.numel(); - - // Launch SYCL kernel - queue.parallel_for(range<1>(numel), [=](id<1> idx) { - auto i = idx[0]; - output_ptr[i] = input_ptr[i] > 0.0f ? input_ptr[i] : 0.0f; - }).wait(); -} - -void relu(torch::Tensor& out, const torch::Tensor& input) { - TORCH_CHECK(input.device().is_xpu(), "input must be a XPU tensor"); - TORCH_CHECK(input.is_contiguous(), "input must be contiguous"); - TORCH_CHECK(input.scalar_type() == torch::kFloat, - "Unsupported data type: ", input.scalar_type()); - - TORCH_CHECK(input.sizes() == out.sizes(), - "Tensors must have the same shape. Got input shape: ", - input.sizes(), " and output shape: ", out.sizes()); - - TORCH_CHECK(input.scalar_type() == out.scalar_type(), - "Tensors must have the same data type. Got input dtype: ", - input.scalar_type(), " and output dtype: ", out.scalar_type()); - - TORCH_CHECK(input.device() == out.device(), - "Tensors must be on the same device. Got input device: ", - input.device(), " and output device: ", out.device()); - - relu_xpu_impl(out, input); -} diff --git a/examples/kernels/extra-data/tests/__init__.py b/examples/kernels/extra-data/tests/__init__.py deleted file mode 100644 index e69de29b..00000000 diff --git a/examples/kernels/extra-data/tests/test_relu.py b/examples/kernels/extra-data/tests/test_relu.py deleted file mode 100644 index 3cf14496..00000000 --- a/examples/kernels/extra-data/tests/test_relu.py +++ /dev/null @@ -1,42 +0,0 @@ -import platform - -import kernels -import pytest -import torch -import torch.nn.functional as F - -extra_data = kernels.get_kernel("kernels-test/extra-data", version=1) - - -@pytest.mark.kernels_ci -def test_relu(): - if platform.system() == "Darwin": - device = torch.device("mps") - elif hasattr(torch, "xpu") and torch.xpu.is_available(): - device = torch.device("xpu") - elif torch.version.cuda is not None and torch.cuda.is_available(): - device = torch.device("cuda") - else: - device = torch.device("cpu") - x = torch.randn(1024, 1024, dtype=torch.float32, device=device) - torch.testing.assert_allclose(F.relu(x), extra_data.relu(x)) - - -@pytest.mark.kernels_ci -def test_relu_layer(): - if platform.system() == "Darwin": - device = torch.device("mps") - elif hasattr(torch, "xpu") and torch.xpu.is_available(): - device = torch.device("xpu") - elif torch.version.cuda is not None and torch.cuda.is_available(): - device = torch.device("cuda") - else: - device = torch.device("cpu") - x = torch.randn(1024, 1024, dtype=torch.float32, device=device) - layer = extra_data.layers.ReLU() - torch.testing.assert_allclose(F.relu(x), layer(x)) - - -@pytest.mark.kernels_ci -def test_data(): - assert extra_data.EASTER_EGG == 42 diff --git a/examples/kernels/extra-data/torch-ext/extra_data/__init__.py b/examples/kernels/extra-data/torch-ext/extra_data/__init__.py deleted file mode 100644 index 684f3558..00000000 --- a/examples/kernels/extra-data/torch-ext/extra_data/__init__.py +++ /dev/null @@ -1,32 +0,0 @@ -import json -from pathlib import Path -from typing import Optional - -import torch - -from ._ops import ops - -from . import layers - - -# This is the regular ReLU, but this example also shows how to embed some -# non-Python data. This can be used for e.g. Triton tuning data. - - -def _read_json() -> dict: - json_path = Path(__file__).parent / "data.json" - with open(json_path, "r") as f: - return json.load(f) - - -EASTER_EGG = _read_json() - - -def relu(x: torch.Tensor, out: Optional[torch.Tensor] = None) -> torch.Tensor: - if out is None: - out = torch.empty_like(x) - ops.relu(out, x) - return out - - -__all__ = ["EASTER_EGG", "relu", "layers"] diff --git a/examples/kernels/extra-data/torch-ext/extra_data/data.json b/examples/kernels/extra-data/torch-ext/extra_data/data.json deleted file mode 100644 index d81cc071..00000000 --- a/examples/kernels/extra-data/torch-ext/extra_data/data.json +++ /dev/null @@ -1 +0,0 @@ -42 diff --git a/examples/kernels/extra-data/torch-ext/extra_data/layers/__init__.py b/examples/kernels/extra-data/torch-ext/extra_data/layers/__init__.py deleted file mode 100644 index 6105a191..00000000 --- a/examples/kernels/extra-data/torch-ext/extra_data/layers/__init__.py +++ /dev/null @@ -1,11 +0,0 @@ -import torch -import torch.nn as nn - -from .._ops import ops - - -class ReLU(nn.Module): - def forward(self, x: torch.Tensor) -> torch.Tensor: - out = torch.empty_like(x) - ops.relu(out, x) - return out diff --git a/examples/kernels/extra-data/torch-ext/torch_binding.cpp b/examples/kernels/extra-data/torch-ext/torch_binding.cpp deleted file mode 100644 index 1765d92d..00000000 --- a/examples/kernels/extra-data/torch-ext/torch_binding.cpp +++ /dev/null @@ -1,19 +0,0 @@ -#include - -#include "registration.h" -#include "torch_binding.h" - -TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) { - ops.def("relu(Tensor! out, Tensor input) -> ()"); -#if defined(CPU_KERNEL) - ops.impl("relu", torch::kCPU, &relu); -#elif defined(CUDA_KERNEL) || defined(ROCM_KERNEL) - ops.impl("relu", torch::kCUDA, &relu); -#elif defined(METAL_KERNEL) - ops.impl("relu", torch::kMPS, relu); -#elif defined(XPU_KERNEL) - ops.impl("relu", torch::kXPU, &relu); -#endif -} - -REGISTER_EXTENSION(TORCH_EXTENSION_NAME) diff --git a/examples/kernels/extra-data/torch-ext/torch_binding.h b/examples/kernels/extra-data/torch-ext/torch_binding.h deleted file mode 100644 index 3bcf2904..00000000 --- a/examples/kernels/extra-data/torch-ext/torch_binding.h +++ /dev/null @@ -1,5 +0,0 @@ -#pragma once - -#include - -void relu(torch::Tensor &out, torch::Tensor const &input); \ No newline at end of file diff --git a/examples/kernels/flake.nix b/examples/kernels/flake.nix index 6a1cf836..8c916408 100644 --- a/examples/kernels/flake.nix +++ b/examples/kernels/flake.nix @@ -107,12 +107,6 @@ drv = sys: out: out.packages.${sys}.redistributable.${"tvm-ffi${tvmFfiVersion}-${cudaVersion}-${sys}"}; } - { - name = "extra-data"; - path = ./extra-data; - drv = - sys: out: out.packages.${sys}.redistributable.${"torch${torchVersion}-cxx11-${cudaVersion}-${sys}"}; - } { name = "relu-kernel-cpu"; path = ./relu; diff --git a/nix-builder/tests/Dockerfile.test-kernel b/nix-builder/tests/Dockerfile.test-kernel index 7329b5cb..371907cd 100644 --- a/nix-builder/tests/Dockerfile.test-kernel +++ b/nix-builder/tests/Dockerfile.test-kernel @@ -79,9 +79,7 @@ COPY cutlass-gemm-tvm-ffi-kernel ./cutlass-gemm-tvm-ffi-kernel COPY relu-triton-kernel ./relu-triton-kernel COPY gemm-triton-autotune-kernel ./gemm-triton-autotune-kernel COPY silu-and-mul-kernel ./silu-and-mul-kernel -COPY extra-data ./extra-data COPY cpp20-symbols-kernel ./cpp20-symbols-kernel -COPY examples/kernels/extra-data/tests ./extra_data_tests COPY examples/kernels/relu/tests ./relu_tests COPY examples/kernels/relu-tvm-ffi/tests ./relu_tvm_ffi_tests COPY examples/kernels/cutlass-gemm/tests ./cutlass_gemm_tests diff --git a/nix-builder/tests/run-tests.sh b/nix-builder/tests/run-tests.sh index 7374aba1..eb3a98b9 100644 --- a/nix-builder/tests/run-tests.sh +++ b/nix-builder/tests/run-tests.sh @@ -2,7 +2,6 @@ set -euo pipefail # Expand to build variant directories. -EXTRA_DATA_PATH=$(echo extra-data/torch*) RELU_PATH=$(echo relu-kernel/torch*) RELU_TORCH_STABLE_ABI_PATH=$(echo relu-torch-stable-abi-kernel/torch*) RELU_TVM_FFI_PATH=$(echo relu-tvm-ffi-kernel/tvm-ffi*) @@ -14,8 +13,8 @@ SILU_MUL_PATH=$(echo silu-and-mul-kernel/torch*) RELU_CPU_PATH=$(echo relu-kernel-cpu/torch*) CPP20_SYMBOLS_PATH=$(echo cpp20-symbols-kernel/torch*) -LOCAL_KERNELS="kernels-test/extra-data=${EXTRA_DATA_PATH}:kernels-test/relu=${RELU_PATH}:kernels-test/relu-torch-stable-abi=${RELU_TORCH_STABLE_ABI_PATH}:kernels-test/relu-tvm-ffi=${RELU_TVM_FFI_PATH}:kernels-test/cutlass-gemm=${CUTLASS_PATH}:kernels-test/cutlass-gemm-tvm-ffi=${CUTLASS_TVM_FFI_PATH}" \ - .venv/bin/pytest extra_data_tests relu_tests relu_tvm_ffi_tests cutlass_gemm_tests cutlass_gemm_tvm_ffi_tests +LOCAL_KERNELS="kernels-test/relu=${RELU_PATH}:kernels-test/relu-torch-stable-abi=${RELU_TORCH_STABLE_ABI_PATH}:kernels-test/relu-tvm-ffi=${RELU_TVM_FFI_PATH}:kernels-test/cutlass-gemm=${CUTLASS_PATH}:kernels-test/cutlass-gemm-tvm-ffi=${CUTLASS_TVM_FFI_PATH}" \ + .venv/bin/pytest relu_tests relu_tvm_ffi_tests cutlass_gemm_tests cutlass_gemm_tvm_ffi_tests LOCAL_KERNELS="kernels-test/relu-triton=${RELU_TRITON_PATH}" \ .venv/bin/pytest relu_triton_tests