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

Question

I use a cuda kernel to do a sigmoid activation on a thrust vector:

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 );

where my kernel is:

__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 );  
}

How can I use thrust::replace with a functor / predicate to do the same?

The examples I have seen are too simplistic to demonstrate such use:

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

Or

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

Answer 1

In the "Thrust Quick Start Guide" on page 8-9 there is an example on how to create your own functions for transform.

I came up with a solution, but notice this will not run on the host side because you use CUDA intrinsics.

Code

#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;
}