sunwen/Aurora
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Cuda matrix compare and value getter and setter

Browse Source
This commit is contained in:
kradchen
2023-11-21 13:13:28 +08:00
parent 4edb2d133d
commit aaf8c1b193
5 changed files with 474 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

190
src/CudaMatrixPrivate.cu
View File

@@ -5,9 +5,9 @@
#include <thrust/execution_policy.h>
using namespace thrust::placeholders;
struct PowOperator: public thrust::unary_function<float, float>{
struct PowOp: public thrust::unary_function<float, float>{
float exponent;
PowOperator(float v):exponent(v) {}
PowOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
@@ -18,6 +18,114 @@ struct PowOperator: public thrust::unary_function<float, float>{
}
};
struct CompareGOp: public thrust::unary_function<float, float>{
float exponent;
CompareGOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
__host__ __device__
float operator()(const float& x) {
return (exponent<x?1.0:.0);
}
};
struct CompareGEOp: public thrust::unary_function<float, float>{
float exponent;
CompareGEOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
__host__ __device__
float operator()(const float& x) {
return (exponent<=x?1.0:.0);
}
};
struct CompareEOp: public thrust::unary_function<float, float>{
float exponent;
CompareEOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
__host__ __device__
float operator()(const float& x) {
return (exponent==x?1.0:.0);
}
};
struct CompareNEOp: public thrust::unary_function<float, float>{
float exponent;
CompareNEOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
__host__ __device__
float operator()(const float& x) {
return (exponent!=x?1.0:.0);
}
};
struct CompareLOp: public thrust::unary_function<float, float>{
float exponent;
CompareLOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
__host__ __device__
float operator()(const float& x) {
return (exponent>x?1.0:.0);
}
};
struct CompareLEOp: public thrust::unary_function<float, float>{
float exponent;
CompareLEOp(float v):exponent(v) {}
void setExponent(float v){
exponent = v;
}
__host__ __device__
float operator()(const float& x) {
return (exponent>=x?1.0:.0);
}
};
struct CompareAGOp{
__host__ __device__
float operator()(const float& x,const float& y) {
return x>y?1:0;
}
};
struct CompareAGEOp{
__host__ __device__
float operator()(const float& x,const float& y) {
return x>=y?1:0;
}
};
struct CompareAEOp{
__host__ __device__
float operator()(const float& x,const float& y) {
return x==y?1:0;
}
};
struct CompareANEOp{
__host__ __device__
float operator()(const float& x,const float& y) {
return x!=y?1:0;
}
};
void unaryAdd(float* in1, float* in2, float* out, unsigned long length)
{
thrust::plus<float> op;
@@ -88,7 +196,83 @@ void unaryPow(float* in1, float N,float* out, unsigned long length){
thrust::transform(thrust::device,in1,in1+length,out,op);
return;
}
thrust::transform(thrust::device,in1,in1+length,out,PowOperator(N));
thrust::transform(thrust::device,in1,in1+length,out,PowOp(N));
}
void unaryCompare(float* in1, const float& in2, float* out, unsigned long length, int type){
switch (type)
{
case G:
thrust::transform(thrust::device,in1,in1+length,out,CompareGOp(in2));
break;
case GE:
thrust::transform(thrust::device,in1,in1+length,out,CompareGEOp(in2));
break;
case E:
thrust::transform(thrust::device,in1,in1+length,out,CompareEOp(in2));
break;
case NE:
thrust::transform(thrust::device,in1,in1+length,out,CompareNEOp(in2));
break;
case LE:
thrust::transform(thrust::device,in1,in1+length,out,CompareLEOp(in2));
break;
case L:
thrust::transform(thrust::device,in1,in1+length,out,CompareLOp(in2));
break;
default:
break;
}
}
void unaryCompare(const float& in1, float* in2, float* out, unsigned long length, int type){
switch (type)
{
case G:
thrust::transform(thrust::device,in2,in2+length,out,CompareLOp(in1));
break;
case GE:
thrust::transform(thrust::device,in2,in2+length,out,CompareLEOp(in1));
break;
case E:
thrust::transform(thrust::device,in2,in2+length,out,CompareEOp(in1));
break;
case NE:
thrust::transform(thrust::device,in2,in2+length,out,CompareNEOp(in1));
break;
case LE:
thrust::transform(thrust::device,in2,in2+length,out,CompareGEOp(in1));
break;
case L:
thrust::transform(thrust::device,in2,in2+length,out,CompareGOp(in1));
break;
default:
break;
}
}
void unaryCompare(float* in1, float* in2, float* out, unsigned long length, int type){
switch (type)
{
case G:
thrust::transform(thrust::device,in1,in1+length,in2,out,CompareAGOp());
break;
case GE:
thrust::transform(thrust::device,in1,in1+length,in2,out,CompareAGEOp());
break;
case E:
thrust::transform(thrust::device,in1,in1+length,in2,out,CompareAEOp());
break;
case NE:
thrust::transform(thrust::device,in1,in1+length,in2,out,CompareANEOp());
break;
case LE:
thrust::transform(thrust::device,in2,in2+length,in1,out,CompareAGEOp());
break;
case L:
thrust::transform(thrust::device,in2,in2+length,in1,out,CompareAGOp());
break;
default:
break;
}
}
void thrustFill(float* aBegin, float* aEnd, float aValue)