推力::设备向量使用推力::替换或推力::转换与自定义函子/谓词

Question

我使用 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(Y.begin(), Y.end(), 1, 10);

或者

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

Answer 1

在第 8-9 页的“Thrust 快速入门指南”中，有一个关于如何创建自己的变换函数的示例。

我想出了一个解决方案，但请注意这不会在主机端运行，因为您使用的是 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;
}