Merge pull request #179 from CVCUDA/release-patch-v0.10.0

Merge patched release code for v0.10.0-beta into main

.gitattributes

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2022-2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2022-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: Apache-2.0
 #
 # Licensed under the Apache License, Version 2.0 (the "License");

.github/workflows/codeql.yml

@@ -1,3 +1,18 @@
+# SPDX-FileCopyrightText: Copyright (c) 2022-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
 name: "CodeQL"
 
 on:

.gitignore

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2022-2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2022-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: Apache-2.0
 #
 # Licensed under the Apache License, Version 2.0 (the "License");

CMakeLists.txt

@@ -23,7 +23,7 @@ endif()
 
 project(cvcuda
         LANGUAGES C CXX
-        VERSION 0.9.0
+        VERSION 0.10.0
         DESCRIPTION "CUDA-accelerated Computer Vision algorithms"
 )
 

LICENSE.md

@@ -87,18 +87,18 @@ END OF TERMS AND CONDITIONS
 
 # APPENDIX: How to apply the Apache License to your work
 
-To apply the Apache License to your work, attach the following boilerplate notice, with the fields enclosed by brackets "[]" replaced with your own identifying information. (Don't include the brackets!) The text should be enclosed in the appropriate comment syntax for the file format. We also recommend that a file or class name and description of purpose be included on the same "printed page" as the copyright notice for easier identification within third-party archives.
+To apply the Apache License to your work, attach the following boilerplate notice, with the fields enclosed by brackets "[ ]" replaced with your own identifying information. (Don't include the brackets!) The text should be enclosed in the appropriate comment syntax for the file format. We also recommend that a file or class name and description of purpose be included on the same "printed page" as the copyright notice for easier identification within third-party archives.
 
-   Copyright (c) 2022-2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+    Copyright (c) [yyyy] [name of copyright owner]
 
-   Licensed under the Apache License, Version 2.0 (the "License");
-   you may not use this file except in compliance with the License.
-   You may obtain a copy of the License at
+    Licensed under the Apache License, Version 2.0 (the "License");
+    you may not use this file except in compliance with the License.
+    You may obtain a copy of the License at
 
-     http://www.apache.org/licenses/LICENSE-2.0
+      http://www.apache.org/licenses/LICENSE-2.0
 
-   Unless required by applicable law or agreed to in writing, software
-   distributed under the License is distributed on an "AS IS" BASIS,
-   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-   See the License for the specific language governing permissions and
-   limitations under the License.
+    Unless required by applicable law or agreed to in writing, software
+    distributed under the License is distributed on an "AS IS" BASIS,
+    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+    See the License for the specific language governing permissions and
+    limitations under the License.

README.md

@@ -18,7 +18,7 @@
 
 [![License](https://img.shields.io/badge/License-Apache_2.0-yellogreen.svg)](https://opensource.org/licenses/Apache-2.0)
 
-![Version](https://img.shields.io/badge/Version-v0.9.0--beta-blue)
+![Version](https://img.shields.io/badge/Version-v0.10.0--beta-blue)
 
 ![Platform](https://img.shields.io/badge/Platform-linux--64_%7C_win--64_wsl2%7C_aarch64-gray)
 
@@ -61,7 +61,9 @@ To get a local copy up and running follow these steps.
 - Only one CUDA version (CUDA 11.x or CUDA 12.x) of CV-CUDA packages (Debian packages, tarballs, Python Wheels) can be installed at a time. Please uninstall all packages from a given CUDA version before installing packages from a different version.
 - Documentation built on Ubuntu 20.04 needs an up-to-date version of sphinx (`pip install --upgrade sphinx`) as well as explicitly parsing the system's default python version ` ./ci/build_docs path/to/build -DPYTHON_VERSIONS="<py_ver>"`.
 - Python bindings installed via Debian packages and Python tests fail with Numpy 2.0. We recommend using an older version of Numpy (e.g. 1.26) until we have implemented a fix.
-- The Resize and RandomResizedCrop operators incorrectly interpolate pixel values near the boundary of an image or tensor when using linear and cubic interpolation. This will be fixed in an upcoming release.
+- The Resize and RandomResizedCrop operators incorrectly interpolate pixel values near the boundary of an image or tensor when using cubic interpolation. This will be fixed in an upcoming release.
+- Cache/resource management introduced in v0.10 add micro-second-level overhead to Python operator calls. Based on the performance analysis of our Python samples, we expect the production- and pipeline-level impact to be negligible. CUDA kernel and C++ call performance is not affected. We aim to investigate and reduce this overhead further in a future release.​
+- Sporadic Pybind11-deallocation crashes have been reported in long-lasting multi-threaded Python pipelines with externally allocated memory (eg wrapped Pytorch buffers). We are evaluating an upgrade of Pybind11 (currently using 2.10) as a potential fix in an upcoming release.
 
 ### Installation
 

bench/BenchAdaptiveThreshold.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchAverageBlur.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchBilateralFilter.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchBrightnessContrast.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchColorTwist.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchComposite.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchCopyMakeBorder.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchCvtColor.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
@@ -21,63 +21,158 @@
 
 #include <nvbench/nvbench.cuh>
 
-template<typename T>
-inline void CvtColor(nvbench::state &state, nvbench::type_list<T>)
+#include <map>
+#include <stdexcept>
+#include <tuple>
+
+inline static std::tuple<NVCVColorConversionCode, NVCVImageFormat, NVCVImageFormat> StringToFormats(
+    const std::string &str)
+{
+    // clang-format off
+    static const std::map<std::string, std::tuple<NVCVColorConversionCode, NVCVImageFormat, NVCVImageFormat>> codeMap{
+        {     "RGB2BGR", {NVCV_COLOR_RGB2BGR,      NVCV_IMAGE_FORMAT_RGB8,  NVCV_IMAGE_FORMAT_BGR8 }},
+        {    "RGB2RGBA", {NVCV_COLOR_RGB2RGBA,     NVCV_IMAGE_FORMAT_RGB8,  NVCV_IMAGE_FORMAT_RGBA8}},
+        {    "RGBA2RGB", {NVCV_COLOR_RGBA2RGB,     NVCV_IMAGE_FORMAT_RGBA8, NVCV_IMAGE_FORMAT_RGB8 }},
+        {    "RGB2GRAY", {NVCV_COLOR_RGB2GRAY,     NVCV_IMAGE_FORMAT_RGB8,  NVCV_IMAGE_FORMAT_Y8   }},
+        {    "GRAY2RGB", {NVCV_COLOR_GRAY2RGB,     NVCV_IMAGE_FORMAT_Y8,    NVCV_IMAGE_FORMAT_RGB8 }},
+        {     "RGB2HSV", {NVCV_COLOR_RGB2HSV,      NVCV_IMAGE_FORMAT_RGB8,  NVCV_IMAGE_FORMAT_HSV8 }},
+        {     "HSV2RGB", {NVCV_COLOR_HSV2RGB,      NVCV_IMAGE_FORMAT_HSV8,  NVCV_IMAGE_FORMAT_RGB8 }},
+        {     "RGB2YUV", {NVCV_COLOR_RGB2YUV,      NVCV_IMAGE_FORMAT_RGB8,  NVCV_IMAGE_FORMAT_YUV8 }},
+        {     "YUV2RGB", {NVCV_COLOR_YUV2RGB,      NVCV_IMAGE_FORMAT_YUV8,  NVCV_IMAGE_FORMAT_RGB8 }},
+        {"RGB2YUV_NV12", {NVCV_COLOR_RGB2YUV_NV12, NVCV_IMAGE_FORMAT_RGB8,  NVCV_IMAGE_FORMAT_NV12 }},
+        {"YUV2RGB_NV12", {NVCV_COLOR_YUV2RGB_NV12, NVCV_IMAGE_FORMAT_NV12,  NVCV_IMAGE_FORMAT_RGB8 }},
+    };
+    // clang-format on
+
+    if (auto it = codeMap.find(str); it != codeMap.end())
+    {
+        return it->second;
+    }
+    else
+    {
+        throw std::invalid_argument("Unrecognized color code");
+    }
+}
+
+template<typename BT>
+inline float BytesPerPixel(NVCVImageFormat imgFormat)
+{
+    switch (imgFormat)
+    {
+#define CVCUDA_BYTES_PER_PIXEL_CASE(FORMAT, BYTES) \
+    case FORMAT:                                   \
+        return BYTES * sizeof(BT)
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_RGB8, 3);
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_BGR8, 3);
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_HSV8, 3);
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_RGBA8, 4);
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_YUV8, 3);
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_NV12, 1.5f);
+        CVCUDA_BYTES_PER_PIXEL_CASE(NVCV_IMAGE_FORMAT_Y8, 1);
+#undef CVCUDA_BYTES_PER_PIXEL_CASE
+    default:
+        throw std::invalid_argument("Unrecognized format");
+    }
+}
+
+// Adapted from src/util/TensorDataUtils.hpp
+inline static nvcv::Tensor CreateTensor(int numImages, int imgWidth, int imgHeight, const nvcv::ImageFormat &imgFormat)
+{
+    if (imgFormat == NVCV_IMAGE_FORMAT_NV12 || imgFormat == NVCV_IMAGE_FORMAT_NV12_ER
+        || imgFormat == NVCV_IMAGE_FORMAT_NV21 || imgFormat == NVCV_IMAGE_FORMAT_NV21_ER)
+    {
+        int height420 = (imgHeight * 3) / 2;
+        if (height420 % 3 != 0 || imgWidth % 2 != 0)
+        {
+            throw std::invalid_argument("Invalid height");
+        }
+
+        return nvcv::Tensor(numImages, {imgWidth, height420}, nvcv::ImageFormat(NVCV_IMAGE_FORMAT_Y8));
+    }
+    else
+    {
+        return nvcv::Tensor(numImages, {imgWidth, imgHeight}, imgFormat);
+    }
+}
+
+template<typename BT>
+inline void CvtColor(nvbench::state &state, nvbench::type_list<BT>)
 try
 {
     long3 shape    = benchutils::GetShape<3>(state.get_string("shape"));
     long  varShape = state.get_int64("varShape");
+    std::tuple<NVCVColorConversionCode, NVCVImageFormat, NVCVImageFormat> formats
+        = StringToFormats(state.get_string("code"));
 
-    using BT = typename nvcv::cuda::BaseType<T>;
-
-    int ch = nvcv::cuda::NumElements<T>;
+    NVCVColorConversionCode code = std::get<0>(formats);
+    nvcv::ImageFormat       inFormat{std::get<1>(formats)};
+    nvcv::ImageFormat       outFormat{std::get<2>(formats)};
 
-    NVCVColorConversionCode code = ch == 3 ? NVCV_COLOR_BGR2RGB : NVCV_COLOR_BGRA2RGBA;
-
-    state.add_global_memory_reads(shape.x * shape.y * shape.z * sizeof(T));
-    state.add_global_memory_writes(shape.x * shape.y * shape.z * sizeof(T));
+    state.add_global_memory_reads(shape.x * shape.y * shape.z * BytesPerPixel<BT>(inFormat));
+    state.add_global_memory_writes(shape.x * shape.y * shape.z * BytesPerPixel<BT>(outFormat));
 
     cvcuda::CvtColor op;
 
-    // clang-format off
-
     if (varShape < 0) // negative var shape means use Tensor
     {
-        nvcv::Tensor src({{shape.x, shape.y, shape.z, ch}, "NHWC"}, benchutils::GetDataType<BT>());
-        nvcv::Tensor dst({{shape.x, shape.y, shape.z, ch}, "NHWC"}, benchutils::GetDataType<BT>());
+        nvcv::Tensor src = CreateTensor(shape.x, shape.z, shape.y, inFormat);
+        nvcv::Tensor dst = CreateTensor(shape.x, shape.z, shape.y, outFormat);
 
         benchutils::FillTensor<BT>(src, benchutils::RandomValues<BT>());
 
-        state.exec(nvbench::exec_tag::sync, [&op, &src, &dst, &code](nvbench::launch &launch)
-        {
-            op(launch.get_stream(), src, dst, code);
-        });
+        state.exec(nvbench::exec_tag::sync,
+                   [&op, &src, &dst, &code](nvbench::launch &launch) { op(launch.get_stream(), src, dst, code); });
     }
     else // zero and positive var shape means use ImageBatchVarShape
     {
-        nvcv::ImageBatchVarShape src(shape.x);
-        nvcv::ImageBatchVarShape dst(shape.x);
+        if (inFormat.chromaSubsampling() != nvcv::ChromaSubsampling::CSS_444
+            || outFormat.chromaSubsampling() != nvcv::ChromaSubsampling::CSS_444)
+        {
+            state.skip("Skipping formats that have subsampled planes for the varshape benchmark");
+        }
+
+        std::vector<nvcv::Image>          imgSrc;
+        std::vector<nvcv::Image>          imgDst;
+        nvcv::ImageBatchVarShape          src(shape.x);
+        nvcv::ImageBatchVarShape          dst(shape.x);
+        std::vector<std::vector<uint8_t>> srcVec(shape.x);
 
-        benchutils::FillImageBatch<T>(src, long2{shape.z, shape.y}, long2{varShape, varShape},
-                                      benchutils::RandomValues<T>());
-        dst.pushBack(src.begin(), src.end());
+        auto randomValuesU8 = benchutils::RandomValues<uint8_t>();
 
-        state.exec(nvbench::exec_tag::sync, [&op, &src, &dst, &code](nvbench::launch &launch)
+        for (int i = 0; i < shape.x; i++)
         {
-            op(launch.get_stream(), src, dst, code);
-        });
+            imgSrc.emplace_back(nvcv::Size2D{(int)shape.z, (int)shape.y}, inFormat);
+            imgDst.emplace_back(nvcv::Size2D{(int)shape.z, (int)shape.y}, outFormat);
+
+            int srcRowStride = imgSrc[i].size().w * inFormat.planePixelStrideBytes(0);
+            int srcBufSize   = imgSrc[i].size().h * srcRowStride;
+            srcVec[i].resize(srcBufSize);
+            for (int idx = 0; idx < srcBufSize; idx++)
+            {
+                srcVec[i][idx] = randomValuesU8();
+            }
+
+            auto imgData = imgSrc[i].exportData<nvcv::ImageDataStridedCuda>();
+            CUDA_CHECK_ERROR(cudaMemcpy2D(imgData->plane(0).basePtr, imgData->plane(0).rowStride, srcVec[i].data(),
+                                          srcRowStride, srcRowStride, imgSrc[i].size().h, cudaMemcpyHostToDevice));
+        }
+        src.pushBack(imgSrc.begin(), imgSrc.end());
+        dst.pushBack(imgDst.begin(), imgDst.end());
+
+        state.exec(nvbench::exec_tag::sync,
+                   [&op, &src, &dst, &code](nvbench::launch &launch) { op(launch.get_stream(), src, dst, code); });
     }
 }
 catch (const std::exception &err)
 {
     state.skip(err.what());
 }
 
-// clang-format on
-
-using CvtColorTypes = nvbench::type_list<uchar3, uchar4>;
+using BaseTypes = nvbench::type_list<uint8_t>;
 
-NVBENCH_BENCH_TYPES(CvtColor, NVBENCH_TYPE_AXES(CvtColorTypes))
-    .set_type_axes_names({"InOutDataType"})
-    .add_string_axis("shape", {"1x1080x1920"})
+NVBENCH_BENCH_TYPES(CvtColor, NVBENCH_TYPE_AXES(BaseTypes))
+    .set_type_axes_names({"BaseType"})
+    .add_string_axis("shape", {"1x1080x1920", "64x720x1280"})
+    .add_string_axis("code", {"RGB2BGR", "RGB2RGBA", "RGBA2RGB", "RGB2GRAY", "GRAY2RGB", "RGB2HSV", "HSV2RGB",
+                              "RGB2YUV", "YUV2RGB", "RGB2YUV_NV12", "YUV2RGB_NV12"})
     .add_int64_axis("varShape", {-1, 0});

bench/BenchErase.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchFlip.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchGaussian.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchGaussianNoise.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchHistogramEq.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchInpaint.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchJointBilateralFilter.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchLabel.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

bench/BenchLaplacian.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  *
  * Licensed under the Apache License, Version 2.0 (the "License");