diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 22089023f2..e8a49e9624 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -23705,9 +23705,6 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
     const int n_ctx_orig = ((int32_t *) dst->op_params)[4];
     const int n_offs     = ((int32_t *) dst->op_params)[15];
 
-    // when dst aliases src0, the channels outside the rotated window already hold the correct data
-    const int inplace = extrad->data_device == extra0->data_device && offsetd == offset0;
-
     float freq_base;
     float freq_scale;
     float ext_factor;
@@ -23827,13 +23824,11 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
     if (is_mrope && !is_vision) {
         CL_CHECK(clSetKernelArg(kernel, 34, sizeof(int), &is_imrope));
     }
-    // norm and neox have n_offs and inplace after beta_slow, mrope has them after is_imrope
+    // norm and neox have n_offs after beta_slow, mrope has it after is_imrope
     if (!is_mrope && !is_vision) {
         CL_CHECK(clSetKernelArg(kernel, 33, sizeof(int), &n_offs));
-        CL_CHECK(clSetKernelArg(kernel, 34, sizeof(int), &inplace));
     } else if (is_mrope && !is_vision) {
         CL_CHECK(clSetKernelArg(kernel, 35, sizeof(int), &n_offs));
-        CL_CHECK(clSetKernelArg(kernel, 36, sizeof(int), &inplace));
     }
 
     size_t global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03};
diff --git a/ggml/src/ggml-opencl/kernels/rope.cl b/ggml/src/ggml-opencl/kernels/rope.cl
index 77a72397f4..27fdbbbc4f 100644
--- a/ggml/src/ggml-opencl/kernels/rope.cl
+++ b/ggml/src/ggml-opencl/kernels/rope.cl
@@ -76,8 +76,7 @@ kernel void kernel_rope_norm_f32(
         float attn_factor,
         float beta_fast,
         float beta_slow,
-        int n_offs,
-        int inplace
+        int n_offs
 ) {
     src0 = (global void*)((global char*)src0 + offset0);
     src1 = (global int*)((global char*)src1 + offset1);
@@ -115,9 +114,6 @@ kernel void kernel_rope_norm_f32(
             dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
             dst_data[1] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
         } else {
-            if (inplace) {
-                continue;
-            }
             global float * src      = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
             global float * dst_data = (global float *)((global char *)  dst + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 
@@ -161,8 +157,7 @@ kernel void kernel_rope_norm_f16(
         float attn_factor,
         float beta_fast,
         float beta_slow,
-        int n_offs,
-        int inplace
+        int n_offs
 ) {
     src0 = (global void*)((global char*)src0 + offset0);
     src1 = (global int*)((global char*)src1 + offset1);
@@ -200,9 +195,6 @@ kernel void kernel_rope_norm_f16(
             dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
             dst_data[1] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
         } else {
-            if (inplace) {
-                continue;
-            }
             global half * src      = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
             global half * dst_data = (global half *)((global char *)  dst + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 
@@ -246,8 +238,7 @@ kernel void kernel_rope_neox_f32(
         float attn_factor,
         float beta_fast,
         float beta_slow,
-        int n_offs,
-        int inplace
+        int n_offs
 ) {
     src0 = (global void*)((global char*)src0 + offset0);
     src1 = (global int*)((global char*)src1 + offset1);
@@ -285,9 +276,6 @@ kernel void kernel_rope_neox_f32(
             dst_data[0]        = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
             dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
         } else {
-            if (inplace) {
-                continue;
-            }
             global float * const src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
             global float * dst_data  = (global float *)((global char *)  dst + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 
@@ -331,8 +319,7 @@ kernel void kernel_rope_neox_f16(
         float attn_factor,
         float beta_fast,
         float beta_slow,
-        int n_offs,
-        int inplace
+        int n_offs
 ) {
     src0 = (global void*)((global char*)src0 + offset0);
     src1 = (global int*)((global char*)src1 + offset1);
@@ -370,9 +357,6 @@ kernel void kernel_rope_neox_f16(
             dst_data[0]        = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
             dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
         } else {
-            if (inplace) {
-                continue;
-            }
             global half * const src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
             global half * dst_data  = (global half *)((global char *)  dst + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 
@@ -418,8 +402,7 @@ kernel void kernel_rope_multi_f32(
         float beta_slow,
         int4 sections,
         int  is_imrope,
-        int n_offs,
-        int inplace
+        int n_offs
 ) {
     src0 = (global void*)((global char*)src0 + offset0);
     src1 = (global int*)((global char*)src1 + offset1);
@@ -487,9 +470,6 @@ kernel void kernel_rope_multi_f32(
             dst_data[0]        = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
             dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
         } else {
-            if (inplace) {
-                continue;
-            }
             global float * const src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
             global float * dst_data  = (global float *)((global char *)  dst + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 
@@ -535,8 +515,7 @@ kernel void kernel_rope_multi_f16(
         float beta_slow,
         int4 sections,
         int  is_imrope,
-        int n_offs,
-        int inplace
+        int n_offs
 ) {
     src0 = (global void*)((global char*)src0 + offset0);
     src1 = (global int*)((global char*)src1 + offset1);
@@ -604,9 +583,6 @@ kernel void kernel_rope_multi_f16(
             dst_data[0]        = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
             dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
         } else {
-            if (inplace) {
-                continue;
-            }
             global half * const src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
             global half * dst_data  = (global half *)((global char *)  dst + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
 
diff --git a/ggml/src/ggml-sycl/rope.cpp b/ggml/src/ggml-sycl/rope.cpp
index 1345cdb393..b6d22559d1 100644
--- a/ggml/src/ggml-sycl/rope.cpp
+++ b/ggml/src/ggml-sycl/rope.cpp
@@ -45,8 +45,7 @@ static void rope_norm(const T *x, D *dst, const int ne00, const int ne01,
                       const float freq_scale, const float ext_factor,
                       const float attn_factor, const rope_corr_dims corr_dims,
                       const float theta_scale, const float *freq_factors,
-                      const int64_t *row_indices, const int set_rows_stride,
-                      const bool inplace) {
+                      const int64_t *row_indices, const int set_rows_stride) {
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     const int i0 = 2 * (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
                         item_ct1.get_local_id(1));
@@ -80,9 +79,6 @@ static void rope_norm(const T *x, D *dst, const int ne00, const int ne01,
         }
     };
     if (i0 < n_offs || i0 >= n_offs + n_dims) {
-        if (inplace) {
-            return;
-        }
         store_coaelsced(x[ix + 0], x[ix + 1]);
         return;
     }
@@ -114,8 +110,7 @@ static void rope_neox(const T *x, D *dst, const int ne00, const int ne01,
                       const float freq_scale, const float ext_factor,
                       const float attn_factor, const rope_corr_dims corr_dims,
                       const float theta_scale, const float *freq_factors,
-                      const int64_t *row_indices, const int set_rows_stride,
-                      const bool inplace) {
+                      const int64_t *row_indices, const int set_rows_stride) {
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     const int i0 = 2 * (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
                         item_ct1.get_local_id(1));
@@ -140,9 +135,6 @@ static void rope_neox(const T *x, D *dst, const int ne00, const int ne01,
     }
 
     if (i0 < n_offs || i0 >= n_offs + n_dims) {
-        if (inplace) {
-            return;
-        }
         dst[idst + i0 / 2 + 0] = ggml_sycl_cast<D>(x[ix + i0 / 2 + 0]);
         dst[idst + i0 / 2 + 1] = ggml_sycl_cast<D>(x[ix + i0 / 2 + 1]);
 
@@ -177,8 +169,7 @@ static void rope_multi(const T *x, T *dst, const int ne00, const int ne01,
                        const float freq_scale, const float ext_factor,
                        const float attn_factor, const rope_corr_dims corr_dims,
                        const float theta_scale, const float *freq_factors,
-                       const mrope_sections sections, const bool is_imrope,
-                       const bool inplace) {
+                       const mrope_sections sections, const bool is_imrope) {
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     const int i0 = 2 * (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
                         item_ct1.get_local_id(1));
@@ -198,9 +189,6 @@ static void rope_multi(const T *x, T *dst, const int ne00, const int ne01,
     const int ix = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03;
 
     if (i0 < n_offs || i0 >= n_offs + n_dims) {
-        if (inplace) {
-            return;
-        }
         dst[idst + i0 / 2 + 0] = x[ix + i0 / 2 + 0];
         dst[idst + i0 / 2 + 1] = x[ix + i0 / 2 + 1];
 
@@ -317,7 +305,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
                const float freq_base, const float ext_factor,
                const float attn_factor, const rope_corr_dims corr_dims,
                const float *freq_factors, const int64_t *row_indices,
-               const int set_rows_stride, const bool inplace, dpct::queue_ptr stream) {
+               const int set_rows_stride, dpct::queue_ptr stream) {
     GGML_ASSERT(ne00 % 2 == 0);
     const dpct::dim3 block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
     const int n_blocks_x =
@@ -334,8 +322,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
                 rope_norm<forward, false>(
                     x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
                     n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
-                    theta_scale, freq_factors, row_indices, set_rows_stride,
-                    inplace);
+                    theta_scale, freq_factors, row_indices, set_rows_stride);
             });
     } else {
         stream->parallel_for(
@@ -345,8 +332,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
                 rope_norm<forward, true>(
                     x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
                     n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
-                    theta_scale, freq_factors, row_indices, set_rows_stride,
-                    inplace);
+                    theta_scale, freq_factors, row_indices, set_rows_stride);
             });
     }
 }
@@ -360,7 +346,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
                const float freq_base, const float ext_factor,
                const float attn_factor, const rope_corr_dims corr_dims,
                const float *freq_factors, const int64_t *row_indices,
-               const int set_rows_stride, const bool inplace, dpct::queue_ptr stream) {
+               const int set_rows_stride, dpct::queue_ptr stream) {
     GGML_ASSERT(ne00 % 2 == 0);
     const dpct::dim3 block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
     const int n_blocks_x =
@@ -377,8 +363,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
                 rope_neox<forward, false>(
                     x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
                     n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
-                    theta_scale, freq_factors, row_indices, set_rows_stride,
-                    inplace);
+                    theta_scale, freq_factors, row_indices, set_rows_stride);
             });
     } else {
         stream->parallel_for(
@@ -388,8 +373,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
                 rope_neox<forward, true>(
                     x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
                     n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
-                    theta_scale, freq_factors, row_indices, set_rows_stride,
-                    inplace);
+                    theta_scale, freq_factors, row_indices, set_rows_stride);
             });
     }
 }
@@ -403,7 +387,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
                 const float freq_base, const float ext_factor,
                 const float attn_factor, const rope_corr_dims corr_dims,
                 const float *freq_factors, const mrope_sections sections,
-                const bool is_imrope, const bool inplace, dpct::queue_ptr stream) {
+                const bool is_imrope, dpct::queue_ptr stream) {
     GGML_ASSERT(ne00 % 2 == 0);
     const dpct::dim3 block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
     const int n_blocks_x =
@@ -420,7 +404,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
                 rope_multi<forward, false, T>(
                     x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
                     n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
-                    theta_scale, freq_factors, sections, is_imrope, inplace);
+                    theta_scale, freq_factors, sections, is_imrope);
             });
     } else {
         stream->parallel_for(
@@ -430,7 +414,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
                 rope_multi<forward, true, T>(
                     x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
                     n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
-                    theta_scale, freq_factors, sections, is_imrope, inplace);
+                    theta_scale, freq_factors, sections, is_imrope);
             });
     }
 }
@@ -524,9 +508,6 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
     const int n_offs = ((int32_t *)dst->op_params)[15];
     mrope_sections sections;
 
-    // when dst aliases src0, the channels outside the rotated window already hold the correct data
-    const bool inplace = dst_d == src0->data;
-
     float freq_base;
     float freq_scale;
     float ext_factor;
@@ -576,19 +557,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
                 (const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01,
                 s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
                 ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                set_rows_stride, inplace, stream);
+                set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
             rope_neox_sycl<forward, float, sycl::half>(
                 (const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02,
                 s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
                 freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
-                row_indices, set_rows_stride, inplace, stream);
+                row_indices, set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
             rope_neox_sycl<forward, sycl::half, sycl::half>(
                 (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
                 ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
                 freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
-                row_indices, set_rows_stride, inplace, stream);
+                row_indices, set_rows_stride, stream);
         } else {
             GGML_ABORT("Fatal error: Tensor type unsupported!");
         }
@@ -599,13 +580,13 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
                                      ne00, ne01, ne02, s01, s02, s03, s1, s2,
                                      s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
                                      ext_factor, attn_factor, corr_dims,
-                                     freq_factors, sections, is_imrope, inplace, stream);
+                                     freq_factors, sections, is_imrope, stream);
         } else if (src0->type == GGML_TYPE_F16) {
             rope_multi_sycl<forward>(
                 (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
                 ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
                 freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
-                sections, is_imrope, inplace, stream);
+                sections, is_imrope, stream);
         } else {
             GGML_ABORT("Fatal error: Tensor type unsupported!");
         }
@@ -633,19 +614,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
                 (const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01,
                 s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
                 ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
-                set_rows_stride, inplace, stream);
+                set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
             rope_norm_sycl<forward, float, sycl::half>(
                 (const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02,
                 s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
                 freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
-                row_indices, set_rows_stride, inplace, stream);
+                row_indices, set_rows_stride, stream);
         } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
             rope_norm_sycl<forward, sycl::half, sycl::half>(
                 (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
                 ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
                 freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
-                row_indices, set_rows_stride, inplace, stream);
+                row_indices, set_rows_stride, stream);
         } else {
             GGML_ABORT("Fatal error: Tensor type unsupported!");
         }
diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl b/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl
index 0afa211044..6ff53088c4 100644
--- a/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl
+++ b/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl
@@ -168,14 +168,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let i_dst_row = params.offset_dst + i3 * params.stride_dst3 + i2 * params.stride_dst2 + i1 * params.stride_dst1;
 
     if ((i0 < params.n_offs || i0 >= params.n_offs + params.n_dims) && !is_vision) {
-#ifdef INPLACE
-        return;
-#else
         let i_src = i_src_row + i0;
         let i_dst = i_dst_row + i0;
         rotate(i_dst, i_dst + 1, f32(src0[i_src]), f32(src0[i_src + 1]));
         return;
-#endif
     }
 
     let iw = i0 - params.n_offs; // relative idx