diff --git a/.github/workflows/cmake-linux-amd64.yml b/.github/workflows/cmake-linux-amd64.yml index 8e6f91ce..ba22d2cd 100644 --- a/.github/workflows/cmake-linux-amd64.yml +++ b/.github/workflows/cmake-linux-amd64.yml @@ -18,12 +18,14 @@ jobs: matrix: #host,cuda,cuda_version include: - host_compiler: "g++-13" - cuda_toolkit: "13.3" + cuda_version: "13.3" + opencv_version: "opencv-4.14.0-cuda133" - host_compiler: "clang++-21" - cuda_toolkit: "13.3" + cuda_version: "13.3" + opencv_version: "opencv-4.14.0-cuda133" steps: - - uses: actions/checkout@v4 + - uses: actions/checkout@v6 with: submodules: recursive - name: Set reusable strings @@ -31,13 +33,13 @@ jobs: id: strings run: | echo "build-output-dir=${{github.workspace}}/build" >> "$GITHUB_OUTPUT" - + echo "PATH=$HOME/cmake-4.4.0-linux-x86_64/bin/:$PATH" >> "$GITHUB_ENV" + echo CUDACXX="/usr/local/cuda-${{matrix.cuda_version}}/bin/nvcc" >> "$GITHUB_ENV" + echo CC="${{matrix.host_compiler}}" >> "$GITHUB_ENV" + echo CXX="${{matrix.host_compiler}}">> "$GITHUB_ENV" + - name: Configure CMake - run: | - export PATH=/home/cudeiro/cmake-4.2.1-linux-x86_64/bin:/usr/lib/llvm-21/bin/:$PATH - export CUDACXX="/usr/local/cuda-${{matrix.cuda_version}}/bin/nvcc" - export CC="${{matrix.host_compiler}}" - export CXX="${{matrix.host_compiler}}" + run: | cmake -G "Ninja" -B ${{steps.strings.outputs.build-output-dir}} -DOPENCV_DIR="/usr/local/${{matrix.opencv_version}}" -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} - name: Build diff --git a/.github/workflows/cmake-linux-arm64.yml b/.github/workflows/cmake-linux-arm64.yml index 92ea8e64..a4543ec3 100644 --- a/.github/workflows/cmake-linux-arm64.yml +++ b/.github/workflows/cmake-linux-arm64.yml @@ -18,13 +18,14 @@ jobs: matrix: #host,cuda,cuda_version include: - host_compiler: "g++-13" - cuda_toolkit: "13.3" + cuda_version: "13.3" + opencv_version: "opencv-4.14.0-cuda133" - host_compiler: "clang++-21" - cuda_toolkit: "13.3" - + cuda_version: "13.3" + opencv_version: "opencv-4.14.0-cuda133" # steps: - - uses: actions/checkout@v4 + - uses: actions/checkout@v6 with: submodules: recursive @@ -32,14 +33,16 @@ jobs: # Turn repeated input strings (such as the build output directory) into step outputs. These step outputs can be used throughout the workflow file. id: strings run: | - echo "build-output-dir=${{github.workspace}}/build" >> "$GITHUB_OUTPUT" - + echo "build-output-dir=${{github.workspace}}/build" >> "$GITHUB_OUTPUT" + echo "PATH=$HOME/cmake-4.4.0-linux-aarch64/bin/:$PATH" >> "$GITHUB_ENV" + echo CUDACXX="/usr/local/cuda-${{matrix.cuda_version}}/bin/nvcc" >> "$GITHUB_ENV" + echo CC="${{matrix.host_compiler}}" >> "$GITHUB_ENV" + echo CXX="${{matrix.host_compiler}}">> "$GITHUB_ENV" + + + - name: Configure CMake run: | - export PATH=/home/cudeiro/cmake-4.2.1-linux-aarch64/bin:/usr/lib/llvm-21/bin/:$PATH - export CUDACXX="/usr/local/cuda-${{matrix.cuda_version}}/bin/nvcc" - export CC="${{matrix.host_compiler}}" - export CXX="${{matrix.host_compiler}}" cmake -G "Ninja" -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} - name: Build diff --git a/.github/workflows/cmake-windows-amd64.yml b/.github/workflows/cmake-windows-amd64.yml index 0ac7831f..66c48ba2 100644 --- a/.github/workflows/cmake-windows-amd64.yml +++ b/.github/workflows/cmake-windows-amd64.yml @@ -19,17 +19,20 @@ jobs: - host_compiler: "cl" msvc_env_version: "14.44" cuda_version: "13.0" + opencv_version: "opencv-4.14.0.cuda13.0.msvc1444" - host_compiler: "cl" msvc_env_version: "14.51" - cuda_version: "13.3" + cuda_version: "13.3" + opencv_version: "opencv-4.14.0.cuda13.3.msvc1451" - host_compiler: "clang-cl" msvc_env_version: "14.51" cuda_version: "13.3" + opencv_version: "opencv-4.14.0.cuda13.3.msvc1451" steps: - - uses: actions/checkout@v4 + - uses: actions/checkout@v6 with: submodules: recursive diff --git a/fkl b/fkl index 98acf9b5..89d4288b 160000 --- a/fkl +++ b/fkl @@ -1 +1 @@ -Subproject commit 98acf9b560cce33ca1e76fd97d5dedf40c09d43d +Subproject commit 89d4288bc28bda87c0bde3dc51d818589ec43b16 diff --git a/tests/batchread/test_circularbatchread_x_write3D.cu b/tests/batchread/test_circularbatchread_x_write3D.cu index 34138d5f..5cc2693e 100644 --- a/tests/batchread/test_circularbatchread_x_write3D.cu +++ b/tests/batchread/test_circularbatchread_x_write3D.cu @@ -24,14 +24,10 @@ #include #include -template +template bool testCircularTensorcvGS() { using TensorOT = typename fk::VectorTraits::base; - constexpr uint BATCH = 15; - constexpr uint WIDTH =128; - constexpr uint HEIGHT = 128; constexpr uint COLOR_PLANES = CV_MAT_CN(IT); - constexpr int ITERS = 100; cvGS::CircularTensor myTensor(WIDTH, HEIGHT); fk::Tensor h_myTensor(WIDTH, HEIGHT, BATCH, COLOR_PLANES, fk::MemType::HostPinned); @@ -263,14 +259,37 @@ bool testOldestFirstCircularTensorcvGS_noSplit() { return correct; } -int launch() { - int returnValue = 0; - if (testCircularTensorcvGS()) { - std::cout << "testCircularTensorcvGS OK" << std::endl; +template +bool launchTestCircularTensorcvGS() { + if (testCircularTensorcvGS()) { + std::cout << "testCircularTensorcvGS<" << WIDTH << ", " << HEIGHT << ", " << BATCH << ", " << ITERS << ", " << IT << ", " << OT << "> OK" << std::endl; + return true; } else { - std::cout << "testCircularTensorcvGS Failed!" << std::endl; - returnValue = -1; + std::cout << "testCircularTensorcvGS<" << WIDTH << ", " << HEIGHT << ", " << BATCH << ", " << ITERS << ", " << IT << ", " << OT << "> Failed!" << std::endl; + return false; } +} + +int launch() { + int returnValue = 0; + + bool correct{true}; + correct &= launchTestCircularTensorcvGS<128, 128, 2, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 3, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 4, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 5, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 6, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 7, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 8, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 9, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 10, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 11, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 12, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 13, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 14, 100, CV_8UC3, CV_32FC3>(); + correct &= launchTestCircularTensorcvGS<128, 128, 15, 100, CV_8UC3, CV_32FC3>(); + returnValue = correct ? returnValue : -1; + if (testTransposedCircularTensorcvGS()) { std::cout << "testTransposedCircularTensorcvGS OK" << std::endl; } else { diff --git a/tests/resize/test_fused_resize.cu b/tests/resize/test_fused_resize.cu index 91dc1a15..52f2d2fe 100644 --- a/tests/resize/test_fused_resize.cu +++ b/tests/resize/test_fused_resize.cu @@ -17,8 +17,9 @@ #include "tests/testsCommon.cuh" #include -#include -#include +#include +#include +#include struct PerPlaneSequenceSelector { FK_HOST_DEVICE_FUSE uint at(const uint& index) { @@ -27,148 +28,107 @@ struct PerPlaneSequenceSelector { }; void testComputeWhatYouSeePlusHorizontalFusion(char* buffer, const uint& NUM_ELEMS_X, const uint& NUM_ELEMS_Y) { + using namespace fk; + Stream fk_stream; - cudaStream_t stream; - gpuErrchk(cudaStreamCreate(&stream)); - fk::Stream fk_stream{stream}; + constexpr Size down(1920, 1080); - constexpr fk::Size down(1920, 1080); - cv::Mat h_result(down.height, down.width, CV_8UC4); - cv::Mat nv12Image(cv::Size(NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2)), CV_8UC1, buffer); + Image nv12Image(NUM_ELEMS_X, NUM_ELEMS_Y); + memcpy(nv12Image.getData().ptrPinned().data, buffer, NUM_ELEMS_X * (NUM_ELEMS_Y + (NUM_ELEMS_Y / 2))); + Ptr2D rgbaImage(down.width, down.height); + Ptr2D rgbaImageBig(NUM_ELEMS_X, NUM_ELEMS_Y); + nv12Image.upload(fk_stream); - uchar* d_dataSource; - size_t sourcePitch; - gpuErrchk(cudaMallocPitch(&d_dataSource, &sourcePitch, NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2))); - fk::RawImage d_nv12Image{ {d_dataSource, {NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2), (uint)sourcePitch}}, NUM_ELEMS_X, NUM_ELEMS_Y}; - fk::Ptr2D d_rgbaImage(down.width, down.height); - fk::Ptr2D d_rgbaImageBig(NUM_ELEMS_X, NUM_ELEMS_Y); - - gpuErrchk(cudaMemcpy2DAsync(d_nv12Image.data.data, d_nv12Image.data.dims.pitch, - nv12Image.data, nv12Image.step, - NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2), cudaMemcpyHostToDevice, stream)); constexpr int CAMERAS = 4; constexpr int OUTPUTS = 1; for (int i = 0; i < CAMERAS; i++) { - fk::Read> read{ {d_nv12Image}}; - fk::Unary> cvtColor{}; - fk::Write> write{ d_rgbaImageBig.ptr() }; - fk::executeOperations>(fk_stream, read, cvtColor, write); - - fk::Read> read2{ {d_rgbaImageBig.ptr()}}; - fk::Unary> cvtColor2{}; - fk::Write> write2{ d_rgbaImageBig.ptr() }; - fk::executeOperations>(fk_stream, read2, cvtColor2, write2); + const auto read = ReadYUV::build(nv12Image); + const auto cvtColor = ConvertYUVToRGB::build(); + const auto addAlpha = AddOpaqueAlpha::build(); + const auto cast = SaturateCast::build(); + const auto write = PerThreadWrite::build(rgbaImageBig.ptr()); + executeOperations>(fk_stream, read, cvtColor, addAlpha, cast, write); + + const auto read2 = PerThreadRead::build(rgbaImageBig.ptr()); + const auto cvtColor2 = VectorReorder::build(); + const auto write2 = PerThreadWrite::build(rgbaImageBig.ptr()); + executeOperations>(fk_stream, read2, cvtColor2, write2); } for (int i = 0; i < OUTPUTS; i++) { - auto read3 = fk::Resize::build(d_rgbaImageBig.ptr(), down, 0., 0.); - fk::Unary> convertTo3{}; - fk::Write> write3{ d_rgbaImage.ptr() }; - fk::executeOperations>(fk_stream, read3, convertTo3, write3); + const auto read3 = Resize::build(rgbaImageBig.ptr(), down, 0., 0.); + const auto convertTo3 = SaturateCast::build(); + const auto write3 = PerThreadWrite::build(rgbaImage.ptr()); + executeOperations>(fk_stream, read3, convertTo3, write3); } - gpuErrchk(cudaMemcpy2DAsync(h_result.data, h_result.step, - d_rgbaImage.ptr().data, d_rgbaImage.dims().pitch, - down.width * sizeof(uchar4), down.height, cudaMemcpyDeviceToHost, stream)); - gpuErrchk(cudaStreamSynchronize(stream)); - - const auto readBackOp = fk::fuse(fk::Read>{d_nv12Image}, - fk::Unary>{}); - const fk::Size srcSize(NUM_ELEMS_X, NUM_ELEMS_Y); - const auto readOp = - fk::Resize::build(readBackOp, down); - auto convertOp = fk::Unary>{}; - auto colorConvert = fk::Unary>{}; - - fk::Write> writesTensor; - fk::Tensor myTensor(down.width, down.height, OUTPUTS); - writesTensor.params = myTensor; - - auto OpSeqTensor = fk::buildOperationSequence(readOp, convertOp, colorConvert, writesTensor); + fk_stream.sync(); - dim3 block = dim3(32,8); - dim3 grid((uint)ceil((float)down.width / (float)block.x), - (uint)ceil((float)down.height / (float)block.y), - (uint)OUTPUTS); + const auto readBackOp = ReadYUV::build(nv12Image).then( + ConvertYUVToRGB::build()); + const Size srcSize(NUM_ELEMS_X, NUM_ELEMS_Y); + const auto readOp = Resize::build(readBackOp, down); + const auto addAlpha = AddOpaqueAlpha::build(); + auto convertOp = SaturateCast::build(); + auto colorConvert = VectorReorder::build(); - const typename fk::DivergentBatchTransformDPP::DPPDetails details{}; + Tensor myTensor(down.width, down.height, OUTPUTS); + const auto writesTensor = TensorWrite::build(myTensor); - fk::launchDivergentBatchTransformDPP_Kernel<<>>(details, OpSeqTensor); - - gpuErrchk(cudaStreamSynchronize(stream)); + auto OpSeqTensor = buildOperationSequence(readOp, addAlpha, convertOp, colorConvert, writesTensor); - gpuErrchk(cudaFree(d_dataSource)); - - gpuErrchk(cudaStreamDestroy(stream)); + executeOperations>(fk_stream, + OpSeqTensor); + fk_stream.sync(); } -void testComputeWhatYouSee(char* buffer, const uint& NUM_ELEMS_X, const uint& NUM_ELEMS_Y) { - - cudaStream_t stream; - gpuErrchk(cudaStreamCreate(&stream)); - fk::Stream fk_stream{ stream }; - - constexpr fk::Size down(1920, 1080); - cv::Mat h_result(down.height, down.width, CV_8UC4); - cv::Mat nv12Image(cv::Size(NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2)), CV_8UC1, buffer); - - uchar* d_dataSource; - size_t sourcePitch; - gpuErrchk(cudaMallocPitch(&d_dataSource, &sourcePitch, NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2))); - fk::RawImage d_nv12Image{ {d_dataSource, { NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2), (uint)sourcePitch}}, NUM_ELEMS_X, NUM_ELEMS_Y }; - fk::Ptr2D d_rgbaImage(down.width, down.height); - fk::Ptr2D d_rgbaImageBig(NUM_ELEMS_X, NUM_ELEMS_Y); - - gpuErrchk(cudaMemcpy2DAsync(d_nv12Image.data.data, d_nv12Image.data.dims.pitch, - nv12Image.data, nv12Image.step, - NUM_ELEMS_X, NUM_ELEMS_Y + (NUM_ELEMS_Y / 2), cudaMemcpyHostToDevice, stream)); - - fk::Read> read{ d_nv12Image }; - fk::Unary> cvtColor{}; - fk::Write> write{ d_rgbaImageBig.ptr() }; - fk::executeOperations>(fk_stream, read, cvtColor, write); - - fk::Read> read2{ d_rgbaImageBig.ptr() }; - fk::Unary> cvtColor2{}; - fk::Write> write2{ d_rgbaImageBig.ptr() }; - fk::executeOperations>(fk_stream, read2, cvtColor2, write2); - - auto read3 = fk::Resize::build(d_rgbaImageBig.ptr(), down, 0., 0.); - fk::Unary> convertTo3{}; - fk::Write> write3{ d_rgbaImage.ptr() }; - fk::executeOperations>(fk_stream, read3, convertTo3, write3); - - gpuErrchk(cudaMemcpy2DAsync(h_result.data, h_result.step, - d_rgbaImage.ptr().data, d_rgbaImage.dims().pitch, - down.width * sizeof(uchar4), down.height, cudaMemcpyDeviceToHost, stream)); - gpuErrchk(cudaStreamSynchronize(stream)); - - const auto readOpInstance = fk::fuse(fk::Read>{d_nv12Image}, - fk::Unary>{}); - const auto readOp = fk::Resize::build(readOpInstance, down); - auto convertOp = fk::Unary>{}; - auto colorConvert = fk::Unary>{}; - auto writeOp = fk::Write>{ d_rgbaImage.ptr() }; - fk::executeOperations>(fk_stream, readOp, convertOp, colorConvert, writeOp); - gpuErrchk(cudaMemcpy2DAsync(h_result.data, h_result.step, - d_rgbaImage.ptr().data, d_rgbaImage.dims().pitch, - down.width * sizeof(uchar4), down.height, cudaMemcpyDeviceToHost, stream)); - - gpuErrchk(cudaStreamSynchronize(stream)); - - gpuErrchk(cudaFree(d_dataSource)); - - gpuErrchk(cudaStreamDestroy(stream)); +void testComputeWhatYouSee(char *buffer, const uint &NUM_ELEMS_X, const uint &NUM_ELEMS_Y) { + using namespace fk; + Stream fk_stream; + + constexpr Size down(1920, 1080); + + Image nv12Image(NUM_ELEMS_X, NUM_ELEMS_Y); + memcpy(nv12Image.getData().ptrPinned().data, buffer, NUM_ELEMS_X * (NUM_ELEMS_Y + (NUM_ELEMS_Y / 2))); + Ptr2D rgbaImage(down.width, down.height); + Ptr2D rgbaImageBig(NUM_ELEMS_X, NUM_ELEMS_Y); + nv12Image.upload(fk_stream); + + const auto read = ReadYUV::build(nv12Image); + const auto cvtColor = ConvertYUVToRGB::build(); + const auto addAlpha = AddOpaqueAlpha::build(); + const auto cast = SaturateCast::build(); + const auto write = PerThreadWrite::build(rgbaImageBig.ptr()); + executeOperations>(fk_stream, read, cvtColor, addAlpha, cast, write); + + const auto read2 = PerThreadRead::build(rgbaImageBig.ptr()); + const auto cvtColor2 = VectorReorder::build(); + const auto write2 = PerThreadWrite::build(rgbaImageBig.ptr()); + executeOperations>(fk_stream, read2, cvtColor2, write2); + + const auto read3 = Resize::build(rgbaImageBig.ptr(), down, 0., 0.); + const auto convertTo3 = SaturateCast::build(); + const auto write3 = PerThreadWrite::build(rgbaImage.ptr()); + executeOperations>(fk_stream, read3, convertTo3, write3); + + rgbaImageBig.download(fk_stream); + rgbaImage.download(fk_stream); + fk_stream.sync(); + + const auto readOpInstance = ReadYUV::build(nv12Image).then( + ConvertYUVToRGB::build()); + const auto readOp = Resize::build(readOpInstance, down); + const auto convertOp = SaturateCast::build(); + const auto colorConvert = VectorReorder::build(); + const auto writeOp = PerThreadWrite::build(rgbaImage.ptr()); + executeOperations>(fk_stream, readOp, addAlpha, convertOp, colorConvert, writeOp); + rgbaImage.download(fk_stream); } int launch() { int returnValue = 0; - cv::cuda::Stream cv_stream; - - cv::Mat::setDefaultAllocator(cv::cuda::HostMem::getAllocator(cv::cuda::HostMem::AllocType::PAGE_LOCKED)); - - const std::string filePath{ "" }; + const std::string filePath{""}; std::ifstream file(filePath, std::ios::binary | std::ios::ate); std::streamsize size = file.tellg(); file.seekg(0, std::ios::beg);