Thank you Alex for your responses, it gave me an idea on how to debug it and at the end I managed to do it. Given that maybe it could be instructional for someone else,  here it is how it went down.
Instead of going in the setup.py direction, I made a CMakeLists.txt file which looks like this:
cmake_minimum_required(VERSION 3.10 FATAL_ERROR)
# here I decide the name of the library
project(nms_cmake LANGUAGES CXX CUDA)
set(CMAKE_CXX_STANDARD 14)

# I make sure to point to the correct libraries
set(CMAKE_PREFIX_PATH /home/michele/pytorch_cpp/libtorch-shared-with-deps-1.9.1+cu102/libtorch)

# I make sure that it has all the python and torch directories available
set(TORCH_INCLUDE_DIRS
   /home/michele/anaconda3/envs/with_cudatoolkit/include/
   /home/michele/anaconda3/envs/with_cudatoolkit/include/torch/csrc/api/include)

# I make sure that both Python and Torch are found, also all the py-dev stuff
find_package(Torch REQUIRED)
find_package(Python REQUIRED COMPONENTS Development)

# add all the sources to the library
# List all your code files here
add_library(nms_cmake SHARED
  main.cu
)

target_link_libraries(nms_cmake "${TORCH_LIBRARIES}")

# To really make sure that it is using the correct language and gcc to make stuff.
target_compile_options(nms_cmake PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-ccbin g++-7>)


main.cu  at this stage was just a combination of the previous nms_kernel.cu and nms_cuda.cpp scripts and full of unknown syntactical errors.
After creating the build fonder and launching:
 cmake -DCMAKE_PREFIX_PATH=$PWD/home/michele/pytorch_cpp/libtorch-shared-with-deps-1.9.1+cu102/libtorch ..

and

make

I could see may errors in the file and after some debugging it ended up like this:
#include <c10/cuda/CUDAException.h>
#include <torch/torch.h>
#include <torch/all.h>
#include <torch/python.h>
#include <torch/extension.h>

using namespace at;

#define DIVUP(m,n) ((m) / (n) + ((m) % (n) > 0))
int const threadsPerBlock = sizeof(unsigned long long) * 8;


__device__ inline float devIoU(float const * const a, float const * const b) {
  float left = fmaxf(a[0], b[0]), right = fminf(a[2], b[2]);
  float top = fmaxf(a[1], b[1]), bottom = fminf(a[3], b[3]);
  float front = fmaxf(a[4], b[4]), back = fminf(a[5], b[5]);

  float width = fmaxf(right - left + 1, 0.f), height = fmaxf(bottom - top + 1, 0.f), depth = fmaxf(back - front + 1, 0.f);
  float interS = width * height * depth;
  float Sa = (a[2] - a[0] + 1) * (a[3] - a[1] + 1) * (a[5] - a[4] + 1);
  float Sb = (b[2] - b[0] + 1) * (b[3] - b[1] + 1) * (b[5] - b[4] + 1);
  //printf("IoU 3D %f \n", interS / (Sa + Sb - interS));

  return interS / (Sa + Sb - interS);
}

__global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
                           const float *dev_boxes, unsigned long long *dev_mask) {
  const int row_start = blockIdx.y;
  const int col_start = blockIdx.x;

  // if (row_start > col_start) return;

  const int row_size =
        fminf(n_boxes - row_start * threadsPerBlock, threadsPerBlock);
  const int col_size =
        fminf(n_boxes - col_start * threadsPerBlock, threadsPerBlock);

  __shared__ float block_boxes[threadsPerBlock * 7];
  if (threadIdx.x < col_size) {
    block_boxes[threadIdx.x * 7 + 0] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 0];
    block_boxes[threadIdx.x * 7 + 1] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 1];
    block_boxes[threadIdx.x * 7 + 2] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 2];
    block_boxes[threadIdx.x * 7 + 3] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 3];
    block_boxes[threadIdx.x * 7 + 4] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 4];
    block_boxes[threadIdx.x * 7 + 5] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 5];
    block_boxes[threadIdx.x * 7 + 6] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 7 + 6];
  }
  __syncthreads();

  if (threadIdx.x < row_size) {
    const int cur_box_idx = threadsPerBlock * row_start + threadIdx.x;
    const float *cur_box = dev_boxes + cur_box_idx * 7;
    int i = 0;
    unsigned long long t = 0;
    int start = 0;
    if (row_start == col_start) {
      start = threadIdx.x + 1;
    }
    for (i = start; i < col_size; i++) {
      if (devIoU(cur_box, block_boxes + i * 7) > nms_overlap_thresh) {
        t |= 1ULL << i;
      }
    }
    const int col_blocks = DIVUP(n_boxes, threadsPerBlock);
    dev_mask[cur_box_idx * col_blocks + col_start] = t;
  }
}


void _nms(int boxes_num, float * boxes_dev,
          unsigned long long * mask_dev, float nms_overlap_thresh) {


  dim3 blocks(DIVUP(boxes_num, threadsPerBlock),
              DIVUP(boxes_num, threadsPerBlock));
  dim3 threads(threadsPerBlock);
  nms_kernel<<<blocks, threads>>>(boxes_num,
                                  nms_overlap_thresh,
                                  boxes_dev,
                                  mask_dev);
}


int64_t gpu_nms(at::Tensor keep, at::Tensor boxes, at::Tensor num_out,  double nms_overlap_thresh){
    // params boxes: (N, 7) [x, y, z, dx, dy, dz, heading]
    // params keep: (N)

    int boxes_num = boxes.size(0);
    float * boxes_data = boxes.data_ptr<float>();
    long * keep_data = keep.data_ptr<long>();

    const int col_blocks = DIVUP(boxes_num, threadsPerBlock);

    //unsigned long long mask_data[boxes_num * col_blocks];
    unsigned long long *mask_data = NULL;
    cudaMalloc((void**)&mask_data, boxes_num * col_blocks * sizeof(unsigned long long));
    _nms(boxes_num, boxes_data, mask_data, nms_overlap_thresh);

    //unsigned long long mask_cpu[boxes_num * col_blocks];
    //unsigned long long * mask_cpu = new unsigned long long [boxes_num * col_blocks];
    std::vector<unsigned long long> mask_cpu(boxes_num * col_blocks);

    //printf("boxes_num=%d, col_blocks=%d\n", boxes_num, col_blocks);
    cudaMemcpy(&mask_cpu[0], mask_data, boxes_num * col_blocks * sizeof(unsigned long long),
                           cudaMemcpyDeviceToHost);

    cudaFree(mask_data);

    unsigned long long remv_cpu[col_blocks];
    cudaMemset(remv_cpu, 0, col_blocks * sizeof(unsigned long long));

    int num_to_keep = 0;
    for (int i = 0; i < boxes_num; i++){
        int nblock = i / threadsPerBlock;
        int inblock = i % threadsPerBlock;

        if (!(remv_cpu[nblock] & (1ULL << inblock))){
            keep_data[num_to_keep++] = i;
            unsigned long long *p = &mask_cpu[0] + i * col_blocks;
            for (int j = nblock; j < col_blocks; j++){
                remv_cpu[j] |= p[j];
            }
        }
    }
    if ( cudaSuccess != cudaGetLastError() ) printf( "Error!\n" );

    return num_to_keep;
}

TORCH_LIBRARY(pytorch_cmake_example, m) {
  m.def("gpu_nms(Tensor keep, Tensor boxes, Tensor num_out, float nms_overlap_thresh) -> int num_to_keep");
  m.impl("gpu_nms", c10::DispatchKey::CUDA, TORCH_FN(gpu_nms));
  //c10::DispatchKey::CPU is also an option
}

To load the library you can then import it in a python file as have a possibly working 3D NMS as follows:
import torch
torch.ops.load_library("build/libnms_cmake.so")

def nms_gpu(dets, thresh):
  """
  dets has to be a tensor
  """

  scores = dets[:, -1]
  order = scores.sort(0, descending=True)[1]
  dets = dets[order].contiguous()

  keep = torch.LongTensor(dets.size(0))
  num_out = torch.LongTensor(1)
  torch.ops.libnms_cmake.gpu_nms(keep, dets, num_out, thresh)
  return order[keep[:num_out[0]].cuda()].contiguous()

Thanks for the help.