Sample 2 - Baseline auction bidding kernel — first numbers

The first CUDA version assigns one thread per unmatched vertex in U. Each thread scans its adjacency slice, finds the best and second-best net value, and emits a bid.

__global__ void bid_kernel(const int* __restrict__ rowPtr,
                           const int* __restrict__ colIdx,
                           const float* __restrict__ w,
                           const float* __restrict__ price,
                           const int* __restrict__ active,
                           int nActive, float eps,
                           int* bidTarget, float* bidValue)
{
  int t = blockIdx.x * blockDim.x + threadIdx.x;
  if (t >= nActive) return;
  int u = active[t];

  float best = -INFINITY, second = -INFINITY;
  int bestV = -1;
  for (int e = rowPtr[u]; e < rowPtr[u + 1]; ++e) {
    float val = w[e] - price[colIdx[e]];
    if (val > best) { second = best; best = val; bestV = colIdx[e]; }
    else if (val > second) { second = val; }
  }
  bidTarget[t] = bestV;
  bidValue[t]  = best - second + eps;
}

Profile

Metric	Naive	After warp-per-vertex
Achieved occupancy	22%	61%
Global load efficiency	18%	74%
Kernel time (1e5 vertices)	41.3 ms	9.8 ms

The naive version is load-imbalanced: one thread walks an entire adjacency list, so a hub vertex with 4000 neighbours serialises its whole warp. Switching to one warp per vertex with a __shfl_down_sync reduction for the best/second-best pair fixed both the divergence and the coalescing, since consecutive lanes now read consecutive colIdx entries.

Takeaway: on CSR traversals, the mapping granularity matters more than any micro-optimisation inside the loop.

Still open: the assignment phase is atomics-bound and does not scale past ~4 SMs worth of contention. Next step is a segmented sort of bids by target vertex.