Add cuda uniqueByRows and unitest.

src/Function1D.cu

@@ -7,6 +7,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdlib>
+#include <sys/types.h>
 #include <thrust/device_vector.h>
 #include <thrust/transform.h>
 #include <thrust/iterator/constant_iterator.h>
@@ -1480,3 +1481,112 @@ CudaMatrix Aurora::createCudaVectorMatrix(float aStartValue, float aStepValue, f
 
     return Matrix::copyFromRawData(matrixData.data(), 1, matrixData.size()).toDeviceMatrix();
 }
+
+struct compareMatrixByRows
+{
+    compareMatrixByRows(unsigned int aSize)
+        : mSize(aSize)
+    {
+    };
+    unsigned int mSize;
+    __host__ __device__
+    bool operator()(const float* aVector1, const float* aVector2) const
+    {
+        for(unsigned int i=0; i<mSize; ++i)
+        {
+            if(aVector1[i] < aVector2[i])
+            {
+                return true;
+            }
+            else if(aVector1[i] > aVector2[i])
+            {
+                return false;
+            }
+        }
+        return false;
+    }
+};
+
+struct uniqueMatrixByRows
+{
+    uniqueMatrixByRows(unsigned int aSize)
+        : mSize(aSize)
+    {
+    };
+    unsigned int mSize;
+    __host__ __device__
+    bool operator()(const float* aVector1, const float* aVector2) const
+    {
+        for(unsigned int i=0; i<mSize; ++i)
+        {
+            if(aVector1[i] != aVector2[i])
+            {
+                return false;
+            }
+        }
+        return true;
+    }
+};
+
+struct equalMatrixByRows
+{
+    equalMatrixByRows(unsigned int aSize, float* aValue)
+        : mSize(aSize)
+        , mValue(aValue)
+    {
+    };
+    unsigned int mSize;
+    float* mValue;
+    __host__ __device__
+    bool operator()(const float* aVector) const
+    {
+        for(unsigned int i=0; i<mSize; ++i)
+        {
+            if(aVector[i] != mValue[i])
+            {
+                return false;
+            }
+        }
+        return true;
+    }
+};
+
+
+
+CudaMatrix Aurora::uniqueByRows(const CudaMatrix& aMatrix, CudaMatrix& aIndexResult)
+{
+    CudaMatrix transposeMatrix = transpose(aMatrix);
+    size_t rows = transposeMatrix.getDimSize(0);
+    size_t columns = transposeMatrix.getDimSize(1);
+    thrust::device_vector<float*> vector(columns);
+    float* indexResult = new float[columns];
+    std::vector<float*> vectorCopy(columns);
+    for(unsigned int i=0; i<columns; ++i)
+    {
+        vector[i] = transposeMatrix.getData() + i*rows;
+        vectorCopy[i] = vector[i];
+    }
+    thrust::sort(thrust::device, vector.begin(), vector.begin() + columns, compareMatrixByRows(rows));
+    auto end = thrust::unique(thrust::device, vector.begin(), vector.end(), uniqueMatrixByRows(rows));
+
+    for(unsigned int i=0; i<transposeMatrix.getDimSize(1); ++i)
+    {
+        auto index = thrust::find_if(thrust::device, vector.begin(), end, equalMatrixByRows(rows, vectorCopy[i]));
+        indexResult[i] = index - vector.begin();
+    }
+    float* indexData = nullptr;
+    cudaMalloc((void**)&indexData, sizeof(float) * transposeMatrix.getDimSize(1));
+    cudaMemcpy(indexData, indexResult, sizeof(float) * transposeMatrix.getDimSize(1), cudaMemcpyHostToDevice);
+    aIndexResult = CudaMatrix::fromRawData(indexData, transposeMatrix.getDimSize(1));
+
+    columns = end - vector.begin();
+    float* resultData = nullptr;
+    cudaMalloc((void**)&resultData, sizeof(float) * rows * columns);
+    for(unsigned int i=0; i<columns; ++i)
+    {
+        cudaMemcpy(resultData + i*rows, vector[i], sizeof(float) * rows, cudaMemcpyDeviceToDevice);
+    }
+
+    delete [] indexResult;
+    return transpose(CudaMatrix::fromRawData(resultData, rows, columns));
+}

src/Function1D.cuh

@@ -87,6 +87,8 @@ namespace Aurora
 
     CudaMatrix createCudaVectorMatrix(float aStartValue, float aStepValue, float aEndValue);
 
+    CudaMatrix uniqueByRows(const CudaMatrix& aMatrix, CudaMatrix& aIndexResult);
+
      /**
      * 将所有nan值设置为特定值
      * @attention 直接在原数据上进行修改！

test/Function1D_Cuda_Test.cpp

@@ -1121,3 +1121,31 @@ TEST_F(Function1D_Cuda_Test, deleteColumn) {
         EXPECT_FLOAT_AE(result1[i], result2[i]);
     }
 }
+
+TEST_F(Function1D_Cuda_Test, uniqueByRows) {
+    auto matrixHost = Aurora::Matrix::fromRawData(new float[12]{6,2,2,2,4,7,7,7,1,8,9,8},4,3,1,Aurora::Normal);
+    Aurora::Matrix indexHost;
+    auto matrixDevice = matrixHost.toDeviceMatrix();
+    Aurora::CudaMatrix indexDevice;
+    auto result1 = Aurora::uniqueByRows(matrixHost, indexHost);
+    auto result2 = Aurora::uniqueByRows(matrixDevice, indexDevice).toHostMatrix();
+    auto indexResult1 = indexHost;
+    auto indexResult2 = indexDevice.toHostMatrix();
+    EXPECT_FLOAT_AE(result1.getDimSize(0),result2.getDimSize(0));
+    EXPECT_FLOAT_AE(result1.getDimSize(1),result2.getDimSize(1));
+    EXPECT_FLOAT_AE(result1.getDataSize(),result2.getDataSize());
+
+    EXPECT_FLOAT_AE(indexResult1.getDimSize(0),indexResult2.getDimSize(0));
+    EXPECT_FLOAT_AE(indexResult1.getDimSize(1),indexResult2.getDimSize(1));
+    EXPECT_FLOAT_AE(indexResult1.getDataSize(),indexResult2.getDataSize());
+    
+    for(int i=0;i<result1.getDataSize();++i)
+    {
+        EXPECT_FLOAT_AE(result1[i], result2[i]);
+    }
+
+    for(int i=0;i<result1.getDataSize();++i)
+    {
+        EXPECT_FLOAT_AE(indexResult1[i], indexResult2[i]);
+    }
+}