sum_axis and min cudaMalloc/cudaFree on every call, which syncs the device

[CanReader]

cuda/kernels.cu:709 and :773:

CUDA_CHECK(cudaMalloc(&d_shape, ndim * sizeof(int)));
...
cudaFree(d_shape);

cudaFree is synchronising, it drains the device. So every sum_axis (which
means every broadcast backward, since reduce_to goes through it) and every min
costs a full device sync plus a driver round trip. That's precisely the cost the
CudaBuffer caching allocator was built to avoid.

fastnn_cuda_min also allocates a full n element scratch buffer per call.

Pass the shape through a by value struct the way permute_nd already does (up to
8 dims), and put scratch through CudaBuffer.

Minor, in the same function: int axis_size = 1; is dead and the first two
assignments to output_size are never read.