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

Add cuda matrix function1d function compare set

Browse Source
This commit is contained in:
kradchen
2023-11-28 09:46:11 +08:00
parent a598de6ea3
commit 7d879c17d4
4 changed files with 639 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

267
src/Function1D.cu
View File

@@ -6,9 +6,12 @@
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/complex.h>
#include <cuda_runtime.h>
using namespace Aurora;
using namespace thrust::placeholders;
namespace
{
@@ -247,7 +250,7 @@ CudaMatrix Aurora::abs(const CudaMatrix& aMatrix)
float* data = nullptr;
cudaMalloc((void**)&data, sizeof(float) * size);
int blocksPerGrid = (size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
absKernel<<<blocksPerGrid, THREADS_PER_BLOCK>>>(aMatrix.getData(), data, size, aMatrix.isComplex());
absKernel<<<blocksPerGrid,THREADS_PER_BLOCK >>>(aMatrix.getData(), data, size, aMatrix.isComplex());
cudaDeviceSynchronize();
return Aurora::CudaMatrix::fromRawData(data, aMatrix.getDimSize(0), aMatrix.getDimSize(1), aMatrix.getDimSize(2));
}
@@ -313,6 +316,268 @@ CudaMatrix Aurora::sign(const CudaMatrix&& aMatrix)
return Aurora::CudaMatrix::fromRawData(data, aMatrix.getDimSize(0), aMatrix.getDimSize(1), aMatrix.getDimSize(2), aMatrix.getValueType());
}
void Aurora::compareSet(CudaMatrix& aValueMatrix,float compareValue, float newValue,CompareOp op)
{
switch (op)
{
case GT:
{
auto lambda = [=] __host__ __device__ (const float& x){
return x>compareValue?newValue:x;
};
thrust::transform(thrust::device,aValueMatrix.getData(),aValueMatrix.getData()+aValueMatrix.getDataSize(),aValueMatrix.getData(),lambda);
break;
}
case NG:{
auto lambda = [=] __host__ __device__ (const float& x){
return x<=compareValue?newValue:x;
};
thrust::transform(thrust::device,aValueMatrix.getData(),aValueMatrix.getData()+aValueMatrix.getDataSize(),aValueMatrix.getData(),lambda);
break;
}
case EQ:{
auto lambda = [=] __host__ __device__ (const float& x){
return x==compareValue?newValue:x;
};
thrust::transform(thrust::device,aValueMatrix.getData(),aValueMatrix.getData()+aValueMatrix.getDataSize(),aValueMatrix.getData(),lambda);
break;
}
case NE:{
auto lambda = [=] __host__ __device__ (const float& x){
return x!=compareValue?newValue:x;
};
thrust::transform(thrust::device,aValueMatrix.getData(),aValueMatrix.getData()+aValueMatrix.getDataSize(),aValueMatrix.getData(),lambda);
break;
}
case NL:{
auto lambda = [=] __host__ __device__ (const float& x){
return x>=compareValue?newValue:x;
};
thrust::transform(thrust::device,aValueMatrix.getData(),aValueMatrix.getData()+aValueMatrix.getDataSize(),aValueMatrix.getData(),lambda);
break;
}
case LT:{
auto lambda = [=] __host__ __device__ (const float& x){
return x<compareValue?newValue:x;
};
thrust::transform(thrust::device,aValueMatrix.getData(),aValueMatrix.getData()+aValueMatrix.getDataSize(),aValueMatrix.getData(),lambda);
break;
}
default:
break;
}
}
void Aurora::compareSet(CudaMatrix& aValueMatrix,CudaMatrix& aCompareMatrix,float compareValue, float newValue,CompareOp op){
switch (op)
{
case GT:
{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x>compareValue?newValue:y;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aValueMatrix.getDataSize(),
aValueMatrix.getData(), aValueMatrix.getData(),
lambda);
break;
}
case NG:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x<=compareValue?newValue:y;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aValueMatrix.getDataSize(),
aValueMatrix.getData(), aValueMatrix.getData(),
lambda);
break;
}
case EQ:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x==compareValue?newValue:y;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aValueMatrix.getDataSize(),
aValueMatrix.getData(), aValueMatrix.getData(),
lambda);
break;
}
case NE:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x!=compareValue?newValue:y;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aValueMatrix.getDataSize(),
aValueMatrix.getData(), aValueMatrix.getData(),
lambda);
break;
}
case NL:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x>=compareValue?newValue:y;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aValueMatrix.getDataSize(),
aValueMatrix.getData(), aValueMatrix.getData(),
lambda);
break;
}
case LT:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x<compareValue?newValue:y;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aValueMatrix.getDataSize(),
aValueMatrix.getData(), aValueMatrix.getData(),
lambda);
break;
}
default:
break;
}
}
void Aurora::compareSet(CudaMatrix& aDesAndCompareMatrix,CudaMatrix& aOtherCompareMatrix, float newValue,CompareOp op){
switch (op)
{
case GT:
{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x>y?newValue:x;
};
thrust::transform(thrust::device,aDesAndCompareMatrix.getData(),
aDesAndCompareMatrix.getData()+aDesAndCompareMatrix.getDataSize(),
aOtherCompareMatrix.getData(), aDesAndCompareMatrix.getData(),
lambda);
break;
}
case NG:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x<=y?newValue:x;
};
thrust::transform(thrust::device,aDesAndCompareMatrix.getData(),
aDesAndCompareMatrix.getData()+aDesAndCompareMatrix.getDataSize(),
aOtherCompareMatrix.getData(), aDesAndCompareMatrix.getData(),
lambda);
break;
}
case EQ:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x==y?newValue:x;
};
thrust::transform(thrust::device,aDesAndCompareMatrix.getData(),
aDesAndCompareMatrix.getData()+aDesAndCompareMatrix.getDataSize(),
aOtherCompareMatrix.getData(), aDesAndCompareMatrix.getData(),
lambda);
break;
}
case NE:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x!=y?newValue:x;
};
thrust::transform(thrust::device,aDesAndCompareMatrix.getData(),
aDesAndCompareMatrix.getData()+aDesAndCompareMatrix.getDataSize(),
aOtherCompareMatrix.getData(), aDesAndCompareMatrix.getData(),
lambda);
break;
}
case NL:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x>=y?newValue:x;
};
thrust::transform(thrust::device,aDesAndCompareMatrix.getData(),
aDesAndCompareMatrix.getData()+aDesAndCompareMatrix.getDataSize(),
aOtherCompareMatrix.getData(), aDesAndCompareMatrix.getData(),
lambda);
break;
}
case LT:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x<y?newValue:x;
};
thrust::transform(thrust::device,aDesAndCompareMatrix.getData(),
aDesAndCompareMatrix.getData()+aDesAndCompareMatrix.getDataSize(),
aOtherCompareMatrix.getData(), aDesAndCompareMatrix.getData(),
lambda);
break;
}
default:
break;
}
}
void Aurora::compareSet(CudaMatrix& aCompareMatrix,float compareValue, CudaMatrix& aNewValueMatrix,CompareOp op)
{
switch (op)
{
case GT:
{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x>compareValue?y:x;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aCompareMatrix.getDataSize(),
aNewValueMatrix.getData(), aCompareMatrix.getData(),
lambda);
break;
}
case NG:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x<=compareValue?y:x;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aCompareMatrix.getDataSize(),
aNewValueMatrix.getData(), aCompareMatrix.getData(),
lambda);
break;
}
case EQ:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x==compareValue?y:x;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aCompareMatrix.getDataSize(),
aNewValueMatrix.getData(), aCompareMatrix.getData(),
lambda);
break;
}
case NE:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x!=compareValue?y:x;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aCompareMatrix.getDataSize(),
aNewValueMatrix.getData(), aCompareMatrix.getData(),
lambda);
break;
}
case NL:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x>=compareValue?y:x;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aCompareMatrix.getDataSize(),
aNewValueMatrix.getData(), aCompareMatrix.getData(),
lambda);
break;
}
case LT:{
auto lambda = [=] __host__ __device__ (const float& x, const float& y){
return x<compareValue?y:x;
};
thrust::transform(thrust::device,aCompareMatrix.getData(),
aCompareMatrix.getData()+aCompareMatrix.getDataSize(),
aNewValueMatrix.getData(), aCompareMatrix.getData(),
lambda);
break;
}
default:
break;
}
}
__global__ void repMatKernel(float* aInputData, float* aOutput, unsigned int aInputSize, bool aIsComplex)
{
unsigned int idX = blockIdx.x * blockDim.x + threadIdx.x;

13
src/Function1D.cuh
View File

@@ -2,6 +2,7 @@
#define AURORA_CUDA_FUNCTION1D_H
#include "CudaMatrix.h"
#include "Function1D.h"
namespace Aurora
{
@@ -33,6 +34,8 @@ namespace Aurora
CudaMatrix sqrt(const CudaMatrix&& aMatrix);
CudaMatrix abs(const CudaMatrix& aMatrix);
CudaMatrix abs2(const CudaMatrix& aMatrix);
CudaMatrix abs(const CudaMatrix&& aMatrix);
@@ -47,6 +50,16 @@ namespace Aurora
CudaMatrix repmat3d(const CudaMatrix& aMatrix,int aRowTimes, int aColumnTimes, int aSliceTimes);
CudaMatrix log(const CudaMatrix& aMatrix, int aBaseNum = -1);
// ------compareSet----------------------------------------------------
void compareSet(CudaMatrix& aValueMatrix,float compareValue, float newValue,CompareOp op);
void compareSet(CudaMatrix& aValueMatrix,CudaMatrix& aCompareMatrix,float compareValue, float newValue,CompareOp op);
void compareSet(CudaMatrix& aDesAndCompareMatrix,CudaMatrix& aOtherCompareMatrix, float newValue,CompareOp op);
void compareSet(CudaMatrix& aCompareMatrix,float compareValue, CudaMatrix& aNewValueMatrix,CompareOp op);
}
#endif //AURORA_CUDA_FUNCTION1D_H