I did a few more tests and debugging via printing in the kernels, and it seems for some combinations of input height/width and padding (moreso lack of padding), during the correlation forward and backward, it tries to access outside of the image i.e. index -1.
I’m not entirely sure who is the original author of the correlation kernel, I think the earliest I’ve noticed is nvidia’s PWC-net. Its 404’ing because I figured this out, made the fix and noted that I need at least some padding (and to increase it if I have this problem again) and then merged it back into my master branch.
The bad accesses occur where the boundary checks are (unless otherwise stated in a comment). If I can figure out a more elegant solution, I could make one and submit a PR? I believe there should probably be a way of calculating required minimal required padding.

  
      github.com
  
  
    5had3z/stereo-to-all/blob/master/nnet_training/correlation_package/correlation_cuda_kernel.cu 1
#include "correlation_cuda_kernel.cuh"

#define CUDA_NUM_THREADS 1024
#define THREADS_PER_BLOCK 32

#include <ATen/Dispatch.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

// using at::Half;
#define TensorAcc4R torch::PackedTensorAccessor32<scalar_t,4,torch::RestrictPtrTraits>

template <typename scalar_t>
__global__ void channels_first(const TensorAcc4R input, TensorAcc4R rinput,
	const int channels, const int height, const int width, const int pad_size)
{
	// n (batch size), c (num of channels), y (height), x (width)
	const int n = blockIdx.x;
	const int y = blockIdx.y;
	const int x = blockIdx.z;


  This file has been truncated. show original