thrust::device_vector使用自定义函子/谓词来替换或转换

thrust::device_vector use thrust::replace or thrust::transform with custom functor/predicate

本文关键字：谓词 转换 替换 自定义 device vector thrust      更新时间：2023-10-16

我使用cuda内核对推力向量进行sigmoid激活：

thrust::device_vector<float> output = input;
float * output_ptr = thrust::raw_pointer_cast( output.data() );
sigmoid_activation<<<num_blocks_x,block_threads_x>>>( output_ptr );

我的内核在哪里：

__device__ float sigmoid_function( float input, float skew )
{
    // -X: Neg X
    float x_neg = __fmul_rz( -1.f, input );
    // Y: exponential value
    float exp_val = __expf( x_neg );
    // 1 + exp^(-X)
    float denom = __fadd_rz( 1.f, e_to_x_neg );
     // 1 / 1 + exp^(-X)
    float output  = __fdividef( 1.f, denom );
    if ( skew != 0.0 )
        return _fadd_rz( output, skew );
    else
        return output;
}
__global__ void sigmoid_activation( float * input float skew )
{
    // Iterate Input vector
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    // Update value
    input[x]  = sigmoid_function( input[x], skew );  
}

如何使用thrust:：replace替换为函子/谓词来执行同样的操作？

我看到的例子过于简单，无法证明这种用途：

thrust::replace(Y.begin(), Y.end(), 1, 10);

或

thrust::transform(X.begin(), X.end(), Y.begin(),thrust::negate<int>());

在第8-9页的"Thrust Quick Start Guide"中，有一个关于如何创建自己的转换函数的示例。

我提出了一个解决方案，但请注意，这不会在主机端运行，因为您使用CUDA内部函数。

代码

#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/replace.h>
#include <thrust/functional.h>
#include <iostream>
template<typename T>
struct sigmoid_function
{
  float _skew;
  sigmoid_function(float skew) : _skew(skew) { /*Empty */ }
  typedef T argument_type;
  typedef T result_type;
  __device__ T operator()(const T &x) const {
    float x_neg = __fmul_rz( -1.f, x );
    float exp_val = __expf( x_neg );
    float denom = __fadd_rz( 1.f, __expf(-exp_val) );
    float output  = __fdividef( 1.f, denom );
    if ( _skew != 0.0 )
        return __fadd_rz( output, _skew );
    else
        return output;
  }
};
int main(void) {
    // allocate three device_vectors with 10 elements
    thrust::device_vector<float> X(10);
    // initialize X to 0,1,2,3, ....
    thrust::sequence(X.begin(), X.end());
    // Before
    thrust::copy(X.begin(),X.end(),std::ostream_iterator<float>(std::cout, "n"));
    // Apply
    thrust::transform(X.begin(), X.end(), X.begin(), sigmoid_function<float>(0.1));
    // After
    thrust::copy(X.begin(),X.end(),std::ostream_iterator<float>(std::cout, "n"));
    return 0;
}
相关文章：