CUDA and PyTorch Version

[iris112358]

Hi, thanks for the code. I tried to run the code but had the following error

subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1.

It seems the CUDA version does not match. Could you please let me know the CUDA and PyTorch version you used to run the code?

[Kristen-Z]

Hi, I encountered the same issue. Did you solve this one?

[Pangolin112]

In my case, some changes to the code in .cu file like this would work for cuda 11.8 + pytorch 2.1.2:

#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <tuple>
// …existing includes…

#define THREAD_COUNT 1024

// ---------------------- in your computeHistogram kernel ----------------------
__global__ void computeHistogram(float *tensor, float *histogram, float *minv, float *maxv, unsigned int channels, unsigned int tensorSize, unsigned int nBins) {
  unsigned int index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index < channels * tensorSize) {
    unsigned int channel = index / tensorSize;
    float value = (tensor[index] - minv[channel]) / (maxv[channel] - minv[channel]) * float(nBins);
    // replace ambiguous min(...) with a ternary
    unsigned int uv = (unsigned int)(value);
    unsigned int bin = uv < (nBins - 1) ? uv : (nBins - 1);
    atomicAdd(histogram + (channel * nBins) + bin, 1);
  }
}

// ---------------------- in your buildSortedLinkmap kernel ----------------------
__global__ void buildSortedLinkmap( /*…*/ ) {
  unsigned int index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index < channels * tensorSize) {
    // …compute value as before…
    float value = (tensor[index] - minv[channel]) / (maxv[channel] - minv[channel]) * float(nBins);
    unsigned int uv = (unsigned int)(value);
    unsigned int binIndex = uv < (nBins - 1) ? uv : (nBins - 1);
    // …rest unchanged…
  }
}

// ---------------------- in computeHistogram(…) host function ----------------------
at::Tensor computeHistogram(at::Tensor const &t, unsigned int numBins) {
  at::Tensor unsqueezed = t.cuda();
  if (unsqueezed.ndimension() == 1)         unsqueezed.unsqueeze_(0);
  if (unsqueezed.ndimension() > 2)          unsqueezed = unsqueezed.view({unsqueezed.size(0), -1});
  unsigned int c = unsqueezed.size(0), n = unsqueezed.numel()/c;

  // grab mins/maxs via C++ front-end
  auto mins = std::get<0>(unsqueezed.min(1, true)).cuda();
  auto maxs = std::get<0>(unsqueezed.max(1, true)).cuda();

  at::Tensor h = at::zeros({int(c), int(numBins)}, unsqueezed.options()).cuda();
  computeHistogram<<< (c*n)/THREAD_COUNT + 1, THREAD_COUNT >>>(
    unsqueezed.data_ptr<float>(),
    h.data_ptr<float>(),
    mins.data_ptr<float>(),
    maxs.data_ptr<float>(),
    c, n, numBins
  );
  return h;
}

// ---------------------- in matchHistogram(…) before buildSortedLinkmap <<<…>>> ----------------------
void matchHistogram(at::Tensor &featureMaps, at::Tensor &targetHistogram) {
  // …existing code up to “unsqueezed” …
  at::Tensor unsqueezed = featureMaps;
  if (unsqueezed.ndimension() == 1)         unsqueezed.unsqueeze_(0);
  if (unsqueezed.ndimension() > 2)          unsqueezed = unsqueezed.view({unsqueezed.size(0), -1});
  unsigned int c = unsqueezed.size(0), n = unsqueezed.numel()/c, nBins = targetHistogram.size(1);

  // recompute mins/maxs here too
  auto mins = std::get<0>(unsqueezed.min(1, true)).cuda();
  auto maxs = std::get<0>(unsqueezed.max(1, true)).cuda();

  // …compute linkMap, etc. …
  buildSortedLinkmap<<< (c*n)/THREAD_COUNT + 1, THREAD_COUNT >>>(
    featureMaps.data_ptr<float>(),
    linkMap,
    featuresHistogram.data_ptr<float>(),
    localIndexes,
    randomIndices[featureMaps.numel()].data_ptr<long>(),
    mins.data_ptr<float>(),
    maxs.data_ptr<float>(),
    c, n, nBins
  );
  // …rest unchanged…
}