     for(uint i = 0; i < xs.size(); i++) {
         Shape xs_shape{xs[i].shape()[0], xs[i].shape()[1], 1};
         Shape ys_shape{ys[i].shape()[0], ys[i].shape()[1], 1};
-        xs[i] = xs[i].AsType(Dtype::kFloat32);
-        ys[i] = ys[i].AsType(Dtype::kFloat32);
-        dys[i] = dys[i].AsType(Dtype::kFloat32);
-        xsDesc.push_back(CudnnTensorDescriptor(AsContiguous(xs[i]).Reshape(xs_shape)));
+        xs[i] = AsContiguous(xs[i].AsType(Dtype::kFloat32));
+        ys[i] = AsContiguous(ys[i].AsType(Dtype::kFloat32));
+        dys[i] = AsContiguous(dys[i].AsType(Dtype::kFloat32));
+        xsDesc.push_back(CudnnTensorDescriptor(xs[i].Reshape(xs_shape)));
         xs_desc[i] = *xsDesc[i];
         dxs.push_back(Empty(xs[i].shape(), xs[i].dtype(), xs[i].device()));
-        dxsDesc.push_back(CudnnTensorDescriptor(AsContiguous(dxs[i]).Reshape(xs_shape)));
+        dxsDesc.push_back(CudnnTensorDescriptor(dxs[i].Reshape(xs_shape)));
         dxs_desc[i] = *dxsDesc[i];
-        ysDesc.push_back(CudnnTensorDescriptor(AsContiguous(ys[i]).Reshape(ys_shape)));
+        ysDesc.push_back(CudnnTensorDescriptor(ys[i].Reshape(ys_shape)));
         ys_desc[i] = *ysDesc[i];
-        dysDesc.push_back(CudnnTensorDescriptor(AsContiguous(dys[i]).Reshape(ys_shape)));
+        dysDesc.push_back(CudnnTensorDescriptor(dys[i].Reshape(ys_shape)));
         dys_desc[i] = *dysDesc[i];
     }
     Array dx = Concatenate(dxs, 0);

                &nbDims,
                filterDimA
                );
                ws[index][lin_layer_id] = ws[index][lin_layer_id].AsType(Dtype::kFloat32);
                initGPUData(m_offset, filterDimA[0] * filterDimA[1] * filterDimA[2], (float*)internal::GetRawOffsetData(AsContiguous(ws[index][lin_layer_id])));
                cudnnDestroyFilterDescriptor(linlayermatdesc);
                


    );

    Array dw = Empty({(int)weight_size / 4, 1, 1}, hx.dtype(), hx.device());
    CudnnFilterDescriptor dwDesc{dw};
    
    handle.Call(
        cudnnRNNBackwardWeights,
        rnn_desc,
        xs.size(),
        xs_desc,
        internal::GetRawOffsetData(x),
        *hx_desc,
        internal::GetRawOffsetData(hx),
        ys_desc,
        internal::GetRawOffsetData(y),
        workspace,
        workSize,
        *dwDesc,
