Cuda lambda 与函子的使用答案

【问题标题】：Cuda lambda vs functor usageCuda lambda 与函子的使用
【发布时间】：2022-08-13 23:57:42
【问题描述】：

我在 CUDA 中有一个使用函子的简单函数

struct MT {
    const float _beta1;
    const float _mb1;

    MT(const float beta1, const float mb1) : _beta1(beta1), _mb1(mb1) { }
    
    __device__
    float operator()(const float& op, const float& gradient) {
        return _beta1 * op + _mb1 * gradient;
    }
};


void example(const thrust::device_vector<float>& gradients, thrust::device_vector<float>& d_weights)
{
    thrust::transform(_mt.begin(), _mt.end(), gradients.begin(), _mt.begin(), MT(_beta1, _mb1));
}

然而，这个等价的例子会崩溃（很好地符合 --extended-lambda flat）。是否有另一种标志或不同的方式来表达它以使其运行。函子很好，但 lambda 看起来更整洁。

void example_crash(const thrust::device_vector<float>& gradients, thrust::device_vector<float>& d_weights)
{
    thrust::transform(_mt.begin(), _mt.end(), gradients.begin(), _mt.begin(), [this](const float& op,const float& gradient) { return _beta1 * op + _mb1 * gradient; });
}

错误是

Exception thrown at 0x00007FFA833D4FD9 in Optioniser.exe: Microsoft C++ exception: thrust::system::system_error at memory location 0x00000031ED7FCDD0.
Exception thrown: \'System.Runtime.InteropServices.SEHException\' in AARC.Optimisation.dll
An exception of type \'System.Runtime.InteropServices.SEHException\' occurred in AARC.Optimisation.dll but was not handled in user code
External component has thrown an exception.

我不清楚您从示例函数中获得_beta1、_mb1 和_mt 的位置。 example 实际上是更大类中的方法吗？
对于 lambda 版本，我会期待 [_beta1,_mb1](const float& op,const float& gradient) { return _beta1 * op + _mb1 * gradient; }，即缺少捕获，对吗？
或者，占位符表达式也应该起作用：thrust::transform(_mt.begin(), _mt.end(), gradients.begin(), _mt.begin(), _beta1 * _1 + _mb1 * _2);。
请尝试提供minimal, reproducible example。所以，main()，除非它对您的示例至关重要，否则也没有推力。
推力是这个例子的重点。 std::transform 很好

标签： c++ cuda thrust

【解决方案1】：

你的example 和example_crash 函数对我来说没有意义，因为我不知道_mt 是什么，而且你似乎没有使用d_weights。

如果我们解决了这个问题，那么您的 lambda 至少有几个问题，其中一个是没有 __device__ 装饰（这是必要的，在这里）。

进行各种更改并修复您未显示的内容，这对我有用：

$ cat t2093.cu
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/host_vector.h>
#include <thrust/copy.h>
#include <iostream>

struct MT {
    const float _beta1;
    const float _mb1;

    MT(const float beta1, const float mb1) : _beta1(beta1), _mb1(mb1) { }

    __device__
    float operator()(const float& op, const float& gradient) {
        return _beta1 * op + _mb1 * gradient;
    }
};

const float _beta1 = 1.0f;
const float _mb1 = 1.0f;
void example(const thrust::device_vector<float>& gradients, thrust::device_vector<float>& _mt)
{
    thrust::transform(_mt.begin(), _mt.end(), gradients.begin(), _mt.begin(), MT(_beta1, _mb1));
};

void example_crash(const thrust::device_vector<float>& gradients, thrust::device_vector<float>& _mt)
{
    thrust::transform(_mt.begin(), _mt.end(), gradients.begin(), _mt.begin(), [=] __device__ (const float& op,const float& gradient) { return _beta1 * op + _mb1 * gradient; });
};

const int len = 1000;
int main(){

  thrust::device_vector<float> g1(len, 1.0f);
  thrust::device_vector<float> mt1(len, 2.0f);
  example(g1, mt1);
  thrust::host_vector<float> h_mt1 = mt1;
  thrust::copy_n(h_mt1.begin(), 2, std::ostream_iterator<float>(std::cout, ","));
  std::cout << std::endl;
  thrust::device_vector<float> g2(len, 1.0f);
  thrust::device_vector<float> mt2(len, 2.0f);
  example_crash(g2, mt2);
  thrust::host_vector<float> h_mt2 = mt2;
  thrust::copy_n(h_mt2.begin(), 2, std::ostream_iterator<float>(std::cout, ","));
  std::cout << std::endl;
}
$ nvcc -o t2093 t2093.cu --extended-lambda
$ compute-sanitizer ./t2093
========= COMPUTE-SANITIZER
3,3,
3,3,
========= ERROR SUMMARY: 0 errors
$

【讨论】：