diff --git a/CMakeLists.txt b/CMakeLists.txt index 16fae3f..fa5b445 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,13 +1,15 @@ cmake_minimum_required(VERSION 3.16) project(Aurora) -set(CMAKE_CXX_STANDARD 14) +set(CMAKE_CXX_STANDARD 17) set(CMAKE_INCLUDE_CURRENT_DIR ON) set(Aurora_USE_CUDA ON) if (Aurora_USE_CUDA) set(CMAKE_CUDA_COMPILER /usr/local/cuda/bin/nvcc) + + enable_language(CUDA) find_package(CUDAToolkit REQUIRED) add_definitions(-DUSE_CUDA) @@ -44,7 +46,7 @@ if (Aurora_USE_CUDA) target_include_directories(Aurora PRIVATE ./src /usr/local/cuda/include) set_target_properties(Aurora PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_compile_options(Aurora PRIVATE $<$: - -arch=sm_75 + -arch=sm_75 --expt-extended-lambda >) target_link_libraries(Aurora PRIVATE ${CUDA_RUNTIME_LIBRARY} CUDA::cufft CUDA::cudart) endif(Aurora_USE_CUDA) @@ -70,7 +72,7 @@ if (Aurora_USE_CUDA) target_include_directories(Aurora_Test PRIVATE ./src /usr/local/cuda/include) set_target_properties(Aurora_Test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_compile_options(Aurora_Test PRIVATE $<$: - -arch=sm_75 + -arch=sm_75 --expt-extended-lambda >) target_link_libraries(Aurora_Test PRIVATE ${CUDA_RUNTIME_LIBRARY} CUDA::cufft CUDA::cudart) endif(Aurora_USE_CUDA) diff --git a/src/Function1D.cu b/src/Function1D.cu index d2339f8..fbfd7e0 100644 --- a/src/Function1D.cu +++ b/src/Function1D.cu @@ -6,9 +6,12 @@ #include #include #include +#include +#include #include 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<<>>(aMatrix.getData(), data, size, aMatrix.isComplex()); + absKernel<<>>(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 xcompareValue?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 xy?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 xcompareValue?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