 }
 
 void CUDADevice::pick_fw_impl(
-    const Tensor &x, unsigned dim,
-    const std::vector<unsigned> &ids, Tensor &y) {
+    const Tensor &x, unsigned dim, const std::vector<unsigned> &ids,
+    Tensor &y) {
   const unsigned wy = y.shape().lower_volume(dim);
   const unsigned sy = y.shape().volume();
   const unsigned g1 = GRID_SIZE(sy, dim1_x_);
   const unsigned bs = y.shape().batch();
 
   CUDA_CALL(::cudaSetDevice(dev_id_));
-  ::cudaMemcpy(
-      ids_ptr_.get(), ids.data(), sizeof(unsigned) * ids.size(),
-      cudaMemcpyHostToDevice);
+  CUDA_CALL(::cudaMemcpy(
+        ids_ptr_.get(), ids.data(), sizeof(unsigned) * ids.size(),
+        cudaMemcpyHostToDevice));
   ::pick_fw_dev<<<dim3(g1, bs), dim1_x_>>>(
       CDATA(x), static_cast<const unsigned *>(ids_ptr_.get()),
       wy * x.shape()[dim], wy,

}

__global__ void pick_bw_dev(
    const float *pgy, unsigned wx, unsigned wy, unsigned repeat, float *pgx) {
  const unsigned i = IDX;
  if (i < wy * repeat) {
    ::atomicAdd(pgx + (i / wy) * wx + (i % wy), pgy[i]);
  }
}

__global__ void slice_bw_dev(

void CUDADevice::pick_bw_impl(
    const Tensor &gy, unsigned dim, const std::vector<unsigned>& ids,
    Tensor &gx) {
  const unsigned wy = gy.shape().lower_volume(dim);
  const unsigned sy = gy.shape().volume();
  const unsigned g1 = GRID_SIZE(sy, dim1_x_);
  const unsigned bs = gy.shape().batch();

  CUDA_CALL(::cudaSetDevice(dev_id_));
  CUDA_CALL(::cudaMemcpy(
        ids_ptr_.get(), ids.data(), sizeof(unsigned) * ids.size(),
        cudaMemcpyHostToDevice));
  ::pick_bw_dev<<<dim3(g1, bs), dim1_x_>>>(
      CDATA(gy), static_cast<const unsigned *>(ids_ptr_.get()),
      wy *gx.shape()[dim], wy,
      gx.shape().has_batch() * gx.shape().volume(), ids.size() > 1, sy,
      DATA(gx));
}

void CUDADevice::slice_bw_impl(
