diff --git a/cuda/include/kernels.h b/cuda/include/kernels.h
index 28e14a3..09c066b 100644
--- a/cuda/include/kernels.h
+++ b/cuda/include/kernels.h
@@ -36,7 +36,6 @@ int fastnn_cuda_clamp(const float* a, float min_val, float max_val, float* out,
 int fastnn_cuda_relu(const float* input, float* output, size_t n);
 int fastnn_cuda_relu_backward(const float* grad_output, const float* input, float* grad_input, size_t n);
 int fastnn_cuda_leaky_relu(const float* input, float negative_slope, float* output, size_t n);
-int fastnn_cuda_leaky_relu_backward(const float* grad_output, const float* input, float negative_slope, float* grad_input, size_t n);
 int fastnn_cuda_sigmoid(const float* input, float* output, size_t n);
 int fastnn_cuda_sigmoid_backward(const float* grad_output, const float* output, float* grad_input, size_t n);
 int fastnn_cuda_tanh_forward(const float* input, float* output, size_t n);
@@ -70,73 +69,6 @@ int fastnn_cuda_sum_axis(const float* input, float* output, int* shape, int ndim
 int fastnn_cuda_mean(const float* input, float* output, size_t n);
 int fastnn_cuda_max(const float* input, float* output, size_t n);
 int fastnn_cuda_min(const float* input, float* output, size_t n);
-int fastnn_cuda_argmax(const float* input, int* output, size_t n);
-int fastnn_cuda_argmax_axis(const float* input, int* output, int outer, int axis_size, int inner);
-
-int fastnn_cuda_conv2d_forward(
-    const float* input, const float* weight, const float* bias, float* output,
-    int batch_size, int in_channels, int out_channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w);
-
-int fastnn_cuda_conv2d_backward_data(
-    const float* grad_output, const float* weight, float* grad_input,
-    int batch_size, int in_channels, int out_channels,
-    int input_h, int input_w, int output_h, int output_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w);
-
-int fastnn_cuda_conv2d_backward_weight(
-    const float* input, const float* grad_output, float* grad_weight,
-    int batch_size, int in_channels, int out_channels,
-    int input_h, int input_w, int output_h, int output_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w);
-
-int fastnn_cuda_max_pool2d(
-    const float* input, float* output, int* indices,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w);
-
-int fastnn_cuda_max_pool2d_backward(
-    const float* grad_output, const int* indices, float* grad_input,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int output_h, int output_w);
-
-int fastnn_cuda_avg_pool2d(
-    const float* input, float* output,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w);
-
-int fastnn_cuda_adaptive_avg_pool2d(
-    const float* input, float* output,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int output_h, int output_w);
-
-int fastnn_cuda_batch_norm_forward(
-    const float* input, const float* gamma, const float* beta,
-    float* running_mean, float* running_var,
-    float* output, float* save_mean, float* save_inv_var,
-    int batch_size, int channels, int spatial_size,
-    float momentum, float epsilon, int training);
-
-int fastnn_cuda_batch_norm_backward(
-    const float* grad_output, const float* input,
-    const float* gamma, const float* save_mean, const float* save_inv_var,
-    float* grad_input, float* grad_gamma, float* grad_beta,
-    int batch_size, int channels, int spatial_size);
 
 int fastnn_cuda_layer_norm_forward(
     const float* input, const float* gamma, const float* beta,
@@ -149,45 +81,9 @@ int fastnn_cuda_layer_norm_backward(
     float* grad_input, float* grad_gamma, float* grad_beta,
     int batch_size, int normalized_size);
 
-int fastnn_cuda_rms_norm_forward(
-    const float* input, const float* gamma,
-    float* output, float* rms,
-    int batch_size, int normalized_size, float epsilon);
-
-int fastnn_cuda_dropout_forward(const float* input, float* output, float* mask, float p, size_t n, unsigned long long seed);
-int fastnn_cuda_dropout_backward(const float* grad_output, const float* mask, float* grad_input, float p, size_t n);
-
 int fastnn_cuda_embedding_forward(const int* indices, const float* weight, float* output, int num_indices, int embedding_dim);
 int fastnn_cuda_embedding_backward(const int* indices, const float* grad_output, float* grad_weight, int num_indices, int embedding_dim, int num_embeddings);
 
-int fastnn_cuda_cross_entropy_loss(const float* log_probs, const int* targets, float* loss, float* grad, int batch_size, int num_classes);
-int fastnn_cuda_mse_loss(const float* predictions, const float* targets, float* loss, float* grad, size_t n);
-int fastnn_cuda_binary_cross_entropy(const float* predictions, const float* targets, float* loss, float* grad, size_t n);
-
-int fastnn_cuda_scaled_dot_product_attention(
-    const float* query, const float* key, const float* value,
-    float* output, float* attn_weights,
-    const float* mask,
-    int batch_size, int num_heads, int seq_len_q, int seq_len_k, int head_dim,
-    float scale, int causal);
-
-int fastnn_cuda_sgd_step(float* params, const float* grads, float* velocity,
-                          float lr, float momentum, float weight_decay, float dampening,
-                          int nesterov, size_t n);
-
-int fastnn_cuda_adam_step(float* params, const float* grads,
-                           float* m, float* v,
-                           float lr, float beta1, float beta2, float epsilon,
-                           float weight_decay, int step, int amsgrad, float* v_max,
-                           size_t n);
-
-int fastnn_cuda_fill(float* data, float value, size_t n);
-int fastnn_cuda_copy(const float* src, float* dst, size_t n);
-int fastnn_cuda_arange(float* output, float start, float step, size_t n);
-int fastnn_cuda_where(const float* condition, const float* x, const float* y, float* output, size_t n);
-int fastnn_cuda_gather(const float* input, const int* indices, float* output,
-                        int outer_size, int gather_dim_size, int inner_size, int num_indices);
-
 #ifdef __cplusplus
 }
 #endif
diff --git a/cuda/kernels.cu b/cuda/kernels.cu
index 497bc86..5cc64c7 100644
--- a/cuda/kernels.cu
+++ b/cuda/kernels.cu
@@ -10,7 +10,6 @@
 #define BLOCK_SIZE 256
 #define TILE_SIZE 32
 #define WARP_SIZE 32
-#define ARGMAX_MAX_THREADS 1024
 #define PERMUTE_MAX_DIMS 8
 
 #define CUDA_CHECK(call) do { \
@@ -203,11 +202,6 @@ __global__ void kernel_leaky_relu(const float* input, float neg_slope, float* ou
     }
 }
 
-__global__ void kernel_leaky_relu_backward(const float* grad_output, const float* input, float neg_slope, float* grad_input, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx < n) grad_input[idx] = input[idx] > 0.0f ? grad_output[idx] : neg_slope * grad_output[idx];
-}
-
 __global__ void kernel_sigmoid(const float* input, float* output, size_t n) {
     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
     if (idx < n) output[idx] = 1.0f / (1.0f + expf(-input[idx]));
@@ -267,7 +261,6 @@ __global__ void kernel_silu_backward(const float* grad_output, const float* inpu
 extern "C" int fastnn_cuda_relu(const float* in, float* out, size_t n) { LAUNCH_ELEMENTWISE(kernel_relu, in, out, n); }
 extern "C" int fastnn_cuda_relu_backward(const float* go, const float* in, float* gi, size_t n) { LAUNCH_ELEMENTWISE(kernel_relu_backward, go, in, gi, n); }
 extern "C" int fastnn_cuda_leaky_relu(const float* in, float ns, float* out, size_t n) { LAUNCH_ELEMENTWISE(kernel_leaky_relu, in, ns, out, n); }
-extern "C" int fastnn_cuda_leaky_relu_backward(const float* go, const float* in, float ns, float* gi, size_t n) { LAUNCH_ELEMENTWISE(kernel_leaky_relu_backward, go, in, ns, gi, n); }
 extern "C" int fastnn_cuda_sigmoid(const float* in, float* out, size_t n) { LAUNCH_ELEMENTWISE(kernel_sigmoid, in, out, n); }
 extern "C" int fastnn_cuda_sigmoid_backward(const float* go, const float* out, float* gi, size_t n) { LAUNCH_ELEMENTWISE(kernel_sigmoid_backward, go, out, gi, n); }
 extern "C" int fastnn_cuda_tanh_forward(const float* in, float* out, size_t n) { LAUNCH_ELEMENTWISE(kernel_tanh_forward, in, out, n); }
@@ -726,67 +719,6 @@ extern "C" int fastnn_cuda_min(const float* input, float* output, size_t n) {
     return 0;
 }
 
-__global__ void kernel_argmax(const float* input, int* output, size_t n) {
-    extern __shared__ char shared_mem[];
-    float* sval = (float*)shared_mem;
-    int* sidx = (int*)(shared_mem + blockDim.x * sizeof(float));
-
-    unsigned int tid = threadIdx.x;
-    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-
-    sval[tid] = (i < n) ? input[i] : -FLT_MAX;
-    sidx[tid] = (i < n) ? (int)i : 0;
-    __syncthreads();
-
-    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
-        if (tid < s && sval[tid + s] > sval[tid]) {
-            sval[tid] = sval[tid + s];
-            sidx[tid] = sidx[tid + s];
-        }
-        __syncthreads();
-    }
-
-    if (tid == 0) {
-        output[0] = sidx[0];
-    }
-}
-
-extern "C" int fastnn_cuda_argmax(const float* input, int* output, size_t n) {
-    // One block, one element per thread: only right for n <= ARGMAX_MAX_THREADS
-    // and a power-of-two n.
-    kernel_argmax<<<1, min((int)n, ARGMAX_MAX_THREADS), ARGMAX_MAX_THREADS * (sizeof(float) + sizeof(int))>>>(input, output, n);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_argmax_axis(const float* input, int* output, int outer, int axis_size, int inner) {
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int total = outer * inner;
-    if (idx >= total) return;
-
-    int o = idx / inner;
-    int i = idx % inner;
-
-    float max_val = -FLT_MAX;
-    int max_idx = 0;
-    for (int a = 0; a < axis_size; a++) {
-        float val = input[(o * axis_size + a) * inner + i];
-        if (val > max_val) {
-            max_val = val;
-            max_idx = a;
-        }
-    }
-    output[idx] = max_idx;
-}
-
-extern "C" int fastnn_cuda_argmax_axis(const float* input, int* output, int outer, int axis_size, int inner) {
-    int total = outer * inner;
-    int blocks = div_ceil(total, BLOCK_SIZE);
-    kernel_argmax_axis<<<blocks, BLOCK_SIZE>>>(input, output, outer, axis_size, inner);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
 __global__ void kernel_im2col(
     const float* input, float* col,
     int channels, int height, int width,
@@ -847,477 +779,6 @@ __global__ void kernel_col2im(
     input[idx] = sum;
 }
 
-extern "C" int fastnn_cuda_conv2d_forward(
-    const float* input, const float* weight, const float* bias, float* output,
-    int batch_size, int in_channels, int out_channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w)
-{
-    int out_h = (input_h + 2 * pad_h - kernel_h) / stride_h + 1;
-    int out_w = (input_w + 2 * pad_w - kernel_w) / stride_w + 1;
-
-    int col_size = in_channels * kernel_h * kernel_w * out_h * out_w;
-    float* col;
-    CUDA_CHECK(cudaMalloc(&col, col_size * sizeof(float)));
-
-    for (int b = 0; b < batch_size; b++) {
-        const float* input_b = input + b * in_channels * input_h * input_w;
-        float* output_b = output + b * out_channels * out_h * out_w;
-
-        int blocks = div_ceil(col_size, BLOCK_SIZE);
-        kernel_im2col<<<blocks, BLOCK_SIZE>>>(
-            input_b, col, in_channels, input_h, input_w,
-            kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, out_h, out_w);
-
-        float alpha = 1.0f, beta = 0.0f;
-        int M = out_channels;
-        int N = out_h * out_w;
-        int K = in_channels * kernel_h * kernel_w;
-
-        CUBLAS_CHECK(cublasSgemm(g_cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N,
-                                  N, M, K, &alpha, col, N, weight, K, &beta, output_b, N));
-
-        if (bias != nullptr) {
-            for (int oc = 0; oc < out_channels; oc++) {
-                float bias_val;
-                CUDA_CHECK(cudaMemcpy(&bias_val, bias + oc, sizeof(float), cudaMemcpyDeviceToHost));
-                kernel_add_scalar<<<div_ceil(N, BLOCK_SIZE), BLOCK_SIZE>>>(
-                    output_b + oc * N, bias_val, output_b + oc * N, N);
-            }
-        }
-    }
-
-    cudaFree(col);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-extern "C" int fastnn_cuda_conv2d_backward_data(
-    const float* grad_output, const float* weight, float* grad_input,
-    int batch_size, int in_channels, int out_channels,
-    int input_h, int input_w, int output_h, int output_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w)
-{
-    int col_size = in_channels * kernel_h * kernel_w * output_h * output_w;
-    float* col;
-    CUDA_CHECK(cudaMalloc(&col, col_size * sizeof(float)));
-
-    CUDA_CHECK(cudaMemset(grad_input, 0, batch_size * in_channels * input_h * input_w * sizeof(float)));
-
-    for (int b = 0; b < batch_size; b++) {
-        const float* go_b = grad_output + b * out_channels * output_h * output_w;
-        float* gi_b = grad_input + b * in_channels * input_h * input_w;
-
-        float alpha = 1.0f, beta = 0.0f;
-        int M = in_channels * kernel_h * kernel_w;
-        int N = output_h * output_w;
-        int K = out_channels;
-
-        CUBLAS_CHECK(cublasSgemm(g_cublas_handle, CUBLAS_OP_N, CUBLAS_OP_T,
-                                  N, M, K, &alpha, go_b, N, weight, M, &beta, col, N));
-
-        int total = in_channels * input_h * input_w;
-        int blocks = div_ceil(total, BLOCK_SIZE);
-        kernel_col2im<<<blocks, BLOCK_SIZE>>>(
-            col, gi_b, in_channels, input_h, input_w,
-            kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, output_h, output_w);
-    }
-
-    cudaFree(col);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-extern "C" int fastnn_cuda_conv2d_backward_weight(
-    const float* input, const float* grad_output, float* grad_weight,
-    int batch_size, int in_channels, int out_channels,
-    int input_h, int input_w, int output_h, int output_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w)
-{
-    int col_size = in_channels * kernel_h * kernel_w * output_h * output_w;
-    float* col;
-    CUDA_CHECK(cudaMalloc(&col, col_size * sizeof(float)));
-
-    CUDA_CHECK(cudaMemset(grad_weight, 0, out_channels * in_channels * kernel_h * kernel_w * sizeof(float)));
-
-    for (int b = 0; b < batch_size; b++) {
-        const float* input_b = input + b * in_channels * input_h * input_w;
-        const float* go_b = grad_output + b * out_channels * output_h * output_w;
-
-        int blocks = div_ceil(col_size, BLOCK_SIZE);
-        kernel_im2col<<<blocks, BLOCK_SIZE>>>(
-            input_b, col, in_channels, input_h, input_w,
-            kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, output_h, output_w);
-
-        float alpha = 1.0f, beta = 1.0f;
-        int M = out_channels;
-        int N = in_channels * kernel_h * kernel_w;
-        int K = output_h * output_w;
-
-        CUBLAS_CHECK(cublasSgemm(g_cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N,
-                                  N, M, K, &alpha, col, K, go_b, K, &beta, grad_weight, N));
-    }
-
-    cudaFree(col);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_max_pool2d(
-    const float* input, float* output, int* indices,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w,
-    int out_h, int out_w)
-{
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int total = batch_size * channels * out_h * out_w;
-    if (idx >= total) return;
-
-    int ow = idx % out_w;
-    int oh = (idx / out_w) % out_h;
-    int c = (idx / out_w / out_h) % channels;
-    int b = idx / out_w / out_h / channels;
-
-    const float* input_bc = input + (b * channels + c) * input_h * input_w;
-
-    float max_val = -FLT_MAX;
-    int max_idx = 0;
-
-    for (int kh = 0; kh < kernel_h; kh++) {
-        for (int kw = 0; kw < kernel_w; kw++) {
-            int ih = oh * stride_h - pad_h + kh;
-            int iw = ow * stride_w - pad_w + kw;
-            if (ih >= 0 && ih < input_h && iw >= 0 && iw < input_w) {
-                int input_idx = ih * input_w + iw;
-                float val = input_bc[input_idx];
-                if (val > max_val) {
-                    max_val = val;
-                    max_idx = input_idx;
-                }
-            }
-        }
-    }
-
-    output[idx] = max_val;
-    if (indices != nullptr) indices[idx] = max_idx;
-}
-
-extern "C" int fastnn_cuda_max_pool2d(
-    const float* input, float* output, int* indices,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w)
-{
-    int out_h = (input_h + 2 * pad_h - kernel_h) / stride_h + 1;
-    int out_w = (input_w + 2 * pad_w - kernel_w) / stride_w + 1;
-    int total = batch_size * channels * out_h * out_w;
-    int blocks = div_ceil(total, BLOCK_SIZE);
-    kernel_max_pool2d<<<blocks, BLOCK_SIZE>>>(
-        input, output, indices, batch_size, channels,
-        input_h, input_w, kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, out_h, out_w);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_max_pool2d_backward(
-    const float* grad_output, const int* indices, float* grad_input,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int output_h, int output_w)
-{
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int total = batch_size * channels * output_h * output_w;
-    if (idx >= total) return;
-
-    int c = (idx / output_h / output_w) % channels;
-    int b = idx / output_h / output_w / channels;
-
-    int input_offset = (b * channels + c) * input_h * input_w;
-    atomicAdd(&grad_input[input_offset + indices[idx]], grad_output[idx]);
-}
-
-extern "C" int fastnn_cuda_max_pool2d_backward(
-    const float* grad_output, const int* indices, float* grad_input,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int output_h, int output_w)
-{
-    CUDA_CHECK(cudaMemset(grad_input, 0, batch_size * channels * input_h * input_w * sizeof(float)));
-    int total = batch_size * channels * output_h * output_w;
-    int blocks = div_ceil(total, BLOCK_SIZE);
-    kernel_max_pool2d_backward<<<blocks, BLOCK_SIZE>>>(
-        grad_output, indices, grad_input, batch_size, channels, input_h, input_w, output_h, output_w);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_avg_pool2d(
-    const float* input, float* output,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w,
-    int out_h, int out_w)
-{
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int total = batch_size * channels * out_h * out_w;
-    if (idx >= total) return;
-
-    int ow = idx % out_w;
-    int oh = (idx / out_w) % out_h;
-    int c = (idx / out_w / out_h) % channels;
-    int b = idx / out_w / out_h / channels;
-
-    const float* input_bc = input + (b * channels + c) * input_h * input_w;
-
-    float sum = 0.0f;
-    int count = 0;
-    for (int kh = 0; kh < kernel_h; kh++) {
-        for (int kw = 0; kw < kernel_w; kw++) {
-            int ih = oh * stride_h - pad_h + kh;
-            int iw = ow * stride_w - pad_w + kw;
-            if (ih >= 0 && ih < input_h && iw >= 0 && iw < input_w) {
-                sum += input_bc[ih * input_w + iw];
-                count++;
-            }
-        }
-    }
-    output[idx] = sum / (float)count;
-}
-
-extern "C" int fastnn_cuda_avg_pool2d(
-    const float* input, float* output,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int kernel_h, int kernel_w,
-    int stride_h, int stride_w,
-    int pad_h, int pad_w)
-{
-    int out_h = (input_h + 2 * pad_h - kernel_h) / stride_h + 1;
-    int out_w = (input_w + 2 * pad_w - kernel_w) / stride_w + 1;
-    int total = batch_size * channels * out_h * out_w;
-    int blocks = div_ceil(total, BLOCK_SIZE);
-    kernel_avg_pool2d<<<blocks, BLOCK_SIZE>>>(
-        input, output, batch_size, channels, input_h, input_w,
-        kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, out_h, out_w);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_adaptive_avg_pool2d(
-    const float* input, float* output,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int output_h, int output_w)
-{
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int total = batch_size * channels * output_h * output_w;
-    if (idx >= total) return;
-
-    int ow = idx % output_w;
-    int oh = (idx / output_w) % output_h;
-    int c = (idx / output_w / output_h) % channels;
-    int b = idx / output_w / output_h / channels;
-
-    int ih_start = (oh * input_h) / output_h;
-    int ih_end = ((oh + 1) * input_h + output_h - 1) / output_h;
-    int iw_start = (ow * input_w) / output_w;
-    int iw_end = ((ow + 1) * input_w + output_w - 1) / output_w;
-
-    const float* input_bc = input + (b * channels + c) * input_h * input_w;
-    float sum = 0.0f;
-    int count = 0;
-    for (int ih = ih_start; ih < ih_end; ih++) {
-        for (int iw = iw_start; iw < iw_end; iw++) {
-            sum += input_bc[ih * input_w + iw];
-            count++;
-        }
-    }
-    output[idx] = sum / (float)count;
-}
-
-extern "C" int fastnn_cuda_adaptive_avg_pool2d(
-    const float* input, float* output,
-    int batch_size, int channels,
-    int input_h, int input_w,
-    int output_h, int output_w)
-{
-    int total = batch_size * channels * output_h * output_w;
-    int blocks = div_ceil(total, BLOCK_SIZE);
-    kernel_adaptive_avg_pool2d<<<blocks, BLOCK_SIZE>>>(
-        input, output, batch_size, channels, input_h, input_w, output_h, output_w);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_batch_norm_forward(
-    const float* input, const float* gamma, const float* beta,
-    float* running_mean, float* running_var,
-    float* output, float* save_mean, float* save_inv_var,
-    int batch_size, int channels, int spatial_size,
-    float momentum, float epsilon, int training)
-{
-    int c = blockIdx.x;
-    if (c >= channels) return;
-
-    int n = batch_size * spatial_size;
-
-    if (training) {
-        float mean = 0.0f;
-        for (int i = threadIdx.x; i < n; i += blockDim.x) {
-            int b = i / spatial_size;
-            int s = i % spatial_size;
-            mean += input[(b * channels + c) * spatial_size + s];
-        }
-        __shared__ float smean[BLOCK_SIZE];
-        smean[threadIdx.x] = mean;
-        __syncthreads();
-        for (int s = blockDim.x / 2; s > 0; s >>= 1) {
-            if (threadIdx.x < s) smean[threadIdx.x] += smean[threadIdx.x + s];
-            __syncthreads();
-        }
-        mean = smean[0] / (float)n;
-
-        float var = 0.0f;
-        for (int i = threadIdx.x; i < n; i += blockDim.x) {
-            int b = i / spatial_size;
-            int s = i % spatial_size;
-            float diff = input[(b * channels + c) * spatial_size + s] - mean;
-            var += diff * diff;
-        }
-        __shared__ float svar[BLOCK_SIZE];
-        svar[threadIdx.x] = var;
-        __syncthreads();
-        for (int s = blockDim.x / 2; s > 0; s >>= 1) {
-            if (threadIdx.x < s) svar[threadIdx.x] += svar[threadIdx.x + s];
-            __syncthreads();
-        }
-        var = svar[0] / (float)n;
-
-        float inv_var = rsqrtf(var + epsilon);
-
-        if (threadIdx.x == 0) {
-            save_mean[c] = mean;
-            save_inv_var[c] = inv_var;
-            running_mean[c] = (1.0f - momentum) * running_mean[c] + momentum * mean;
-            running_var[c] = (1.0f - momentum) * running_var[c] + momentum * var * (float)n / (float)(n - 1);
-        }
-        __syncthreads();
-
-        for (int i = threadIdx.x; i < n; i += blockDim.x) {
-            int b = i / spatial_size;
-            int s = i % spatial_size;
-            int idx = (b * channels + c) * spatial_size + s;
-            float normalized = (input[idx] - mean) * inv_var;
-            output[idx] = gamma[c] * normalized + beta[c];
-        }
-    } else {
-        float mean = running_mean[c];
-        float inv_var = rsqrtf(running_var[c] + epsilon);
-        for (int i = threadIdx.x; i < n; i += blockDim.x) {
-            int b = i / spatial_size;
-            int s = i % spatial_size;
-            int idx = (b * channels + c) * spatial_size + s;
-            output[idx] = gamma[c] * (input[idx] - mean) * inv_var + beta[c];
-        }
-    }
-}
-
-extern "C" int fastnn_cuda_batch_norm_forward(
-    const float* input, const float* gamma, const float* beta,
-    float* running_mean, float* running_var,
-    float* output, float* save_mean, float* save_inv_var,
-    int batch_size, int channels, int spatial_size,
-    float momentum, float epsilon, int training)
-{
-    kernel_batch_norm_forward<<<channels, min(batch_size * spatial_size, BLOCK_SIZE)>>>(
-        input, gamma, beta, running_mean, running_var,
-        output, save_mean, save_inv_var,
-        batch_size, channels, spatial_size, momentum, epsilon, training);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_batch_norm_backward(
-    const float* grad_output, const float* input,
-    const float* gamma, const float* save_mean, const float* save_inv_var,
-    float* grad_input, float* grad_gamma, float* grad_beta,
-    int batch_size, int channels, int spatial_size)
-{
-    int c = blockIdx.x;
-    if (c >= channels) return;
-
-    int n = batch_size * spatial_size;
-    float mean = save_mean[c];
-    float inv_var = save_inv_var[c];
-    float g = gamma[c];
-
-    float dg = 0.0f, db = 0.0f;
-    for (int i = threadIdx.x; i < n; i += blockDim.x) {
-        int b = i / spatial_size;
-        int s = i % spatial_size;
-        int idx = (b * channels + c) * spatial_size + s;
-        float xhat = (input[idx] - mean) * inv_var;
-        dg += grad_output[idx] * xhat;
-        db += grad_output[idx];
-    }
-
-    __shared__ float sdg[BLOCK_SIZE], sdb[BLOCK_SIZE];
-    sdg[threadIdx.x] = dg;
-    sdb[threadIdx.x] = db;
-    __syncthreads();
-    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
-        if (threadIdx.x < s) {
-            sdg[threadIdx.x] += sdg[threadIdx.x + s];
-            sdb[threadIdx.x] += sdb[threadIdx.x + s];
-        }
-        __syncthreads();
-    }
-    dg = sdg[0];
-    db = sdb[0];
-
-    if (threadIdx.x == 0) {
-        grad_gamma[c] = dg;
-        grad_beta[c] = db;
-    }
-    __syncthreads();
-
-    float inv_n = 1.0f / (float)n;
-    for (int i = threadIdx.x; i < n; i += blockDim.x) {
-        int b = i / spatial_size;
-        int s = i % spatial_size;
-        int idx = (b * channels + c) * spatial_size + s;
-        float xhat = (input[idx] - mean) * inv_var;
-        grad_input[idx] = g * inv_var * inv_n * ((float)n * grad_output[idx] - db - xhat * dg);
-    }
-}
-
-extern "C" int fastnn_cuda_batch_norm_backward(
-    const float* grad_output, const float* input,
-    const float* gamma, const float* save_mean, const float* save_inv_var,
-    float* grad_input, float* grad_gamma, float* grad_beta,
-    int batch_size, int channels, int spatial_size)
-{
-    kernel_batch_norm_backward<<<channels, min(batch_size * spatial_size, BLOCK_SIZE)>>>(
-        grad_output, input, gamma, save_mean, save_inv_var,
-        grad_input, grad_gamma, grad_beta,
-        batch_size, channels, spatial_size);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
 __global__ void kernel_layer_norm_forward(
     const float* input, const float* gamma, const float* beta,
     float* output, float* mean_out, float* inv_var_out,
@@ -1442,80 +903,6 @@ extern "C" int fastnn_cuda_layer_norm_backward(
     return 0;
 }
 
-__global__ void kernel_rms_norm_forward(
-    const float* input, const float* gamma,
-    float* output, float* rms_out,
-    int batch_size, int normalized_size, float epsilon)
-{
-    int b = blockIdx.x;
-    if (b >= batch_size) return;
-
-    extern __shared__ float shared[];
-    const float* in_row = input + b * normalized_size;
-    float* out_row = output + b * normalized_size;
-
-    float thread_sq = 0.0f;
-    for (int i = threadIdx.x; i < normalized_size; i += blockDim.x)
-        thread_sq += in_row[i] * in_row[i];
-    shared[threadIdx.x] = thread_sq;
-    __syncthreads();
-    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
-        if (threadIdx.x < s) shared[threadIdx.x] += shared[threadIdx.x + s];
-        __syncthreads();
-    }
-    float rms = rsqrtf(shared[0] / (float)normalized_size + epsilon);
-    if (threadIdx.x == 0) rms_out[b] = rms;
-    __syncthreads();
-
-    for (int i = threadIdx.x; i < normalized_size; i += blockDim.x)
-        out_row[i] = gamma[i] * in_row[i] * rms;
-}
-
-extern "C" int fastnn_cuda_rms_norm_forward(
-    const float* input, const float* gamma,
-    float* output, float* rms,
-    int batch_size, int normalized_size, float epsilon)
-{
-    int threads = reduction_threads(normalized_size);
-    kernel_rms_norm_forward<<<batch_size, threads, threads * sizeof(float)>>>(
-        input, gamma, output, rms, batch_size, normalized_size, epsilon);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_dropout_forward(const float* input, float* output, float* mask, float p, size_t n, unsigned long long seed) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx >= n) return;
-
-    curandState state;
-    curand_init(seed, idx, 0, &state);
-    float r = curand_uniform(&state);
-    float m = (r > p) ? 1.0f : 0.0f;
-    float scale = 1.0f / (1.0f - p);
-    mask[idx] = m;
-    output[idx] = input[idx] * m * scale;
-}
-
-__global__ void kernel_dropout_backward(const float* grad_output, const float* mask, float* grad_input, float p, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx >= n) {
-        return;
-    }
-    float scale = 1.0f / (1.0f - p);
-    grad_input[idx] = grad_output[idx] * mask[idx] * scale;
-}
-
-extern "C" int fastnn_cuda_dropout_forward(const float* input, float* output, float* mask, float p, size_t n, unsigned long long seed) {
-    int blocks = div_ceil((int)n, BLOCK_SIZE);
-    kernel_dropout_forward<<<blocks, BLOCK_SIZE>>>(input, output, mask, p, n, seed);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-extern "C" int fastnn_cuda_dropout_backward(const float* grad_output, const float* mask, float* grad_input, float p, size_t n) {
-    LAUNCH_ELEMENTWISE(kernel_dropout_backward, grad_output, mask, grad_input, p, n);
-}
-
 __global__ void kernel_embedding_forward(const int* indices, const float* weight, float* output, int num_indices, int embedding_dim) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     if (idx >= num_indices * embedding_dim) return;
@@ -1555,241 +942,6 @@ extern "C" int fastnn_cuda_embedding_backward(const int* indices, const float* g
     return 0;
 }
 
-__global__ void kernel_cross_entropy_loss(const float* log_probs, const int* targets, float* loss, float* grad,
-                                           int batch_size, int num_classes) {
-    int b = blockIdx.x * blockDim.x + threadIdx.x;
-    if (b >= batch_size) return;
-
-    int target = targets[b];
-    float nll = -log_probs[b * num_classes + target];
-    atomicAdd(loss, nll / (float)batch_size);
-
-    for (int c = 0; c < num_classes; c++) {
-        float prob = expf(log_probs[b * num_classes + c]);
-        grad[b * num_classes + c] = (prob - (c == target ? 1.0f : 0.0f)) / (float)batch_size;
-    }
-}
-
-extern "C" int fastnn_cuda_cross_entropy_loss(const float* log_probs, const int* targets, float* loss, float* grad,
-                                                int batch_size, int num_classes) {
-    CUDA_CHECK(cudaMemset(loss, 0, sizeof(float)));
-    int blocks = div_ceil(batch_size, BLOCK_SIZE);
-    kernel_cross_entropy_loss<<<blocks, BLOCK_SIZE>>>(log_probs, targets, loss, grad, batch_size, num_classes);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_mse_loss(const float* predictions, const float* targets, float* loss, float* grad, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx >= n) return;
-
-    float diff = predictions[idx] - targets[idx];
-    atomicAdd(loss, diff * diff / (float)n);
-    grad[idx] = 2.0f * diff / (float)n;
-}
-
-extern "C" int fastnn_cuda_mse_loss(const float* predictions, const float* targets, float* loss, float* grad, size_t n) {
-    CUDA_CHECK(cudaMemset(loss, 0, sizeof(float)));
-    int blocks = div_ceil((int)n, BLOCK_SIZE);
-    kernel_mse_loss<<<blocks, BLOCK_SIZE>>>(predictions, targets, loss, grad, n);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_bce_loss(const float* predictions, const float* targets, float* loss, float* grad, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx >= n) return;
-
-    float p = fminf(fmaxf(predictions[idx], 1e-7f), 1.0f - 1e-7f);
-    float t = targets[idx];
-    atomicAdd(loss, -(t * logf(p) + (1.0f - t) * logf(1.0f - p)) / (float)n);
-    grad[idx] = (-t / p + (1.0f - t) / (1.0f - p)) / (float)n;
-}
-
-extern "C" int fastnn_cuda_binary_cross_entropy(const float* predictions, const float* targets, float* loss, float* grad, size_t n) {
-    CUDA_CHECK(cudaMemset(loss, 0, sizeof(float)));
-    int blocks = div_ceil((int)n, BLOCK_SIZE);
-    kernel_bce_loss<<<blocks, BLOCK_SIZE>>>(predictions, targets, loss, grad, n);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-extern "C" int fastnn_cuda_scaled_dot_product_attention(
-    const float* query, const float* key, const float* value,
-    float* output, float* attn_weights,
-    const float* mask,
-    int batch_size, int num_heads, int seq_len_q, int seq_len_k, int head_dim,
-    float scale, int causal)
-{
-    int bh = batch_size * num_heads;
-
-    float alpha = scale;
-    float beta = 0.0f;
-
-    float* key_t;
-    CUDA_CHECK(cudaMalloc(&key_t, (size_t)bh * head_dim * seq_len_k * sizeof(float)));
-    fastnn_cuda_transpose_batched(key, key_t, bh, seq_len_k, head_dim);
-
-    // scores = scale * Q K^T
-    CUBLAS_CHECK(cublasSgemmStridedBatched(g_cublas_handle,
-        CUBLAS_OP_N, CUBLAS_OP_N,
-        seq_len_k, seq_len_q, head_dim,
-        &alpha,
-        key_t, seq_len_k, (long long)head_dim * seq_len_k,
-        query, head_dim, (long long)seq_len_q * head_dim,
-        &beta,
-        attn_weights, seq_len_k, (long long)seq_len_q * seq_len_k,
-        bh));
-
-    cudaFree(key_t);
-
-    // `causal` and `mask` are not applied yet.
-    fastnn_cuda_softmax(attn_weights, attn_weights, bh * seq_len_q, seq_len_k);
-
-    alpha = 1.0f;
-    beta = 0.0f;
-    CUBLAS_CHECK(cublasSgemmStridedBatched(g_cublas_handle,
-        CUBLAS_OP_N, CUBLAS_OP_N,
-        head_dim, seq_len_q, seq_len_k,
-        &alpha,
-        value, head_dim, (long long)seq_len_k * head_dim,
-        attn_weights, seq_len_k, (long long)seq_len_q * seq_len_k,
-        &beta,
-        output, head_dim, (long long)seq_len_q * head_dim,
-        bh));
-
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_sgd_step(float* params, const float* grads, float* velocity,
-                                 float lr, float momentum, float weight_decay, float dampening,
-                                 int nesterov, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx >= n) return;
-
-    float g = grads[idx];
-    if (weight_decay != 0.0f) g += weight_decay * params[idx];
-
-    if (momentum != 0.0f) {
-        float v = velocity[idx];
-        v = momentum * v + (1.0f - dampening) * g;
-        velocity[idx] = v;
-        if (nesterov) g = g + momentum * v;
-        else g = v;
-    }
-
-    params[idx] -= lr * g;
-}
-
-extern "C" int fastnn_cuda_sgd_step(float* params, const float* grads, float* velocity,
-                                      float lr, float momentum, float weight_decay, float dampening,
-                                      int nesterov, size_t n) {
-    int blocks = div_ceil((int)n, BLOCK_SIZE);
-    kernel_sgd_step<<<blocks, BLOCK_SIZE>>>(params, grads, velocity, lr, momentum, weight_decay, dampening, nesterov, n);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_adam_step(float* params, const float* grads,
-                                  float* m, float* v,
-                                  float lr, float beta1, float beta2, float epsilon,
-                                  float weight_decay, int step, int amsgrad, float* v_max,
-                                  size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx >= n) return;
-
-    float g = grads[idx];
-
-    // decay is applied to the weights directly, not folded into the gradient
-    if (weight_decay != 0.0f) {
-        params[idx] -= lr * weight_decay * params[idx];
-    }
-
-    m[idx] = beta1 * m[idx] + (1.0f - beta1) * g;
-    v[idx] = beta2 * v[idx] + (1.0f - beta2) * g * g;
-
-    float m_hat = m[idx] / (1.0f - powf(beta1, (float)step));
-    float v_hat = v[idx] / (1.0f - powf(beta2, (float)step));
-
-    if (amsgrad && v_max != nullptr) {
-        v_max[idx] = fmaxf(v_max[idx], v_hat);
-        v_hat = v_max[idx];
-    }
-
-    params[idx] -= lr * m_hat / (sqrtf(v_hat) + epsilon);
-}
-
-extern "C" int fastnn_cuda_adam_step(float* params, const float* grads,
-                                      float* m, float* v,
-                                      float lr, float beta1, float beta2, float epsilon,
-                                      float weight_decay, int step, int amsgrad, float* v_max,
-                                      size_t n) {
-    int blocks = div_ceil((int)n, BLOCK_SIZE);
-    kernel_adam_step<<<blocks, BLOCK_SIZE>>>(params, grads, m, v, lr, beta1, beta2, epsilon, weight_decay, step, amsgrad, v_max, n);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_fill(float* data, float value, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx < n) data[idx] = value;
-}
-
-extern "C" int fastnn_cuda_fill(float* data, float value, size_t n) {
-    LAUNCH_ELEMENTWISE(kernel_fill, data, value, n);
-}
-
-extern "C" int fastnn_cuda_copy(const float* src, float* dst, size_t n) {
-    CUDA_CHECK(cudaMemcpy(dst, src, n * sizeof(float), cudaMemcpyDeviceToDevice));
-    return 0;
-}
-
-__global__ void kernel_arange(float* output, float start, float step, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx < n) output[idx] = start + (float)idx * step;
-}
-
-extern "C" int fastnn_cuda_arange(float* output, float start, float step, size_t n) {
-    int blocks = div_ceil((int)n, BLOCK_SIZE);
-    kernel_arange<<<blocks, BLOCK_SIZE>>>(output, start, step, n);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
-__global__ void kernel_where(const float* condition, const float* x, const float* y, float* output, size_t n) {
-    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx < n) output[idx] = condition[idx] != 0.0f ? x[idx] : y[idx];
-}
-
-extern "C" int fastnn_cuda_where(const float* condition, const float* x, const float* y, float* output, size_t n) {
-    LAUNCH_ELEMENTWISE(kernel_where, condition, x, y, output, n);
-}
-
-__global__ void kernel_gather(const float* input, const int* indices, float* output,
-                               int outer_size, int gather_dim_size, int inner_size, int num_indices) {
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int total = outer_size * num_indices * inner_size;
-    if (idx >= total) return;
-
-    int inner_idx = idx % inner_size;
-    int index_idx = (idx / inner_size) % num_indices;
-    int outer_idx = idx / inner_size / num_indices;
-
-    int gather_idx = indices[index_idx];
-    int input_idx = (outer_idx * gather_dim_size + gather_idx) * inner_size + inner_idx;
-    output[idx] = input[input_idx];
-}
-
-extern "C" int fastnn_cuda_gather(const float* input, const int* indices, float* output,
-                                    int outer_size, int gather_dim_size, int inner_size, int num_indices) {
-    int total = outer_size * num_indices * inner_size;
-    int blocks = div_ceil(total, BLOCK_SIZE);
-    kernel_gather<<<blocks, BLOCK_SIZE>>>(input, indices, output, outer_size, gather_dim_size, inner_size, num_indices);
-    CUDA_CHECK(cudaGetLastError());
-    return 0;
-}
-
 __global__ void kernel_im2col(const float* input, float* cols,
                               int c, int h, int w,
                               int kh, int kw, int sh, int sw, int ph, int pw,
diff --git a/cuda/stubs.c b/cuda/stubs.c
index 7f460fc..ce66987 100644
--- a/cuda/stubs.c
+++ b/cuda/stubs.c
@@ -9,7 +9,6 @@ int fastnn_cuda_memcpy_h2d(float *d, const float *h, size_t n)  { (void)d;(void)
 int fastnn_cuda_memcpy_d2h(float *h, const float *d, size_t n)  { (void)h;(void)d;(void)n; STUB; }
 int fastnn_cuda_memcpy_d2d(float *d, const float *s, size_t n)  { (void)d;(void)s;(void)n; STUB; }
 int fastnn_cuda_memset(float *p, int v, size_t n)               { (void)p;(void)v;(void)n; STUB; }
-int fastnn_cuda_fill(float *p, float v, size_t n)               { (void)p;(void)v;(void)n; STUB; }
 
 int fastnn_cuda_add(const float*a,const float*b,float*o,size_t n)            {(void)a;(void)b;(void)o;(void)n;STUB;}
 int fastnn_cuda_sub(const float*a,const float*b,float*o,size_t n)            {(void)a;(void)b;(void)o;(void)n;STUB;}
@@ -47,8 +46,6 @@ int fastnn_cuda_sum(const float*i,float*o,size_t n)                          {(v
 int fastnn_cuda_mean(const float*i,float*o,size_t n)                         {(void)i;(void)o;(void)n;STUB;}
 int fastnn_cuda_max(const float*i,float*o,size_t n)                          {(void)i;(void)o;(void)n;STUB;}
 int fastnn_cuda_min(const float*i,float*o,size_t n)                          {(void)i;(void)o;(void)n;STUB;}
-int fastnn_cuda_argmax(const float*i,int*o,size_t n)                         {(void)i;(void)o;(void)n;STUB;}
-int fastnn_cuda_argmax_axis(const float*i,int*o,size_t rows,size_t cols,int axis){(void)i;(void)o;(void)rows;(void)cols;(void)axis;STUB;}
 int fastnn_cuda_matmul(const float*a,const float*b,float*c,int m,int k,int n){(void)a;(void)b;(void)c;(void)m;(void)k;(void)n;STUB;}
 int fastnn_cuda_matmul_batched(const float*a,const float*b,float*c,int batch,int m,int k,int n){(void)a;(void)b;(void)c;(void)batch;(void)m;(void)k;(void)n;STUB;}
 int fastnn_cuda_matmul_nt(const float*a,const float*b,float*c,int m,int n,int k){(void)a;(void)b;(void)c;(void)m;(void)n;(void)k;STUB;}
diff --git a/src/cuda/ffi.rs b/src/cuda/ffi.rs
index 75f5b94..8f7c45a 100644
--- a/src/cuda/ffi.rs
+++ b/src/cuda/ffi.rs
@@ -18,7 +18,6 @@ extern "C" {
     pub fn fastnn_cuda_memcpy_d2h(dst: *mut f32, src: *const f32, bytes: usize) -> c_int;
     pub fn fastnn_cuda_memcpy_d2d(dst: *mut f32, src: *const f32, bytes: usize) -> c_int;
     pub fn fastnn_cuda_memset(ptr: *mut f32, value: c_int, bytes: usize) -> c_int;
-    pub fn fastnn_cuda_fill(data: *mut f32, value: f32, n: usize) -> c_int;
 
     pub fn fastnn_cuda_add(a: *const f32, b: *const f32, out: *mut f32, n: usize) -> c_int;
     pub fn fastnn_cuda_sub(a: *const f32, b: *const f32, out: *mut f32, n: usize) -> c_int;
@@ -245,14 +244,6 @@ extern "C" {
     pub fn fastnn_cuda_mean(x: *const f32, out: *mut f32, n: usize) -> c_int;
     pub fn fastnn_cuda_max(x: *const f32, out: *mut f32, n: usize) -> c_int;
     pub fn fastnn_cuda_min(x: *const f32, out: *mut f32, n: usize) -> c_int;
-    pub fn fastnn_cuda_argmax(x: *const f32, out: *mut c_int, n: usize) -> c_int;
-    pub fn fastnn_cuda_argmax_axis(
-        x: *const f32,
-        out: *mut c_int,
-        outer: c_int,
-        axis: c_int,
-        inner: c_int,
-    ) -> c_int;
     pub fn fastnn_cuda_sum_axis(
         x: *const f32,
         out: *mut f32,