基于毕昇编译器的softmax异构算子

已于 2023-10-18 09:25:08 修改
阅读量207 收藏
点赞数
文章标签： c++ gpu算力 深度学习 人工智能
于 2023-08-20 20:01:59 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/Deleter357/article/details/132395300
版权

为了良好的阅读体验，请移步笔者博客博客链接

分析算子

Softmax是非常常见的激活函数，此处不过多赘述。观察其公式：
Softmax ( x ) = e x i ∑ i e x i \text{Softmax}(x) = \frac{e^{x_i}}{\sum_ie^{x_i}} Softmax(x)=∑i​exi​exi​​
首先应该注意到的是分母中的求和，因为向量内元素求和本身就是一个不太适合于向量化的操作，初步分析可知，此处的求和可能是性能的瓶颈。而分子除以分母的运算可以在得到分母（即向量元素之和）的情况下，自然而然的向量化运算。

初步方案（方案一）

  • 向量化方案：因为向量除法是可以自然而然向量化的操作，故暂时先将分母的求和操作放在Host端进行。
  • 分核方案：本算子是在昇腾910B上进行开发的，由于数据对其以及访存粒度的限制，暂时令一个核处理640个元素。

请添加图片描述

初步方案实现

先用标准C++实现一个Softmax算子，作为功能验证和性能对比的基准。

using data_t = float;

std::vector<data_t> softmax(std::vector<data_t> input) {
  data_t sum = 0.0;
  for (auto x : input) {
    sum += expf(x);
  }
  std::vector<data_t> res;
  for (auto x : input) {
    res.push_back(expf(x) / sum);
  }
  return res;
}

接下来开始实现异构算子的逻辑。

std::vector<data_t> ascend_softmax(std::vector<data_t> input) {
  // 首先拿到输入向量的大小
  std::size_t input_sz = input.size();
  // 计算出总的数据量，单位是字节
  std::size_t byte_count = input_sz * sizeof(data_t);

  // 如果输入向量的大小不足一个block，则直接调用Host算子逻辑
  if (byte_count < 32)
    return softmax(input);

  sycl::queue Q(sycl::ascend_selector{});

  // 申请GM上的内存，分别存储输入和结果
  auto input_buf = sycl::malloc_device<data_t>(input_sz, Q);
  auto res_buf = sycl::malloc_device<data_t>(input_sz, Q);

  // host -> GM 的内存搬移
  Q.memcpy(input_buf, input.data(), byte_count);

  // 指定每个核处理的元素个数
  const std::size_t elem_per_group = 640;
  // 计算尾块中的元素个数
  const std::size_t tail_elem_count = input_sz % elem_per_group;
  // 逻辑核的数量，若尾块中存在元素，则多开一个逻辑核
  const std::size_t group_num = (tail_elem_count > 0)
                                    ? ((input_sz / elem_per_group) + 1)
                                    : (input_sz / elem_per_group);

  // 求和计算暂时由Host完成
  data_t sum = 0.0;
  for (auto x : input) {
    sum += expf(x);
  }

  Q.launch<class Softmax>(group_num, [=](sycl::group<1> group) {
    // UB内存申请，分别为输入向量、指数计算结果向量、分母向量、结果向量
    bisheng::vector<data_t, elem_per_group> input_vec;
    bisheng::vector<data_t, elem_per_group> exp_res_vec;
    bisheng::vector<data_t, elem_per_group> divisor_vec(sum);
    bisheng::vector<data_t, elem_per_group> res_vec;
    // 获取group id
    std::size_t group_id = group.get_group_id();

    // GM -> UB 的内存搬移
    input_vec.load(
        sycl::global_ptr<data_t>(input_buf + group_id * elem_per_group).get(), elem_per_group);

    if (tail_elem_count > 0 && group_id == group_num - 1) {
      // 本分支处理存在尾块，且当前是最后一个处理尾块的group
      // 由于尾块大概率是非整block，故采用毕昇变长向量对数据进行操作
      bisheng::vector_view<data_t> input_vec_v(input_vec.data(), tail_elem_count);
      bisheng::vector_view<data_t> exp_res_vec_v(exp_res_vec.data(), tail_elem_count);
      bisheng::vector_view<data_t> divisor_vec_v(divisor_vec.data(), tail_elem_count);
      bisheng::vector_view<data_t> res_vec_v(res_vec.data(), tail_elem_count);

      bisheng::vec_exp(exp_res_vec_v, input_vec_v);
      bisheng::vec_div(res_vec_v, exp_res_vec_v, divisor_vec_v);
    } else {
      // 本分支处理整block的情况
      // 由于指定了每个核处理的元素个数，故此处一定是整block的
      bisheng::vec_exp(exp_res_vec, input_vec);
      bisheng::vec_div(res_vec, exp_res_vec, divisor_vec);
    }

    // UB -> GM 内存搬移
    res_vec.store(
        sycl::global_ptr<data_t>(res_buf + group_id * elem_per_group).get(),
        elem_per_group);
  });

  std::vector<data_t> res(input_sz, 0.0f);
  // GM -> host 内存搬移
  Q.memcpy(res.data(), res_buf, byte_count);
  Q.wait();
  // 释放资源
  sycl::free(input_buf, Q);
  sycl::free(res_buf, Q);

  return res;
}

功能验证

由于之前实现了Host端的逻辑，所以可以用Host端的计算结果来验证异构算子的正确性，测试代码如下。输入向量的数据是从(-1, 1)均的匀分布中取的随机数。

由于设备端会有一定的精度损失，故在测试过程中留有一定的宽容度，相对精度损失在2%以内均认为计算正确。

int main() {
  std::vector<data_t> vec;

  std::random_device rd;
  std::mt19937 gen(rd());
  std::uniform_real_distribution<data_t> urdis(-1, 1);
  for (int i = 0; i < INPUT_COUNT; i++) {
    vec.push_back(urdis(gen));
  }

  std::vector<data_t> host_res = softmax(vec);
  std::vector<data_t> ascend_res = ascend_softmax(vec);

  for (int i = 0; i < host_res.size(); ++i) {
    if (std::fabs(host_res[i] - ascend_res[i]) / host_res[i] > 0.02) {
      std::cout << "Calculation error." << std::endl;
      return EXIT_FAILURE;
    }
  }
  std::cout << "Result correct." << std::endl;

  return EXIT_SUCCESS;
}

性能测试

性能测试使用不同长度的输入向量进行，长度分别为640、6400、64000、640000，数据类型为float。Host端的时间统计使用<time.h>中的结构体，设备端的时间统计利用毕昇C++提供的profiling系列接口统计。仅统计算子计算逻辑部分的时间，并执行5次取平均值后计算加速比。之后的性能测试均为该策略，后面不再赘述。

性能测试结果如下。

测试用例640640064000640000
加速比0.181631.1038322.9683453.784732

可以看到在向量长度达到6400之后才勉强与Host端计算逻辑持平，虽然在长向量的情况下有速度上的提升，但远不是令人满意的效果。

求和方式优化（方案二）

在初步方案中，求和的过程是由Host端完成的，故接下来将求和也尽量的向量化。对于求和无法做到元素级别的并行，故只能将长向量拆分为多个短向量，这些短向量之间的求和操作是并行的，但单个短向量内部的求和，只能是串行的。

由于求和的每一项也需要经过指数运算，故可以把指数运算与求和放在同一个核函数内进行，计算完成后将指数运算结果向量存储下来，则后面做向量除法时就不必再运算一遍了。

总体的异构方案如图所示。

请添加图片描述

方案实现

实现代码如下，只需关注有注释的部分，没有注释的部分与上一方案中的代码相同。

std::vector<float> ascend_softmax(std::vector<float> input) {
  std::size_t input_sz = input.size();
  std::size_t byte_count = input_sz * sizeof(float);

  if (byte_count < 32)
    return softmax(input);

  const std::size_t elem_per_group = 640;
  const std::size_t tail_elem_count = input_sz % elem_per_group;
  const std::size_t group_num = (tail_elem_count > 0)
                                    ? ((input_sz / elem_per_group) + 1)
                                    : (input_sz / elem_per_group);
  // 此处计算出每个group所处理的repeat个数，方便vec_cross_add()接口调用
  const std::size_t repeat_per_group = (elem_per_group * sizeof(float)) / 256;

  sycl::queue Q(sycl::ascend_selector{});

  auto input_buf = sycl::malloc_device<float>(group_num * elem_per_group, Q);
  // 这里在GM上多申请两块内存，用于存放指数运算结果和求和结果
  auto exp_res_buf = sycl::malloc_device<float>(group_num * elem_per_group, Q);
  auto sum_res_buf = sycl::malloc_device<float>(group_num * repeat_per_group, Q);
  auto res_buf = sycl::malloc_device<float>(group_num * elem_per_group, Q);

  // 由于vec_cross_add是按repeat为单位进行求和的，故申请的向量即为group数量乘以每个group的repeat数量
  std::vector<float> sum_res((group_num * repeat_per_group), 0.0f);
  std::vector<float> res(input_sz, 0.0f);

  Q.memcpy(input_buf, input.data(), byte_count);

  // 第一个核函数进行求和与指数运算的操作
  Q.launch<class Summary>(group_num, [=](sycl::group<1> group) {
    bisheng::vector<float, elem_per_group> input_vec;
    bisheng::vector<float, elem_per_group> exp_res_vec;
    bisheng::vector<float, repeat_per_group> sum_res_vec;
    std::size_t group_id = group.get_group_id();

    input_vec.load(
        sycl::global_ptr<float>(input_buf + group_id * elem_per_group).get(),
        elem_per_group);

    if (tail_elem_count > 0 && group_id == group_num - 1) {
      // 同样的处理尾块中存在运算，且为最后一个group的情况
      bisheng::vector_view<float> input_vec_v(input_vec.data(), tail_elem_count);
      bisheng::vector_view<float> exp_res_vec_v(exp_res_vec.data(), tail_elem_count);

      bisheng::vec_exp(exp_res_vec_v, input_vec_v);
      // 由于尾块中的元素大概率不是整repeat，所以采用标量运算的的方式
      for (int i = 0; i < tail_elem_count; ++i)
        sum_res_vec[0] += exp_res_vec_v[i];
      for (int i = 1; i < repeat_per_group; ++i)
        sum_res_vec[i] = 0.0f;
    } else {
      // 整block情况
      bisheng::vec_exp(exp_res_vec, input_vec);
      // 这里不仅确定是整block，也能确定是整repeat，故直接调用接口
      bisheng::vec_cross_add(sum_res_vec.data(), exp_res_vec);
    }

    // UB -> GM 内存搬移，将指数运算结果与求和结果均保存下来
    exp_res_vec.store(
        sycl::global_ptr<float>(exp_res_buf + group_id * elem_per_group).get(),
        elem_per_group);
    sum_res_vec.store(
        sycl::global_ptr<float>(sum_res_buf + group_id * repeat_per_group).get(),
        repeat_per_group);
  });

  Q.memcpy(sum_res.data(), sum_res_buf, group_num * repeat_per_group * sizeof(float));
  Q.wait();

  // 由于vec_cross_add求和后的结果是多个短向量的和
  // 依然是一个向量，故须在Host端进一步计算为标量
  float sum;
  for (auto x : sum_res)
    sum += x;
  
  // 第二个核函数进行向量除法的运算
  Q.launch<class Softmax>(group_num, [=](sycl::group<1> group) {
    // 只需将上个核函数计算到的指数运算结果向量搬移进来即可
    bisheng::vector<float, elem_per_group> exp_res_vec;
    bisheng::vector<float, elem_per_group> divisor_vec(sum);
    bisheng::vector<float, elem_per_group> res_vec;
    std::size_t group_id = group.get_group_id();

    exp_res_vec.load(
        sycl::global_ptr<float>(exp_res_buf + group_id * elem_per_group).get(),
        elem_per_group);

    // 此处分支大同小异，不再赘述
    if (tail_elem_count > 0 && group_id == group_num - 1) {
      bisheng::vector_view<float> exp_res_vec_v(exp_res_vec.data(), tail_elem_count);
      bisheng::vector_view<float> divisor_vec_v(divisor_vec.data(), tail_elem_count);
      bisheng::vector_view<float> res_vec_v(res_vec.data(), tail_elem_count);

      bisheng::vec_div(res_vec_v, exp_res_vec_v, divisor_vec_v);
    } else {
      bisheng::vec_div(res_vec, exp_res_vec, divisor_vec);
    }

    res_vec.store(
        sycl::global_ptr<float>(res_buf + group_id * elem_per_group).get(),
        elem_per_group);
  });

  Q.memcpy(res.data(), res_buf, byte_count);
  Q.wait();

  sycl::free(input_buf, Q);
  sycl::free(exp_res_buf, Q);
  sycl::free(sum_res_buf, Q);
  sycl::free(res_buf, Q);

  return res;
}

功能验证

功能验证与上述方式相同，但在验证过程中出现了较为严重的问题。

在代码执行过程中，出现了计算结果时对时错的情况，在之前的项目开发过程中其实是踩过这样的坑的，所以出现这种情况时也没有特别慌张。这里分享一下定位问题的心路历程，总结一下设备端代码Debug的思路。

由于对毕昇C++和毕昇编译器这套逻辑了解不够充分，所以可能Debug的方式很笨，但只要能De出来Bug，那就是好方法。

  • 首先可以确定的是，计算向量除法的部分一定没有问题，这是方案一里面验证过的。
  • 其次是要确定访存过程中是否存在问题，是不是访问到了一些不该访问的地方。确定访存无误后再进行下一步。
  • 将每一步的计算结果输出，具体查看到底哪一步计算出现了错误。

首先分析访存的问题，可以先将输入向量的长度和数据类型确定下来，然后带入这个向量长度计算每一次访存的范围。当然你也可以写个脚本来帮你完成这一步，但我懒，我选择草稿纸。一波计算后发现，访存并没有什么问题，每一步操作访问的范围也都是它们应该访问的，并没有访问的未定义数据。那么基本可以确定，这个Bug不是我自己的原因，那就看看每一步的计算结果。

因为第二个核函数已经经过了方案一的验证，所以没有过多纠结，分析第一个核函数。第一个核函数进行了两种运算，指数运算和求和运算。但指数运算也在方案一里验证过了，是没有问题的，所以直接就将问题定位在了求和过程中。使用同一个输入向量，分别输出Host算子中的和与vec_cross_add()计算的和。当然这里输出的和均为标量，vec_cross_add()返回的结果已经在Host端相加计算为了标量。

[Debug]: Host sum: 7550.05
[Debug]: Ascend sum: 7090.52
[Error]: Calculation error.

[Debug]: Host sum: 7549.67
[Debug]: Ascend sum: 7549.66
[Debug]: Result correct.

果然！是求和出现了问题，而且是时对时错的。然后有对这个问题进行了更详细的测试，主要是测试了两种情况，也即第一个核函数中的两个分支。

由于尾块是采用for循环计算的，理论上不会出现错误，但为了严谨还是进行了一些测试。将向量长度锁定在320，迫使它只执行尾块的逻辑，结果如下。

[Debug]: Host sum: 387.095
[Debug]: Ascend sum: 387.095
[Debug]: Result correct.

[Debug]: Host sum: 356.134
[Debug]: Ascend sum: 356.134
[Debug]: Result correct.

无论执行多少次，结果都是正确的。那现在基本可以确定是vec_cross_add()接口出现了问题。所以我们对代码进行修改，将vec_cross_add()接口用for循环代替。

此处的结论不正确，vec_cross_add()接口本身没有任何问题，详细测试及Softmax的重新实现详见关于vec_cross_add接口的详细测试 - 亦初 (deleter-d.github.io)

由于使用标量运算，故每个group求和的结果不再是向量，而是标量。所以存放求和结果的内存空间大小需要做一定的调整，这里申请大小为group_num * (32 / sizeof(data_t))的空间。其实理论上，每个group求和的结果只需要一个data_t数据类型的大小即可，但为了按照block为粒度严格分离group的访存空间，所以申请了与group_num个block大小相同的内存空间来存放。

auto sum_res_buf = sycl::malloc_device<data_t>(group_num * (32 / sizeof(data_t)), Q);

具体求和过程则改为如下方式。

if (tail_elem_count > 0 && group_id == group_num - 1) {
      bisheng::vector_view<data_t> input_vec_v(input_vec.data(), tail_elem_count);
      bisheng::vector_view<data_t> exp_res_vec_v(exp_res_vec.data(), tail_elem_count);

      bisheng::vec_exp(exp_res_vec_v, input_vec_v);
      for (int i = 0; i < tail_elem_count; ++i)
        sum_res_buf[group_id * (32 / sizeof(data_t))] += exp_res_vec_v[i];
    } else {
      bisheng::vec_exp(exp_res_vec, input_vec);
      for (int i = 0; i < elem_per_group; ++i) {
        sum_res_buf[group_id * (32 / sizeof(data_t))] += exp_res_vec[i];
      }
    }

进一步测试后，时对时错的问题解决了，到这里就可以说功能验证正确了，可喜可贺！

性能测试

性能测试结果如下。

测试用例640640064000640000
加速比0.2374321.47898813.1960587.72573

对比方案一可以观察到，对求和进行向量化的意义是非常大的，尤其是向量长度变得越来越长后，这种优化的提升尤为明显。现在的加速比可以说是令人比较满意的了。

空间上的优化（方案三）

通过观察方案二中的数据流动，我们可以发现，在空间利用上有些浪费的地方。先来看一下方案二的数据流动方式。

请添加图片描述

注：图中只描述了GM与UB之间的数据流，其中还发生了GM与Host之间的数据搬移。例如求和结果将会搬回Host，计算为标量后，利用该标量对分母向量进行初始化。

观察上述数据流可以发现，在使用exp_res_buf存储指数运算结果的时候，输入input_buf已经失去了作用，且后面也不会再使用其中的数据。同理，在使用res_buf存储最终结果的时候，exp_res_buf也不再使用了，因为指数运算结果此时已经读入了UB中。所以，input_buf、exp_res_buf和res_buf三者是可以合一的。

继续对UB中的内存使用进行分析。

请添加图片描述

初步的理论分析可知，Kernel 1中的input和exp_res_vec可以合一，Kernel 2中的exp_res_vec与res_vec可以合一。我们在此过程中使用了vec_exp(dst, src)和vec_div(dst, src0, src1)接口，这两个接口分别为一元运算和二元运算。

在毕昇C++中，对于基于bisheng::vector类型的通用一元运算函数接口，dst和src可以是同一个bisheng::vector对象，即原址计算。故Kernel 1中的input和exp_res_vec可以合一。而对于二元运算，目标数据和源数据在不同的repeat迭代之间不允许出现地址重叠，虽然有部分接口例外，但我们所使用的vec_div接口并不在这些例外中，故无法将Kernel 2中的exp_res_vec与res_vec合一。

经过上述一系列空间优化，最终的数据流如图所示。

请添加图片描述

方案实现

代码与方案二几乎一致，只是改变了内存搬移的源地址与目的地址，这里就不再放代码了。

功能验证

经过测试，功能验证正确。

性能测试

性能测试结果如下。

测试用例640640064000640000
加速比0.2246131.51239413.3043388.70575

可以观察到，与方案二相比，时间上几乎没有区别。但由于优化了空间利用率，所以使得设备端可以承载更大长度的向量，优化的意义是比较大的。

分核方案的优化（方案四）

下面所有的讨论均已float类型的数据为例。

分核方案的核心思想就是，尽可能利用所有物理核心，并在此基础上令每个核心处理尽可能多的数据。而我们上面采用的方案是临时将每个核处理的元素数量固定为640，这显然不是最优的方案。

首先是尽可能利用所有的物理核心，昇腾910拥有32个物理核心，所以我们要想办法让32个核心都在工作状态，尽量避免一核干活儿，多核围观的滑稽场景。

首先分析一个问题，逻辑核的数量如何确定？假设输入向量长度为 l e n len len，每个逻辑核处理的运算个数为 n n n，在不考虑有尾块的情况下，可以得到逻辑核数量的公式为 g r o u p _ n u m = l e n ÷ n group\_num=len\div n group_num=len÷n。 n n n由用户指定，不是我们可以控制的，故我们只能在 l e n len len和 g r o u p _ n u m group\_num group_num上做文章。为了更好的理解，我们变形一下公式 l e n = g r o u p _ n u m × n len=group\_num\times n len=group_num×n。这样就可以比较直观的看出，我们需要在 g r o u p _ n u m group\_num group_num和 n n n之间做一个权衡。

这个权衡只有两种思考方式：

  • 一种是确定 g r o u p _ n u m group\_num group_num，根据 g r o u p _ n u m group\_num group_num计算得到 n n n。说人话就是把逻辑核的数量定死，然后根据用户给的向量长度计算每个核要处理的元素个数。
  • 另一种是确定 n n n，即把每个逻辑核要处理的元素数量定死，然后根据用户给的向量长度计算逻辑核数量。

情况一

先来考虑第一种情况，即将 g r o u p _ n u m group\_num group_num定死。假设我们就定为与物理核数相同的数量，即32。考虑一个问题，假设输入向量长度非常长，那么拆分成32份后依然非常长，长到 l e n ÷ 32 len\div 32 len÷32个元素的大小超出了UB的承载范围，那么此时算子就会崩溃。

这时候有人就要说了（假装有人要说）：那不能把逻辑核数量写大一点吗？

好！听你的，我们将逻辑核数量定为320，理想状态下，每个物理核将处理10个逻辑核。此时再考虑一种情况，用户给的输入向量非常短，短到没办法分为320份，此时 l e n ÷ 32 len\div 32 len÷32为0。意味着你的每个逻辑核中，要么是处理尾块，要么根本就没有元素，但320个逻辑核依然会开启。这显然是不够合理的。

情况二

再来考虑将 n n n定死的情况，即将每个逻辑核要处理的元素个数定死，其实就是我们上面方案的使用的策略，这里我们暂时考虑n为640的情况。同样考虑一些比较极端的例子，假设输入向量非常长，此时 l e n ÷ n len\div n len÷n即 l e n ÷ 640 len\div 640 len÷640会非常大，即逻辑核的数量会非常多。虽然这样能够充分利用所以物理核，但每个逻辑核的承载能力远不止640个元素，这样就浪费了单个逻辑核的能力，把资源都消耗在调度逻辑核上了。

这时候又有人要说了（依然假装有人说）：那不能把逻辑核处理的元素个数写大一点吗？

好！还是听你的，我们将 n n n定为UB能够承载的上限 m a x max max。这种情况下，我们甚至不用考虑输入向量长度非常短的情况，只考虑向量长度小于 m a x × 31 max\times 31 max×31的情况，即向量长度小于31个物理核同时工作时可以处理的最大元素个数。此时至少会有1个物理核心在看戏，若向量长度进一步缩短，那看戏的物理核只会越来越多。这显然也是不够合理的。

动态方案

分析完两种情况，可以得出一个结论，单纯的确定 n n n与 g r o u p _ n u m group\_num group_num中的任何一个都是不合适的。

那我们应该怎么确定呢？动态确定！

首先确定一个问题，我们这个算子，每个group处理多少数据是UB的上限。经过测试，每个group最多可以处理87360字节的数据，即 87360 ÷ sizeof ( d a t a _ t ) 87360\div \text{sizeof}(data\_t) 87360÷sizeof(data_t)个元素。这个上限并不是所有算子都一样的，因为每个算子在UB上申请内存的情况不同，所以要具体问题具体分析。

我们继续上面那个公式 l e n = g r o u p _ n u m × n len = group\_num\times n len=group_num×n，这里为了通用性，我们换成处理的字节数 b y t e s bytes bytes，而不是元素个数。进而 n = b y t e s ÷ sizeof ( d a t a _ t ) n=bytes\div\text{sizeof}(data\_t) n=bytes÷sizeof(data_t)，那么公式变为 t o t a l _ b y t e s = l e n × sizeof ( d a t a _ t ) = g r o u p _ n u m × ( b y t e s ÷ sizeof ( d a t a _ t ) ) total\_bytes=len\times\text{sizeof}(data\_t) = group\_num\times (bytes\div\text{sizeof}(data\_t)) total_bytes=len×sizeof(data_t)=group_num×(bytes÷sizeof(data_t))。

  • 当 t o t a l _ b y t e s < 32 × 2560 total\_bytes< 32\times 2560 total_bytes<32×2560时，将 b y t e s bytes bytes定为1280，则 g r o u p _ n u m < 32 group\_num<32 group_num<32，此时算子可能无法充分利用所有物理核。
  • 当 t o t a l _ b y t e s ≥ 32 × 2560 total\_bytes\ge 32\times 2560 total_bytes≥32×2560时，将 b y t e s bytes bytes定为2560，则 g r o u p _ n u m ≥ 32 group\_num\ge 32 group_num≥32，这意味着算子将充分利用所有物理核。
  • 当 t o t a l _ b y t e s ≥ 32 × 5120 total\_bytes\ge 32\times 5120 total_bytes≥32×5120时，将 b y t e s bytes bytes定为5120，则 g r o u p _ n u m ≥ 32 group\_num\ge 32 group_num≥32，也会充分利用所有物理核。
  • 当 t o t a l _ b y t e s ≥ 32 × 12800 total\_bytes\ge 32\times 12800 total_bytes≥32×12800时，将 b y t e s bytes bytes定为12800，则 g r o u p _ n u m ≥ 32 group\_num\ge 32 group_num≥32，也会充分利用所有物理核。
  • 当 t o t a l _ b y t e s ≥ 32 × 25600 total\_bytes\ge 32\times 25600 total_bytes≥32×25600时，将 b y t e s bytes bytes定为25600，则 g r o u p _ n u m ≥ 32 group\_num\ge 32 group_num≥32，同样充分利用所有物理核。
  • 当 t o t a l _ b y t e s ≥ 32 × 51200 total\_bytes\ge 32\times 51200 total_bytes≥32×51200时，将 b y t e s bytes bytes定为51200，则 g r o u p _ n u m ≥ 32 group\_num\ge 32 group_num≥32，同样充分利用所有物理核。
  • 当 t o t a l _ b y t e s ≥ 32 × 87360 total\_bytes\ge 32\times 87360 total_bytes≥32×87360时，将 b y t e s bytes bytes定为87360，则 g r o u p _ n u m ≥ 32 group\_num\ge 32 group_num≥32，同样充分利用所有物理核。

采取这种策略，虽然在输入向量总字节数小/于 32 × 2560 32\times2560 32×2560时可能会出现某些物理核不工作的情况，但考虑到实际情况下，输入向量都是比较长的向量，这种偶尔的空闲是可接受的。

这里只是展示一种思路，条件分支中的阈值是可以根据实际情况进行调整的，并不是只能按照2560、5120等阈值进行分割。

方案实现

同样还是注意带注释的地方，其余地方与之前相同。

std::vector<data_t> ascend_softmax(std::vector<data_t> input) {
  std::size_t input_sz = input.size();
  std::size_t byte_count = input_sz * sizeof(data_t);

  if (byte_count < 32)
    return softmax(input);

  // 这里依照上面介绍的策略确定每个逻辑核所处理的元素个数
  std::size_t elem_per_group = 0;
  if (byte_count >= PHYSICAL_CORES * UB_MAX_BYTES)
    elem_per_group = UB_MAX_BYTES / sizeof(data_t);
  else if (byte_count >= PHYSICAL_CORES * 51200)
    elem_per_group = 51200 / sizeof(data_t);
  else if (byte_count >= PHYSICAL_CORES * 25600)
    elem_per_group = 25600 / sizeof(data_t);
  else if (byte_count >= PHYSICAL_CORES * 12800)
    elem_per_group = 12800 / sizeof(data_t);
  else if (byte_count >= PHYSICAL_CORES * 5120)
    elem_per_group = 5120 / sizeof(data_t);
  else if (byte_count >= PHYSICAL_CORES * 2560)
    elem_per_group = 2560 / sizeof(data_t);
  else
    elem_per_group = 1280 / sizeof(data_t);

  const std::size_t tail_elem_count = input_sz % elem_per_group;
  const std::size_t group_num = (tail_elem_count > 0)
                                    ? ((input_sz / elem_per_group) + 1)
                                    : (input_sz / elem_per_group);

  sycl::queue Q(sycl::ascend_selector{});

  auto dev_buf = sycl::malloc_device<data_t>(group_num * elem_per_group, Q);
  auto sum_res_buf = sycl::malloc_device<data_t>(group_num * (32 / sizeof(data_t)), Q);

  std::vector<data_t> sum_res(group_num * (32 / sizeof(data_t)), 0.0f);
  std::vector<data_t> res(input_sz, 0.0f);

  Q.memcpy(dev_buf, input.data(), byte_count);

  Q.launch<class Summary>(group_num, [=](sycl::group<1> group) {
    // 此处直接申请最大空间，因为定义毕昇向量时指定大小必须用常量表达式，大小需要在编译时确定
    // 由于前面采用了动态策略，所以不能直接使用elem_per_group来定义毕昇向量
    // 只需要在使用时控制访存范围即可，第二个核函数同理，不再赘述
    bisheng::vector<data_t, UB_MAX_BYTES / sizeof(data_t)> input_vec;
    std::size_t group_id = group.get_group_id();

    input_vec.load(
        sycl::global_ptr<data_t>(dev_buf + group_id * elem_per_group).get(),
        elem_per_group);

    if (tail_elem_count > 0 && group_id == group_num - 1) {
      bisheng::vector_view<data_t> input_vec_v(input_vec.data(), tail_elem_count);

      bisheng::vec_exp(input_vec_v, input_vec_v);
      for (int i = 0; i < tail_elem_count; ++i)
        sum_res_buf[group_id * (32 / sizeof(data_t))] += input_vec_v[i];
    } else {
      // 由于毕昇向量定义了最大长度，故即使是整block的情况，也需要用变长向量来控制访存范围
      bisheng::vector_view<data_t> input_vec_v(input_vec.data(), elem_per_group);
      bisheng::vec_exp(input_vec_v, input_vec_v);
      for (int i = 0; i < elem_per_group; ++i) {
        sum_res_buf[group_id * (32 / sizeof(data_t))] += input_vec_v[i];
      }
    }

    input_vec.store(
        sycl::global_ptr<data_t>(dev_buf + group_id * elem_per_group).get(),
        elem_per_group);
  });

  Q.memcpy(sum_res.data(), sum_res_buf, group_num * (32 / sizeof(data_t)) * sizeof(data_t));
  Q.wait();

  data_t sum;
  for (int i = 0; i < sum_res.size(); i += 32 / sizeof(data_t))
    sum += sum_res[i];

  Q.launch<class Softmax>(group_num, [=](sycl::group<1> group) {
    bisheng::vector<data_t, UB_MAX_BYTES / sizeof(data_t)> exp_res_vec;
    bisheng::vector<data_t, UB_MAX_BYTES / sizeof(data_t)> divisor_vec(sum);
    bisheng::vector<data_t, UB_MAX_BYTES / sizeof(data_t)> res_vec;
    std::size_t group_id = group.get_group_id();

    exp_res_vec.load(
        sycl::global_ptr<data_t>(dev_buf + group_id * elem_per_group).get(),
        elem_per_group);

    if (tail_elem_count > 0 && group_id == group_num - 1) {
      bisheng::vector_view<data_t> exp_res_vec_v(exp_res_vec.data(), tail_elem_count);
      bisheng::vector_view<data_t> divisor_vec_v(divisor_vec.data(), tail_elem_count);
      bisheng::vector_view<data_t> res_vec_v(res_vec.data(), tail_elem_count);

      bisheng::vec_div(res_vec_v, exp_res_vec_v, divisor_vec_v);
    } else {
      bisheng::vector_view<data_t> exp_res_vec_v(exp_res_vec.data(), elem_per_group);
      bisheng::vector_view<data_t> divisor_vec_v(divisor_vec.data(), elem_per_group);
      bisheng::vector_view<data_t> res_vec_v(res_vec.data(), elem_per_group);
      bisheng::vec_div(res_vec_v, exp_res_vec_v, divisor_vec_v);
    }

    res_vec.store(
        sycl::global_ptr<data_t>(dev_buf + group_id * elem_per_group).get(),
        elem_per_group);
  });

  Q.memcpy(res.data(), dev_buf, byte_count);
  Q.wait();

  sycl::free(dev_buf, Q);
  sycl::free(sum_res_buf, Q);

  return res;
}

功能测试

功能测试验证正确。

性能测试

性能测试结果如下。

测试用例640640064000640000
加速比0.2343731.43694112.3590880.59147

分析了许多，本想着动态分核结果会有惊喜。嘿！您猜怎么着？还真是大惊喜！

在动态分核的策略下，当向量长度总字节数不少于 32 × 2560 32\times 2560 32×2560时，总能保证32个物理核都在工作，而且不至于令逻辑核数量过多，但神奇的事情来了，这种策略成功实现了负优化！！

本方案的性能测试结果看起来还不错，但我们继续增大向量长度，使得总字节数到达划分策略的阈值附近。以float类型的数据为例，令向量长度为698880，此时总字节数为 698880 × 4 = 2795520 = 87360 × 32 698880\times 4=2795520=87360\times 32 698880×4=2795520=87360×32。此时这种策略将分配32个逻辑核，完美贴合物理核数量，每个核心处理87360字节的数据，完美贴合UB承载的上限。惊喜的事情来了，请看加速比

测试用例698880
方案三加速比96.29706
方案四加速比67.26953

什么鬼情况？！？！

我们分别观察一下它们的分核情况。

方案三如下：

[PERMORMANCE]: Host time cost: 93873327 ns
[Debug]: Group num: 1875 Elements per group: 640
[PERMORMANCE]: Ascend time cost: 728001 ns

方案四如下：

[PERMORMANCE]: Host time cost: 94127478 ns
[Debug]: Group num: 55 Elements per group: 21840
[PERMORMANCE]: Ascend time cost: 785999 ns

可以发现两者都充分利用了所有物理核，但方案四的策略使得加速比下降了，反观方案三的1875个逻辑核取得了完胜。但转念一想，是不是让每个逻辑核承载到UB的上限有点过分，那么再来测试一下正常压力下的表现。

还是以float类型数据为例，向量长度为102400，此时总字节数为 102400 × 4 = 409600 = 12800 × 32 102400\times 4=409600=12800\times 32 102400×4=409600=12800×32，此时将分配32个逻辑核，每个核处理12800字节的数据，远不到UB承载的上限。

方案三分核情况如下：

[PERMORMANCE]: Host time cost: 7872144 ns
[Debug]: Group num: 160 Elements per group: 640
[PERMORMANCE]: Ascend time cost: 375999 ns

方案四分核情况如下：

[PERMORMANCE]: Host time cost: 7813204 ns
[Debug]: Group num: 32 Elements per group: 3200
[PERMORMANCE]: Ascend time cost: 400999 ns

加速比如下：

测试用例102400
方案三加速比20.73252
方案四加速比19.4274

依然是有略微的下降，这也排除了UB压力过大的问题。

结论

经过一系列分析，目前能够得出的结论是，尽可能多的逻辑核数量的收益要大于单逻辑核内处理尽可能多的数据。

异构分核的坑还是太多了，踩都踩不完，过程中有很多反直觉的情况，必须靠实验来佐证。

完整代码

最后贴上目前效果最好（方案三）的完整代码，其中包括一些自定义的Debug信息，不用太纠结。

#include <cmath>
#include <iomanip>
#include <iostream>
#include <random>
#include <stdlib.h>
#include <time.h>
#include <vector>

#include <bisheng/bisheng.hpp>
#include <sycl/sycl.hpp>

#define DEBUG
#define DEBUG_HEAD "\033[34m[Debug]: \033[0m"
#define ERROR_HEAD "\033[31m[Error]: \033[0m"
#define PERFORMANCE
#define PERFORMANCE_HEAD "\033[36m[PERMORMANCE]: \033[0m"
#define INPUT_COUNT 102400

std::vector<float> softmax(std::vector<float> input) {
#ifdef DEBUG
  std::cout << DEBUG_HEAD << "The host operator is called.\n";
#endif

#ifdef PERFORMANCE
  struct timespec time;
  clock_gettime(CLOCK_REALTIME, &time);
  auto start_time = time.tv_sec * 1000000000 + time.tv_nsec;
#endif

  float sum = 0.0;
  for (auto x : input) {
    sum += expf(x);
  }

#ifdef DEBUG
  std::cout << DEBUG_HEAD << "Host sum: " << sum << "\n";
#endif

  std::vector<float> res;
  for (auto x : input) {
    res.push_back(expf(x) / sum);
  }

#ifdef PERFORMANCE
  clock_gettime(CLOCK_REALTIME, &time);
  auto end_time = time.tv_sec * 1000000000 + time.tv_nsec;
  std::cout << PERFORMANCE_HEAD << "Host time cost: " << end_time - start_time
            << " ns" << std::endl;
#endif

  return res;
}

std::vector<float> ascend_softmax(std::vector<float> input) {
  std::size_t input_sz = input.size();
  std::size_t byte_count = input_sz * sizeof(float);

  // call the host operator if input isn't enough a full block
  if (byte_count < 32) {
#ifdef DEBUG
    std::cout << DEBUG_HEAD
              << "The input vector is not enough for a full block.\n";
#endif
    return softmax(input);
  }

  // ascend code start
#ifdef DEBUG
  std::cout << DEBUG_HEAD << "The ascend operator is called.\n";
#endif

  // number of elements per group
  const std::size_t elem_per_group = 640;
  // number of elements in tail block
  const std::size_t tail_elem_count = input_sz % elem_per_group;
  // number of groups
  // if tail block is exist, apply for one more group
  const std::size_t group_num = (tail_elem_count > 0)
                                    ? ((input_sz / elem_per_group) + 1)
                                    : (input_sz / elem_per_group);

#ifdef DEBUG
  std::cout << DEBUG_HEAD << "Group num: " << group_num
            << " Elements per group: " << elem_per_group << "\n";
#endif

  sycl::queue Q(sycl::ascend_selector{}, nullptr,
                {sycl::property::queue::enable_profiling()});

  // GM memory allocation
  auto dev_buf = sycl::malloc_device<float>(group_num * elem_per_group, Q);
  auto sum_res_buf =
      sycl::malloc_device<float>(group_num * (32 / sizeof(float)), Q);

  // Host memory allocation
  std::vector<float> sum_res(group_num * (32 / sizeof(float)), 0.0f);
  std::vector<float> res(input_sz, 0.0f);

  // host -> GM
  Q.memcpy(dev_buf, input.data(), byte_count);

#ifdef DEBUG
  std::cout << DEBUG_HEAD << "Kernel function started.\n";
#endif

  sycl::event e0 =
      Q.launch<class Summary>(group_num, [=](sycl::group<1> group) {
        bisheng::vector<float, elem_per_group> input_vec;
        std::size_t group_id = group.get_group_id();

        // GM -> UB
        input_vec.load(
            sycl::global_ptr<float>(dev_buf + group_id * elem_per_group).get(),
            elem_per_group);

        if (tail_elem_count > 0 && group_id == group_num - 1) {
          // if tail block has element and this is the last group
          bisheng::vector_view<float> input_vec_v(input_vec.data(),
                                                  tail_elem_count);

          bisheng::vec_exp(input_vec_v, input_vec_v);
          for (int i = 0; i < tail_elem_count; ++i)
            sum_res_buf[group_id * (32 / sizeof(float))] += input_vec_v[i];
        } else {
          // full block data
          bisheng::vec_exp(input_vec, input_vec);
          for (int i = 0; i < elem_per_group; ++i) {
            sum_res_buf[group_id * (32 / sizeof(float))] += input_vec[i];
          }
        }

        // UB -> GM
        input_vec.store(
            sycl::global_ptr<float>(dev_buf + group_id * elem_per_group).get(),
            elem_per_group);
      });

  // GM -> Host
  Q.memcpy(sum_res.data(), sum_res_buf,
           group_num * (32 / sizeof(float)) * sizeof(float));
  Q.wait();

  float sum;
  for (int i = 0; i < sum_res.size(); i += 32 / sizeof(float))
    sum += sum_res[i];
#ifdef DEBUG
  std::cout << DEBUG_HEAD << "Ascend sum: " << sum << "\n";
#endif

  sycl::event e1 =
      Q.launch<class Softmax>(group_num, [=](sycl::group<1> group) {
        // UB memory of exponent result
        bisheng::vector<float, elem_per_group> exp_res_vec;
        // UB memory of divisor
        bisheng::vector<float, elem_per_group> divisor_vec(sum);
        // UB memory of final result
        bisheng::vector<float, elem_per_group> res_vec;
        std::size_t group_id = group.get_group_id();

        // GM -> UB
        exp_res_vec.load(
            sycl::global_ptr<float>(dev_buf + group_id * elem_per_group).get(),
            elem_per_group);

        if (tail_elem_count > 0 && group_id == group_num - 1) {
          // if tail block has element and this is the last group
          bisheng::vector_view<float> exp_res_vec_v(exp_res_vec.data(),
                                                    tail_elem_count);
          bisheng::vector_view<float> divisor_vec_v(divisor_vec.data(),
                                                    tail_elem_count);
          bisheng::vector_view<float> res_vec_v(res_vec.data(),
                                                tail_elem_count);

          bisheng::vec_div(res_vec_v, exp_res_vec_v, divisor_vec_v);
        } else {
          // full block data
          bisheng::vec_div(res_vec, exp_res_vec, divisor_vec);
        }

        // UB -> GM
        res_vec.store(
            sycl::global_ptr<float>(dev_buf + group_id * elem_per_group).get(),
            elem_per_group);
      });

#ifdef DEBUG
  std::cout << DEBUG_HEAD << "Kernel function finished.\n";
#endif

  // GM -> host
  Q.memcpy(res.data(), dev_buf, byte_count);
  Q.wait();

  sycl::free(dev_buf, Q);
  sycl::free(sum_res_buf, Q);

  // ascend code end

#ifdef PERFORMANCE
  const uint64_t e0_start_time =
      e0.get_profiling_info<sycl::info::event_profiling::command_start>();
  const uint64_t e0_end_time =
      e0.get_profiling_info<sycl::info::event_profiling::command_end>();
  const uint64_t e1_start_time =
      e1.get_profiling_info<sycl::info::event_profiling::command_start>();
  const uint64_t e1_end_time =
      e1.get_profiling_info<sycl::info::event_profiling::command_end>();
  std::cout << PERFORMANCE_HEAD << "Ascend time cost: "
            << (e0_end_time - e0_start_time) + (e1_end_time - e1_start_time)
            << " ns" << std::endl;
#endif

  return res;
}

int main() {
#ifdef DEBUG
  std::cout << DEBUG_HEAD << "Compile succeed" << std::endl;
#endif

  std::vector<float> vec;

  std::random_device rd;
  std::mt19937 gen(rd());
  std::uniform_real_distribution<float> urdis(-1, 1);
  for (int i = 0; i < INPUT_COUNT; i++) {
    vec.push_back(urdis(gen));
  }

  std::vector<float> host_res = softmax(vec);
  std::vector<float> ascend_res = ascend_softmax(vec);

  for (int i = 0; i < host_res.size(); ++i) {
    if (std::fabs(host_res[i] - ascend_res[i]) / host_res[i] > 0.02) {
      std::cout << ERROR_HEAD << "Calculation error." << std::endl;
      return EXIT_FAILURE;
    }
  }
  std::cout << DEBUG_HEAD << "Result correct." << std::endl;

  return EXIT_SUCCESS;
}
GoldPancake
关注
SoftMax函数
且慢
08-30 1210
SOFTMAX回归模型 softmax回归跟线性回归⼀样将输⼊特征与权重做线性叠加。与线性回归的⼀个主要不同在于，softmax回归的输出值个数等于标签⾥的类别数。 softmax函数是一个单层的神经网络，例如由一个例子： 4个输入特征，三个类别 计算方法： 但是这样输出的值并不能代表分类的结果，因为有的输出的值为100，无法表示是三个类别重的哪一类。 softmax运算符（softmax operator）解决了这个问题。 softmax运算符 ...
onlyoffice毕升office
xiehongan123的专栏
03-04 7425
毕升文档| 多人协同编辑 | 在线Office| onlyOffice 毕升文档公网免费使用地址：https://bishengoffice.com 毕升文档免费部署文档：https://bishengoffice.com/apps/blog/free/ 毕升文档免费部署脚本地址：https://gitee.com/ibisheng/deploy.git 毕升文档在线文件服务部分的...
softmax算法详解
热门推荐
星辰的博客
03-09 1万+
softmax简介 在机器学习尤其是深度学习中，softmax是个非常常用而且比较重要的函数，尤其在多分类的场景中使用广泛。他把一些输入映射为0-1之间的实数，并且归一化保证和为1，因此多分类的概率之和也刚好为1。 2.softmax函数的数学形式定义 可以看到，softmax函数解决了从一个原始的n维向量，向一个n维的概率分布映射的问题。那么在多分类问题中，假设分类数是n，模型希望预测的就是某样本在n个分类上的概率分布。如果用深度学习模型进行建模，那么最后输出层的形式是由n个神经元组成的，再把n个神经元
CUDA 算子优化：Softmax
qq_56047026的博客
06-11 1488
把输入映射为0-1之间的实数，并且归一化保证和为1。max非黑即白，soft，输出是每个分类被取到的概率。定义：该元素的softmax值就是该元素的指数和所有元素指数和的比值定义的原因：1.softmax设计的初衷是希望特征对概率的影响是乘性的。2.多类目标分类问题的目标函数常常选为cross-entropy。
SoftMax
qq_38816432的博客
12-02 478
文章目录1. softmax 1. softmax softmax算子一般用于多分类场景，把神经元的输出映射到（0-1）之间的实数，并且归一化保证和为1，从而使得多分类的概率之和也刚好为1。使得经过softmax后的输出为每个类别的概率，并且各个概率的和为1。 softmax的计算过程为：某元素的指数与所有元素指数和的比值。 在实际应用中，需要注意softmax中存在的数据溢出问题，如果softmax的输入V过大，经过指数运算后，很容易出现数据溢出。解决办法是：将V中的每一个数据，都减去V中的最大值：
softmax算子实现与优化
qq_45642593的博客
02-25 584
实现一个包含1
基于Shell脚本的一键部署毕升文档云平台源码
最新发布
09-28
本项目是一款基于Shell脚本一键部署的毕升文档云平台源码，总计136个文件，涵盖37个字体文件、28个图标文件、18个Shell脚本文件、17个PNG图片文件、5个配置文件、4个YAML文件、2个BMP和DOCX/PPTX/xlsx文件。...
Office文档在线预览-在线编辑解决方案-毕升OfficeAPI说明.docx
10-26
Office 文档在线预览-在线编辑解决方案-毕升 OfficeAPI 说明 本文档介绍了毕升 Office API 的集成解决方案，实现 Office 文档在线预览和在线编辑功能。该解决方案可以供其他在线系统（OA、ERP、网盘等）调用，实现...
毕升Office免费私有化部署手册1
08-08
毕升Office免费私有化部署手册1 毕升Office免费私有化部署手册1是毕升Office系列的手册之一，本手册旨在指导用户如何实施毕升Office的免费私有化部署。本手册涵盖了毕升Office的安装、配置、部署等方面的知识点，...
毕升文档云平台包括drive和在线文件服务功能
04-15
毕升文档云平台包括drive和在线文件服务功能。其中drive实现文件的在线管理，组织结构权限，分享，团队协作等，文件的全文检索等功能；在线文件服务能够处word,ppt,excel格式文件的带水印预览以及多人协同编辑，另外...
softmax算法的理解
11-29
通过介绍可以对softmax有一个更好地理解，通过对其深入学习。
Softmax算法原理及实现
abcdefg90876的博客
05-03 849
逻辑回归主要是用于处理二分类问题，如果面对的是多分类问题，如手写字识别，其中有十个类别，这时候就需要对逻辑回归进行推广，且同样任意两个类之间都是线性可分的。SoftmaxReg...
TVM Relay softmax算子调度分析
Eurypterid的博客
08-05 662
tvm relay 算子调度分析
softmax回归及其实现（TensorFlow）
相国大人
03-29 1192
在之前的博文《logistic回归》中，我们简单的提到了softmax回归。本文将首先介绍softmax回归的基本原理。然后比较softmax回归于logistic回归的关联。最后用开源TensorFlow编写算法并应用于手写数字（MNIST）的识别。softmax原理softmax与logistic用TensorFlow实现softmax regression识别手写数字#!/usr/bin/en
神经网络之softmax(作用，工作原理【示例说明】，损失计算)
m0_45267220的博客
03-08 1万+
神经网络之softmax,说明作用，其位于神经网络中的位置，工作原理【示例说明】，损失计算
softmax手动实现
weixin_44223180的博客
09-13 230
def accuracy(y_hat, y): #@save """计算预测正确的数量""" if len(y_hat.shape) > 1 and y_hat.shape[1] > 1: # y_hat.shape返回y_hat的形状，len()返回第一维度的长度 #如果y_hat不是标量且y_hat的列>1 y_hat = y_hat.argmax(axis=1) #按列降维（列相加） cmp = y_hat.t
深层次理解 Sigmoid - Softmax 激活函数
nobles007820的博客
08-17 2651
这两个概念是涉及到神经网络的激活函数，尤其是针对分类问题最后一层的激活函数，sigmoid对应2分类问题，softmax对应多分类问题。 Softmax: soft为“软”或更加温和 反义词为 hard 硬，深度学习稀疏优化算法常用的软阈值和硬阈值算子就是此概念。 另外一个单词为max。如果单用max，则神经网络最后一层输出的是一个非0即1的向量，且只有输出值一个1（对应前一层向量最大的值的位置）. 如上图例子：最后输出的是[1 0 0]. 两个单词合到一起：我们希望分值大的那一项被经...
详解Softmax多分类算法及交叉熵（Cross-Entropy）成本函数（附代码实现、计算过程、公式）
Brandonn_Xu的博客
10-26 2454
Softmax 是一个用于多类别分类问题的激活函数，它通常用于神经网络的输出层。Softmax 函数将一个实数向量转换为概率分布，其中每个元素表示一个类别的概率。简单来说，Softmax的输出代表了一个样本属于某一类别的概率。：Softmax 将输入转换成一个概率分布，其中每个元素表示相应类别的概率。这意味着所有输出元素的总和将等于 1。：Softmax 常用于多类别分类问题，其中有多个互斥的类别需要被分配概率。图中，an代表了该y=n的概率预测。
如何用OpenAI Triton 优化Softmax算子
Hi20240217的博客
02-21 1782
本文介绍了如何用OpenAI Triton优化Softmax算子,从TensorRT Softmax算子的性能测试到Triton Softmax算子的不同实现。
毕升office部署
09-09
毕升Office的部署需要按照以下步骤进行： 1. 首先，确保你的系统中的80和443端口没有被占用。这是因为毕升Office安装完成后自带了nginx，并且会使用这两个端口进行通信。 2. 如果你只需要集成毕升Office来实现...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

GoldPancake

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值