Initial commit

.gitignore

+.idea
+target
+Cargo.lock
\ No newline at end of file

Cargo.toml

+[package]
+name = "candle-rotary"
+version = "0.0.1"
+edition = "2021"
+
+description = "Rotary layer for the candle ML framework."
+keywords = ["tensor", "machine-learning"]
+categories = ["science"]
+license = "MIT OR Apache-2.0"
+readme = "README.md"
+
+[dependencies]
+candle = { version = "*", package = "candle-core", features = ["cuda"]}
+half = { version = "2.3.1", features = ["num-traits"] }
+
+[build-dependencies]
+anyhow = { version = "1", features = ["backtrace"] }
+bindgen_cuda = "0.1.1"
+
+[dev-dependencies]
+anyhow = { version = "1", features = ["backtrace"] }
+candle-nn = { version = "0.3.0", features = ["cuda"] }

LICENSE-APACHE

+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright [yyyy] [name of copyright owner]
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

LICENSE-MIT

+Permission is hereby granted, free of charge, to any
+person obtaining a copy of this software and associated
+documentation files (the "Software"), to deal in the
+Software without restriction, including without
+limitation the rights to use, copy, modify, merge,
+publish, distribute, sublicense, and/or sell copies of
+the Software, and to permit persons to whom the Software
+is furnished to do so, subject to the following
+conditions:
+
+The above copyright notice and this permission notice
+shall be included in all copies or substantial portions
+of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF
+ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED
+TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A
+PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT
+SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION
+OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR
+IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+DEALINGS IN THE SOFTWARE.

README.md

+# Candle Rotary
+
+All files in `kernels` are adapted from https://github.com/vllm-project/vllm/tree/main/csrc and are under the vLLM
+Project copyright.
\ No newline at end of file

build.rs

+// Build script to run nvcc and generate the C glue code for launching the flash-attention kernel.
+// The cuda build time is very long so one can set the CANDLE_FLASH_ATTN_BUILD_DIR environment
+// variable in order to cache the compiled artifacts and avoid recompiling too often.
+use anyhow::{Context, Result};
+use std::path::PathBuf;
+
+const KERNEL_FILES: [&str; 1] = ["kernels/rotary.cu"];
+
+fn main() -> Result<()> {
+    println!("cargo:rerun-if-changed=build.rs");
+    for kernel_file in KERNEL_FILES.iter() {
+        println!("cargo:rerun-if-changed={kernel_file}");
+    }
+    let out_dir = PathBuf::from(std::env::var("OUT_DIR").context("OUT_DIR not set")?);
+    let build_dir = match std::env::var("CANDLE_ROTARY_BUILD_DIR") {
+        Err(_) =>
+        {
+            #[allow(clippy::redundant_clone)]
+            out_dir.clone()
+        }
+        Ok(build_dir) => {
+            let path = PathBuf::from(build_dir);
+            path.canonicalize().expect(&format!(
+                "Directory doesn't exists: {} (the current directory is {})",
+                &path.display(),
+                std::env::current_dir()?.display()
+            ))
+        }
+    };
+
+    let kernels : Vec<_>= KERNEL_FILES.iter().collect();
+    let builder = bindgen_cuda::Builder::default().kernel_paths(kernels).out_dir(build_dir.clone())
+                    .arg("-std=c++17")
+                    .arg("-O3")
+                    .arg("-U__CUDA_NO_HALF_OPERATORS__")
+                    .arg("-U__CUDA_NO_HALF_CONVERSIONS__")
+                    .arg("-U__CUDA_NO_HALF2_OPERATORS__")
+                    .arg("-U__CUDA_NO_BFLOAT16_CONVERSIONS__")
+                    .arg("--expt-relaxed-constexpr")
+                    .arg("--expt-extended-lambda")
+                    .arg("--use_fast_math")
+                    .arg("--ptxas-options=-v")
+                    .arg("--verbose")
+                    .arg("-fPIC")
+                    .arg("--gpu-max-threads-per-block=256");
+
+    let out_file = build_dir.join("librotary.a");
+    builder.build_lib(out_file);
+
+    println!("cargo:rustc-link-search={}", build_dir.display());
+    println!("cargo:rustc-link-lib=rotary");
+    println!("cargo:rustc-link-lib=dylib=cudart");
+    println!("cargo:rustc-link-lib=dylib=stdc++");
+
+    Ok(())
+}

kernels/cuda_compat.h

+#pragma once
+
+#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)
+#else
+  #define VLLM_SHFL_XOR_SYNC(var, lane_mask) __shfl_xor(var, lane_mask)
+#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_DevFuncAttribute_SET_MaxDynamicSharedMemorySize(FUNC, VAL) \
+    cudaFuncSetAttribute(FUNC, cudaFuncAttributeMaxDynamicSharedMemorySize, VAL)
+#else
+  #define VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize(FUNC, VAL) \
+    hipFuncSetAttribute(FUNC, hipFuncAttributeMaxDynamicSharedMemorySize, VAL)
+#endif
+

kernels/rotary.cu

+#include <cuda_fp16.h>
+#include <cuda_bf16.h>
+#include <stdint.h>
+
+#include "cuda_compat.h"
+
+namespace vllm {
+
+template<typename scalar_t, bool IS_NEOX>
+inline __device__ void apply_rotary_embedding(
+  scalar_t* __restrict__ arr,
+  const scalar_t* __restrict__ cos_ptr,
+  const scalar_t* __restrict__ sin_ptr,
+  int rot_offset,
+  int rot_dim)
+{
+  int x_index, y_index;
+  scalar_t cos, sin;
+  if (IS_NEOX) {
+    // GPT-NeoX style rotary embedding.
+    x_index = rot_offset;
+    y_index = rot_dim + rot_offset;
+    cos = VLLM_LDG(cos_ptr + x_index);
+    sin = VLLM_LDG(sin_ptr + x_index);
+  } else {
+    // GPT-J style rotary embedding.
+    x_index = 2 * rot_offset;
+    y_index = 2 * rot_offset + 1;
+    cos = VLLM_LDG(cos_ptr + x_index / 2);
+    sin = VLLM_LDG(sin_ptr + x_index / 2);
+  }
+
+  const scalar_t x = arr[x_index];
+  const scalar_t y = arr[y_index];
+  arr[x_index] = x * cos - y * sin;
+  arr[y_index] = y * cos + x * sin;
+}
+
+template<typename scalar_t, bool IS_NEOX>
+__global__ void rotary_embedding_kernel(
+  scalar_t* __restrict__ query,                    // [num_tokens, num_heads, head_size]
+  scalar_t* __restrict__ key,                      // [num_tokens, num_heads, head_size]
+  const scalar_t* __restrict__ cos_cache,          // [num_tokens, rot_dim]
+  const scalar_t* __restrict__ sin_cache,          // [num_tokens, rot_dim]
+  const int rot_dim,
+  const int64_t query_stride,
+  const int64_t key_stride,
+  const int num_heads,
+  const int num_kv_heads,
+  const int head_size) {
+  // Each thread block is responsible for one token.
+  const int token_idx = blockIdx.x;
+
+  const scalar_t* cos_ptr = cos_cache + token_idx * rot_dim;
+  const scalar_t* sin_ptr = sin_cache + token_idx * rot_dim;
+
+  const int nq = num_heads * rot_dim;
+  for (int i = threadIdx.x; i < nq; i += blockDim.x) {
+    const int head_idx = i / rot_dim;
+    const int64_t token_head = token_idx * query_stride + head_idx * head_size;
+    const int rot_offset = i % rot_dim;
+    apply_rotary_embedding<scalar_t, IS_NEOX>(query + token_head, cos_ptr,
+                                              sin_ptr, rot_offset, rot_dim);
+  }
+
+  const int nk = num_kv_heads * rot_dim;
+  for (int i = threadIdx.x; i < nk; i += blockDim.x) {
+    const int head_idx = i / rot_dim;
+    const int64_t token_head = token_idx * key_stride + head_idx * head_size;
+    const int rot_offset = i % rot_dim;
+    apply_rotary_embedding<scalar_t, IS_NEOX>(key + token_head, cos_ptr,
+                                              sin_ptr, rot_offset, rot_dim);
+  }
+}
+
+} // namespace vllm
+
+#define CALL_ROTARY(T, IS_NEOX)                                               \
+  vllm::rotary_embedding_kernel<T, IS_NEOX><<<grid, block, 0, stream>>>(      \
+    reinterpret_cast<T*>(query),                                              \
+    reinterpret_cast<T*>(key),                                                \
+    reinterpret_cast<T*>(cos_cache),                                          \
+    reinterpret_cast<T*>(sin_cache),                                          \
+    rot_dim,                                                                  \
+    query_stride,                                                             \
+    key_stride,                                                               \
+    num_heads,                                                                \
+    num_kv_heads,                                                             \
+    head_size);
+
+extern "C" void rotary_embedding(
+  void *query,             // [num_tokens, num_heads, head_size]
+  void *key,               // [num_tokens, num_kv_heads, head_size]
+  void *cos_cache,         // [num_tokens, rot_dim]
+  void *sin_cache,         // [num_tokens, rot_dim]
+  int32_t is_neox,
+
+  int32_t head_size,
+  int64_t num_tokens,
+  int32_t rot_dim,
+  int32_t num_heads,
+  int32_t num_kv_heads,
+  int64_t query_stride,
+  int64_t key_stride,
+
+  uint32_t dtype // 0 => f16; 1 => bf16; 2 => f32
+  ) {
+
+  dim3 grid(num_tokens);
+  dim3 block(std::min(num_heads * rot_dim, 512));
+  const cudaStream_t stream = 0;
+  const bool is_neox_bool = is_neox;
+
+  if (is_neox_bool) {
+    if (dtype == 0){
+      CALL_ROTARY(half, true);
+    } else if (dtype == 1) {
+      CALL_ROTARY(__nv_bfloat16, true);
+    } else if (dtype == 2) {
+      CALL_ROTARY(float, true);
+    }
+  } else {
+    if (dtype == 0){
+      CALL_ROTARY(half, false);
+    } else if (dtype == 1) {
+      CALL_ROTARY(__nv_bfloat16, false);
+    } else if (dtype == 2) {
+      CALL_ROTARY(float, false);
+    }
+  }
+}

src/ffi.rs

+use core::ffi::{c_int, c_long, c_void};
+
+extern "C" {
+    pub(crate) fn rotary_embedding(
+        query: *const c_void,
+        key: *const c_void,
+        cos_cache: *const c_void,
+        sin_cache: *const c_void,
+
+        is_neox: c_int,
+
+        head_size: c_int,
+        num_tokens: c_long,
+        rot_dim: c_int,
+        num_heads: c_int,
+        num_kv_heads: c_int,
+        query_stride: c_long,
+        key_stride: c_long,
+
+        dtype: u32,
+    );
+}

src/lib.rs

+mod ffi;
+
+use candle::cuda_backend::cudarc::driver::DevicePtr;
+use candle::{DType, Device, Result, Storage, Tensor};
+use half::{bf16, f16};
+use std::ffi::{c_int, c_long};
+
+fn apply_rotary_<
+    T: candle::cuda_backend::CudaDType + candle::cuda_backend::cudarc::driver::DeviceRepr,
+>(
+    query: &Tensor,
+    key: &Tensor,
+    cos_cache: &Tensor,
+    sin_cache: &Tensor,
+    is_neox: bool,
+) -> Result<()> {
+    let dtype = query.dtype();
+    if key.dtype() != dtype || cos_cache.dtype() != dtype || sin_cache.dtype() != dtype {
+        candle::bail!("apply-rotary expects all tensors to have the same dtype");
+    }
+
+    let internal_type = match dtype {
+        DType::F16 => 0,
+        DType::BF16 => 1,
+        DType::F32 => 2,
+        dtype => candle::bail!("dtype {dtype:?} is not supported"),
+    };
+
+    let (q, q_l) = query.storage_and_layout();
+    let q = match &*q {
+        Storage::Cuda(q) => q,
+        _ => candle::bail!("query must be a cuda tensor"),
+    };
+
+    let (k, k_l) = key.storage_and_layout();
+    let k = match &*k {
+        Storage::Cuda(k) => k,
+        _ => candle::bail!("key must be a cuda tensor"),
+    };
+
+    let (cc, cc_l) = cos_cache.storage_and_layout();
+    let cc = match &*cc {
+        Storage::Cuda(cc) => cc,
+        _ => candle::bail!("cos_cache must be a cuda tensor"),
+    };
+
+    let (sc, sc_l) = sin_cache.storage_and_layout();
+    let sc = match &*sc {
+        Storage::Cuda(sc) => sc,
+        _ => candle::bail!("sin_cache must be a cuda tensor"),
+    };
+
+    let q_rank = q_l.stride().len();
+    let k_rank = k_l.stride().len();
+    let cc_rank = cc_l.stride().len();
+    let sc_rank = sc_l.stride().len();
+
+    if q_rank != 3 || k_rank != 3 {
+        candle::bail!("apply-rotary expects input tensors of rank 3 (k: {q_l:?}, v: {k_l:?})")
+    }
+
+    if cc_rank != 2 || sc_rank != 2 {
+        candle::bail!("apply-rotary expects cache tensors of rank 2 (k: {cc_l:?}, v: {sc_l:?})")
+    }
+
+    // Get cuda slices for all tensors
+    let q = q.as_cuda_slice::<T>()?;
+    let k = k.as_cuda_slice::<T>()?;
+    let cc = cc.as_cuda_slice::<T>()?;
+    let sc = sc.as_cuda_slice::<T>()?;
+
+    // Get cuda views for all tensors
+    let q = q.slice(q_l.start_offset()..);
+    let k = k.slice(k_l.start_offset()..);
+    let cc = cc.slice(cc_l.start_offset()..);
+    let sc = sc.slice(sc_l.start_offset()..);
+
+    let (num_tokens, num_heads, head_size) = q_l.shape().dims3()?;
+    let (num_tokens_kv, num_kv_heads, head_size_kv) = k_l.shape().dims3()?;
+
+    if (num_tokens, head_size) != (num_tokens_kv, head_size_kv) {
+        candle::bail!("shape mismatch q {:?} and k {:?}", q_l.shape(), k_l.shape())
+    }
+
+    let rot_dim = cc_l.dims()[1];
+    if (num_tokens, rot_dim) != cc_l.shape().dims2()? {
+        candle::bail!(
+            "shape mismatch cos_cache {:?}, expected {:?}",
+            cc_l.shape(),
+            (num_tokens, rot_dim)
+        )
+    }
+
+    if (num_tokens, rot_dim) != sc_l.shape().dims2()? {
+        candle::bail!(
+            "shape mismatch sin_cache {:?}, expected {:?}",
+            sc_l.shape(),
+            (num_tokens, rot_dim)
+        )
+    }
+
+    let query_stride = q_l.stride()[0];
+    let key_stride = k_l.stride()[0];
+
+    let q_ptr = *q.device_ptr() as *const core::ffi::c_void;
+    let k_ptr = *k.device_ptr() as *const core::ffi::c_void;
+    let cc_ptr = *cc.device_ptr() as *const core::ffi::c_void;
+    let sc_ptr = *sc.device_ptr() as *const core::ffi::c_void;
+
+    let neox = if is_neox { 1 } else { 0 };
+
+    unsafe {
+        ffi::rotary_embedding(
+            q_ptr,
+            k_ptr,
+            cc_ptr,
+            sc_ptr,
+            neox,
+            head_size as c_int,
+            num_tokens as c_long,
+            rot_dim as c_int,
+            num_heads as c_int,
+            num_kv_heads as c_int,
+            query_stride as c_long,
+            key_stride as c_long,
+            internal_type,
+        )
+    }
+    Ok(())
+}
+
+pub fn inv_freqs(dim: usize, base: f32, device: &Device) -> Result<Tensor> {
+    let inv_freq: Vec<_> = (0..dim)
+        .step_by(2)
+        .map(|i| 1f32 / base.powf(i as f32 / dim as f32))
+        .collect();
+    let inv_freq_len = inv_freq.len();
+    Tensor::from_vec(inv_freq, (1, inv_freq_len), device)
+}
+
+pub fn cos_sin(length: usize, inv_freqs: &Tensor, dtype: DType) -> Result<(Tensor, Tensor)> {
+    let t = Tensor::arange(0u32, length as u32, inv_freqs.device())?
+        .to_dtype(DType::F32)?
+        .reshape((length, 1))?;
+    let freqs = t.matmul(&inv_freqs)?;
+    let cos = freqs.cos()?.to_dtype(dtype)?;
+    let sin = freqs.sin()?.to_dtype(dtype)?;
+    Ok((cos, sin))
+}
+
+/// Apply Rotary position encoding inplace
+///
+/// # Arguments
+///
+/// * `query` - Query tensor of shape `(num_tokens, num_heads, head_size)`.
+/// * `key` - Key tensor of shape `(num_tokens, num_kv_heads, head_size)`.
+/// * `cos_cache` - Aligned cache of shape `(num_tokens, rot_dim)`
+/// * `sin_cache` - Aligned cache of shape `(num_tokens, rot_dim)`
+/// * `is_neox` - Use neox encoding instead of gpt-j style rotary
+pub fn apply_rotary_inplace(
+    query: &Tensor,
+    key: &Tensor,
+    cos_cache: &Tensor,
+    sin_cache: &Tensor,
+    is_neox: bool,
+) -> Result<()> {
+    match key.dtype() {
+        DType::F16 => apply_rotary_::<f16>(query, key, cos_cache, sin_cache, is_neox),
+        DType::BF16 => apply_rotary_::<bf16>(query, key, cos_cache, sin_cache, is_neox),
+        DType::F32 => apply_rotary_::<f32>(query, key, cos_cache, sin_cache, is_neox),
+        dt => {
+            candle::bail!("apply_rotary is only supported for f32, f16 and bf16 ({dt:?})")
+        }
+    }
+}

tests/rotary_tests.rs

+use anyhow::Result;
+use candle::{DType, Device, Tensor, D};
+
+fn to_vec3_round(t: Tensor, digits: i32) -> Result<Vec<Vec<Vec<f32>>>> {
+    let b = 10f32.powi(digits);
+    let t = t.to_vec3::<f32>()?;
+    let t = t
+        .iter()
+        .map(|t| {
+            t.iter()
+                .map(|t| t.iter().map(|t| f32::round(t * b) / b).collect())
+                .collect()
+        })
+        .collect();
+    Ok(t)
+}
+
+fn rotate_half(xs: &Tensor) -> candle::error::Result<Tensor> {
+    let last_dim = xs.dim(D::Minus1)?;
+    let xs1 = xs.narrow(D::Minus1, 0, last_dim / 2)?;
+    let xs2 = xs.narrow(D::Minus1, last_dim / 2, last_dim / 2)?;
+    Tensor::cat(&[&xs2.neg()?, &xs1], D::Minus1)
+}
+
+fn freqs(rot_dim: usize, seqlen: usize, dev: &Device) -> candle::error::Result<Tensor> {
+    let inv_freq: Vec<_> = (0..rot_dim)
+        .step_by(2)
+        .map(|i| 1f32 / 10000f32.powf(i as f32 / rot_dim as f32))
+        .collect();
+    let inv_freq_len = inv_freq.len();
+    let inv_freq = Tensor::from_vec(inv_freq, (1, inv_freq_len), dev)?;
+    let t = Tensor::arange(0u32, seqlen as u32, dev)?
+        .to_dtype(DType::F32)?
+        .reshape((seqlen, 1))?;
+    t.matmul(&inv_freq)
+}
+
+fn apply_rotary_emb_qkv(
+    q: &Tensor,
+    k: &Tensor,
+    cos: &Tensor,
+    sin: &Tensor,
+) -> Result<(Tensor, Tensor)> {
+    let cos = cos.unsqueeze(1)?; // (seq_len, 1, dim)
+    let sin = sin.unsqueeze(1)?; // (seq_len, 1, dim)
+    let q_embed = (q.broadcast_mul(&cos)? + rotate_half(q)?.broadcast_mul(&sin))?;
+    let k_embed = (k.broadcast_mul(&cos)? + rotate_half(k)?.broadcast_mul(&sin))?;
+    Ok((q_embed, k_embed))
+}
+
+#[test]
+fn rotary() -> Result<()> {
+    let device = Device::new_cuda(0)?;
+
+    let seqlen = 12;
+    let num_heads = 8;
+    let rot_dim = 64;
+
+    let q = Tensor::randn(0.0, 1.0, (seqlen, num_heads, rot_dim), &device)?.to_dtype(DType::F32)?;
+    let k = Tensor::randn(0.0, 1.0, (seqlen, num_heads, rot_dim), &device)?.to_dtype(DType::F32)?;
+
+    let (expected_q, expected_k) = {
+        let freqs = freqs(rot_dim, seqlen, &device)?;
+        let freqs = Tensor::cat(&[&freqs, &freqs], D::Minus1)?;
+        apply_rotary_emb_qkv(&q, &k, &freqs.cos()?, &freqs.sin()?)
+    }?;
+
+    // Create inv freqs
+    let inv_freqs = candle_rotary::inv_freqs(rot_dim, 10000f32, &device)?;
+    // Create an over-sized cos sin cache like you would usually do
+    let (cos, sin) = candle_rotary::cos_sin(32, &inv_freqs, DType::F32)?;
+    // Positions for seqlen
+    let position_ids = Tensor::arange(0, seqlen as u32, &device)?;
+    // Filter cos and sin
+    let cos = cos.index_select(&position_ids, 0)?;
+    let sin = sin.index_select(&position_ids, 0)?;
+
+    // Inplace
+    candle_rotary::apply_rotary_inplace(&q, &k, &cos, &sin, true)?;
+
+    assert_eq!(to_vec3_round(expected_q, 3)?, to_vec3_round(q, 3)?);
+    assert_eq!(to_vec3_round(expected_k, 3)?, to_vec3_round(k, 3)?);
+
+    Ok(())
+}