Add functions to 1D(nan, finite, padding, not)

This commit is contained in:
kradchen
2023-11-28 14:26:46 +08:00
parent 1716b139c4
commit 3438121b0c
4 changed files with 178 additions and 5 deletions

View File

@@ -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"<<std::endl;
return;
}
if (aMatrix.getDataSize()>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<CudaMatrix&&>(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)
{

View File

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

View File

@@ -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<float>::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<result1.getDataSize() * result1.getValueType(); ++i)
{
EXPECT_FLOAT_AE(result1[i], result2.getValue(i));
}
result1 = Aurora::auroraNot(hostMatrix);
result2 = Aurora::auroraNot(hostMatrix.toDeviceMatrix());
EXPECT_EQ(result2.getDataSize(), result1.getDataSize());
EXPECT_EQ(result2.getValueType(), result1.getValueType());
for(size_t i=0; i<result1.getDataSize() * result1.getValueType(); ++i)
{
EXPECT_FLOAT_AE(result1[i], result2.getValue(i));
}
}
TEST_F(Function1D_Cuda_Test, compareSet)
{

View File

@@ -560,8 +560,10 @@ TEST_F(Function1D_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);
Aurora::padding(ma,20,1);
auto result = ma.getData();
EXPECT_FLOAT_AE(ma[19],1);
EXPECT_FLOAT_AE(ma[20],1);
Aurora::padding(ma,10,-1);
EXPECT_FLOAT_AE(ma[10],-1);
}
TEST_F(Function1D_Test, convertfp16tofloat) {