From 3438121b0c1cee4900ead39f245cc2325e043952 Mon Sep 17 00:00:00 2001 From: kradchen Date: Tue, 28 Nov 2023 14:26:46 +0800 Subject: [PATCH] Add functions to 1D(nan, finite, padding, not) --- src/Function1D.cu | 72 ++++++++++++++++++++++++++++++++ src/Function1D.cuh | 27 ++++++++++-- test/Function1D_Cuda_Test.cpp | 78 +++++++++++++++++++++++++++++++++++ test/Function1D_Test.cpp | 6 ++- 4 files changed, 178 insertions(+), 5 deletions(-) diff --git a/src/Function1D.cu b/src/Function1D.cu index 7dd5825..ca12bbf 100644 --- a/src/Function1D.cu +++ b/src/Function1D.cu @@ -326,6 +326,78 @@ CudaMatrix Aurora::sign(const CudaMatrix&& aMatrix) return Aurora::CudaMatrix::fromRawData(data, aMatrix.getDimSize(0), aMatrix.getDimSize(1), aMatrix.getDimSize(2), aMatrix.getValueType()); } +void Aurora::nantoval(CudaMatrix& aMatrix,float val){ + auto lambda = [=] __host__ __device__ (const float& x){ + return ::isnan(x)?val:x; + }; + thrust::transform(thrust::device,aMatrix.getData(), + aMatrix.getData()+aMatrix.getDataSize(),aMatrix.getData(),lambda); +} + +CudaMatrix Aurora::isnan(const CudaMatrix& aMatrix){ + size_t size = aMatrix.getDataSize() * aMatrix.getValueType(); + float* data = nullptr; + + cudaMalloc((void**)&data, sizeof(float) * size); + auto lambda = [=] __host__ __device__ (const float& x){ + return ::isnan(x)?1.0:0; + }; + thrust::transform(thrust::device,aMatrix.getData(),aMatrix.getData()+aMatrix.getDataSize(), + data,lambda); + return Aurora::CudaMatrix::fromRawData(data, aMatrix.getDimSize(0), + aMatrix.getDimSize(1), aMatrix.getDimSize(2), aMatrix.getValueType()); +} + +CudaMatrix Aurora::isfinite(const CudaMatrix& aMatrix){ + size_t size = aMatrix.getDataSize() * aMatrix.getValueType(); + float* data = nullptr; + cudaMalloc((void**)&data, sizeof(float) * size); + auto lambda = [=] __host__ __device__ (const float& x){ + return ::isfinite(x)?1.0:0; + }; + thrust::transform(thrust::device,aMatrix.getData(),aMatrix.getData()+aMatrix.getDataSize(), + data,lambda); + return Aurora::CudaMatrix::fromRawData(data, aMatrix.getDimSize(0), + aMatrix.getDimSize(1), aMatrix.getDimSize(2), aMatrix.getValueType()); +} + +void Aurora::padding(CudaMatrix& aMatrix, int aIndex, float aValue){ + if(aMatrix.isNull() || !aMatrix.isVector() || aMatrix.isComplex()) + { + std::cerr<<"padding only support real vector"<aIndex){ + aMatrix.setValue(aIndex, aValue); + return; + } + //长度不足需补齐 + size_t size = (aIndex+1) ; + float* data = nullptr; + cudaMalloc((void**)&data, sizeof(float) * size); + cudaMemcpy(data, aMatrix.getData(), aMatrix.getDataSize(), cudaMemcpyDeviceToDevice); + thrust::fill_n(thrust::device, data+aMatrix.getDataSize(),size-aMatrix.getDataSize(),aValue); + aMatrix=CudaMatrix::fromRawData(data,size,1,1,aMatrix.getValueType()); +} + + +CudaMatrix Aurora::auroraNot(const CudaMatrix& aMatrix){ + return auroraNot(std::forward(aMatrix.deepCopy())); +} + +CudaMatrix Aurora::auroraNot(CudaMatrix&& aMatrix){ + size_t size = aMatrix.getDataSize() * aMatrix.getValueType(); + float* data = nullptr; + cudaMalloc((void**)&data, sizeof(float) * size); + auto lambda = [=] __host__ __device__ (const float& x){ + return x<=0?1.0:0; + }; + thrust::transform(thrust::device,aMatrix.getData(),aMatrix.getData()+aMatrix.getDataSize(), + data,lambda); + 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) { diff --git a/src/Function1D.cuh b/src/Function1D.cuh index 5db2eb7..8d3dfd6 100644 --- a/src/Function1D.cuh +++ b/src/Function1D.cuh @@ -77,11 +77,32 @@ namespace Aurora CudaMatrix intersect(const CudaMatrix& aMatrix1, const CudaMatrix& aMatrix2); - CudaMatrix intersect(const CudaMatrix& aMatrix1, const CudaMatrix& aMatrix2, CudaMatrix& aIa); + CudaMatrix intersect(const CudaMatrix& aMatrix1, const CudaMatrix& aMatrix2, CudaMatrix& aIa); /** + * 将所有nan值设置为特定值 + * @attention 直接在原数据上进行修改！ + * @param aMatrix 向量 + * @param val 指定值 + */ + void nantoval(CudaMatrix& aMatrix,float val); + + CudaMatrix isnan(const CudaMatrix& aMatrix); + + CudaMatrix isfinite(const CudaMatrix& aMatrix); + + /** + * 使用特定值补齐矩阵，默认为设置原数据矩阵数据到制定长度索引的所有值为制定值 + * @attention 直接在原数据上进行修改！不支持复数！ + * @param aMatrix 向量 + * @param aIndex 长度索引 + * @param aValue 指定值 + */ + void padding(CudaMatrix& aMatrix, int aIndex, float aValue); + + CudaMatrix auroraNot(const CudaMatrix& aMatrix); + + CudaMatrix auroraNot(CudaMatrix&& aMatrix); // ------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); diff --git a/test/Function1D_Cuda_Test.cpp b/test/Function1D_Cuda_Test.cpp index 20fb373..f5d0664 100644 --- a/test/Function1D_Cuda_Test.cpp +++ b/test/Function1D_Cuda_Test.cpp @@ -319,6 +319,84 @@ TEST_F(Function1D_Cuda_Test, log) } } +TEST_F(Function1D_Cuda_Test, nanAndFinite) +{ + float n=std::nan(""); + float infinite=std::numeric_limits::infinity(); + + float* data = new float[9]{infinite,3,n,2,2,1,4,4,7}; + auto matrix = Aurora::Matrix::fromRawData(data, 9,1,1).toDeviceMatrix(); + auto r = Aurora::isnan(matrix); + EXPECT_FLOAT_EQ(r.getValue(0),0); + EXPECT_FLOAT_EQ(r.getValue(1),0); + EXPECT_FLOAT_EQ(r.getValue(2),1); + + r = Aurora::isfinite(matrix); + EXPECT_FLOAT_EQ(r.getValue(0),0); + EXPECT_FLOAT_EQ(r.getValue(1),1); + EXPECT_FLOAT_EQ(r.getValue(2),0); + + Aurora::nantoval(matrix,1); + EXPECT_FLOAT_EQ(matrix.getValue(2),1); + + float* data2 = new float[1]{n}; + auto matrix2 = Aurora::Matrix::fromRawData(data2, 1,1,1).toDeviceMatrix(); + r = Aurora::isnan(matrix2); + EXPECT_FLOAT_EQ(r.getValue(0),1); + + r = Aurora::isfinite(matrix2); + EXPECT_FLOAT_EQ(r.getValue(0),0); + + Aurora::nantoval(matrix2,1); + EXPECT_FLOAT_EQ(matrix2.getValue(0),1); +} + +TEST_F(Function1D_Cuda_Test, isnanOrfinite) +{ + float *input = new float[18]{10,2,1,3,4,4,16,3,1,2,15,-2,1,-3,4,-4,1,-3}; + auto ma = Aurora::Matrix::fromRawData(input,18,1,1).toDeviceMatrix(); + Aurora::padding(ma,20,1); + EXPECT_FLOAT_AE(ma.getValue(19),1); + EXPECT_FLOAT_AE(ma.getValue(20),1); + Aurora::padding(ma,10,-1); + EXPECT_FLOAT_AE(ma.getValue(10),-1); +} + +TEST_F(Function1D_Cuda_Test, padding) +{ + float *input = new float[18]{10,2,1,3,4,4,16,3,1,2,15,-2,1,-3,4,-4,1,-3}; + auto ma = Aurora::Matrix::fromRawData(input,18,1,1).toDeviceMatrix(); + Aurora::padding(ma,20,1); + EXPECT_FLOAT_AE(ma.getValue(19),1); + EXPECT_FLOAT_AE(ma.getValue(20),1); + Aurora::padding(ma,10,-1); + EXPECT_FLOAT_AE(ma.getValue(10),-1); +} + +TEST_F(Function1D_Cuda_Test, not) +{ + Aurora::Matrix hostMatrix = Aurora::Matrix::fromRawData(new float[8]{-1.1,2.2,-3.3,0,-5.5,6.6,0.7,-0.8}, 2,4); + Aurora::CudaMatrix deviceMatrix = hostMatrix.toDeviceMatrix(); + + auto result1 = Aurora::auroraNot(hostMatrix); + auto result2 = Aurora::auroraNot(deviceMatrix); + EXPECT_EQ(result2.getDataSize(), result1.getDataSize()); + EXPECT_EQ(result2.getValueType(), result1.getValueType()); + for(size_t i=0; i