2023-10-31 14:35:29 +08:00
|
|
|
#include <CudaMatrixPrivate.cuh>
|
|
|
|
|
#include <math.h>
|
|
|
|
|
#include <thrust/transform.h>
|
|
|
|
|
#include <thrust/functional.h>
|
|
|
|
|
#include <thrust/execution_policy.h>
|
|
|
|
|
using namespace thrust::placeholders;
|
|
|
|
|
|
2023-11-01 14:31:29 +08:00
|
|
|
struct PowOperator: public thrust::unary_function<float, float>{
|
2023-10-31 14:35:29 +08:00
|
|
|
float exponent;
|
|
|
|
|
PowOperator(float v):exponent(v) {}
|
|
|
|
|
void setExponent(float v){
|
|
|
|
|
exponent = v;
|
|
|
|
|
}
|
|
|
|
|
__host__ __device__
|
|
|
|
|
float operator()(const float& x) {
|
|
|
|
|
return powf(x, exponent);
|
|
|
|
|
}
|
|
|
|
|
};
|
|
|
|
|
|
|
|
|
|
void unaryAdd(float* in1, float* in2, float* out, unsigned long length)
|
|
|
|
|
{
|
|
|
|
|
thrust::plus<float> op;
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,in2,out,op);
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unaryAdd(float* in1, const float& in2, float* out, unsigned long length)
|
|
|
|
|
{
|
2023-11-01 14:31:29 +08:00
|
|
|
thrust::transform(thrust::device,in1,in1+length,out,in2 + _1);
|
2023-10-31 14:35:29 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unaryMul(float* in1, float* in2, float* out, unsigned long length)
|
|
|
|
|
{
|
|
|
|
|
thrust::multiplies<float> op;
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,in2,out,op);
|
|
|
|
|
}
|
|
|
|
|
|
2023-11-01 14:31:29 +08:00
|
|
|
void unaryMul(float* in1, const float& in2, float* out, unsigned long length)
|
|
|
|
|
{
|
|
|
|
|
thrust::transform(thrust::device,in1, in1+length, out, in2 * _1);
|
|
|
|
|
}
|
|
|
|
|
|
2023-10-31 14:35:29 +08:00
|
|
|
void unaryNeg(float* in1, float* out, unsigned long length){
|
|
|
|
|
thrust::negate<float> op;
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,out,op);
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unarySub(float* in1, float* in2, float* out, unsigned long length){
|
|
|
|
|
thrust::minus<float> op;
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,in2,out,op);
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unaryDiv(float* in1, float* in2, float* out, unsigned long length){
|
|
|
|
|
thrust::divides<float> op;
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,in2,out,op);
|
|
|
|
|
}
|
|
|
|
|
|
2023-11-01 14:31:29 +08:00
|
|
|
void unarySub(const float& in1, float* in2, float* out, unsigned long length){
|
|
|
|
|
thrust::transform(thrust::device,in2,in2+length,out,in1-_1);
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unaryDiv(const float& in1, float* in2, float* out, unsigned long length){
|
|
|
|
|
thrust::transform(thrust::device,in2,in2+length,out,in1/_1);
|
|
|
|
|
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unarySub(float* in1, const float& in2, float* out, unsigned long length){
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,out,_1-in2);
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void unaryDiv(float* in1, const float& in2, float* out, unsigned long length){
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,out,_1/in2);
|
|
|
|
|
}
|
|
|
|
|
|
2023-10-31 14:35:29 +08:00
|
|
|
void unaryPow(float* in1, float N,float* out, unsigned long length){
|
|
|
|
|
if (N == 0.0f)
|
|
|
|
|
{
|
|
|
|
|
thrust::fill(out,out+length,0);
|
|
|
|
|
return;
|
|
|
|
|
}
|
|
|
|
|
if (N == 1.0f)
|
|
|
|
|
{
|
|
|
|
|
thrust::copy(in1,in1+length,out);
|
|
|
|
|
return;
|
|
|
|
|
}
|
|
|
|
|
if (N == 2.0f){
|
|
|
|
|
thrust::square<float> op;
|
|
|
|
|
thrust::transform(thrust::device,in1,in1+length,out,op);
|
|
|
|
|
return;
|
|
|
|
|
}
|
2023-11-01 14:31:29 +08:00
|
|
|
thrust::transform(thrust::device,in1,in1+length,out,PowOperator(N));
|
2023-10-31 14:35:29 +08:00
|
|
|
}
|
|
|
|
|
|
2023-11-09 17:56:13 +08:00
|
|
|
void thrustFill(float* aBegin, float* aEnd, float aValue)
|
|
|
|
|
{
|
|
|
|
|
thrust::fill(thrust::device, aBegin, aEnd, aValue);
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
void thrustFill(float* aBegin, float* aEnd, std::complex<float> aValue)
|
|
|
|
|
{
|
|
|
|
|
thrust::fill(thrust::device, (std::complex<float>*)aBegin, (std::complex<float>*)aEnd, aValue);
|
|
|
|
|
}
|
|
|
|
|
|