mmcv中的focal loss CUDA实现注释

template <typename T>
__global__ void sigmoid_focal_loss_forward_cuda_kernel(
    const int nthreads, const T* input, const int64_t* target, const T* weight,
    T* output, const T gamma, const T alpha, const int num_classes) {
// nthreads就是outputsize，等于num_total_anchors*num_classes
// const T*就是tensor的连续内存首地址
// input的连续内存长度为num_total_anchors*num_classes，target的连续内存长度为num_total_anchors。
  CUDA_1D_KERNEL_LOOP(index, nthreads) {
    // index等于blockIdx.x * blockDim.x + threadIdx.x，即线程索引

    // 因为index就是对应tensor(num_total_anchors,num_classes)的一个元素
    int n = index / num_classes; // n就是预测框的序号
    // c的范围是[0,num_classes-1],就是预测框预测的第几个类别序号
    // 下面的input[index]就是预测框对第c类的预测得分
    int c = index % num_classes; // c的范围是[0,num_classes-1],就是预测框预测的第几个类别序号

    // 第n个预测框对应的gt类别label
    int64_t t = target[n]; // 获得anchor n 的target label
    T flag_p = (t == c); // 表示正样本
    T flag_n = (t != c); // 表示负样本

    // p = sigmoid(x) = 1. / 1. + expf(-x)
    T p = (T)1. / ((T)1. + expf(-input[index]));

    // (1 - p)**gamma * log(p) 正样本的focal loss权重
    T term_p = pow(((T)1. - p), gamma) * log(max(p, (T)FLT_MIN));
    // p**gamma * log(1 - p) 负样本的focal loss权重
    T term_n = pow(p, gamma) * log(max((T)1. - p, (T)FLT_MIN));

    output[index] = (T)0.; //计算结果放到output tensor中
    output[index] += -flag_p * alpha * term_p;
    output[index] += -flag_n * ((T)1. - alpha) * term_n;
    if (weight != NULL) {
      output[index] *= weight[t];
    }
  }
}