[FEAT] Update NCU Profile Script and RelWithDebInfo Build Mode

Author: jamesnulliu

.clangd

@@ -46,7 +46,7 @@ Diagnostics:
       readability-identifier-naming.PrivateMemberCase: aNy_CasE
       readability-identifier-naming.PublicMemberCase: aNy_CasE
       readability-identifier-naming.NamespaceCase: lower_case
-      readability-identifier-naming.EnumCase: camelBack
+      readability-identifier-naming.EnumCase: CamelCase
       readability-identifier-naming.ClassCase: CamelCase
       readability-identifier-naming.StructCase: CamelCase
       readability-identifier-naming.FunctionCase: camelBack

.gitignore

@@ -160,4 +160,6 @@ cython_debug/
 
 /pmpp/_torch_ops
 vcpkg_installed
-/tmp
+/tmp
+/output/*
+!.gitkeep

README.md

@@ -5,6 +5,8 @@
 
 ## 1. Environment
 
+### 1.1. Method 1: Use Docker Image
+
 The simplest way is to use my docker image [jamesnulliu/deeplearning:latest](https://hub.docker.com/r/jamesnulliu/deeplearning) which contains all the softwares you need to build the project:
 
 ```bash
@@ -13,6 +15,8 @@ docker pull jamesnulliu/deeplearning:latest
 
 > Check my blog: [Docker Container with Nvidia GPU Support](https://jamesnulliu.github.io/blogs/docker-container-with-nvidia-gpu-support) if you need any help.
 
+### 1.2. Method 2: Setup Environment Manually
+
 Or if you are planing to setup your own environment, here are some tips:
 
 You should install all the softwares with corresponding versions listed bellow:
@@ -27,7 +31,7 @@ You should install all the softwares with corresponding versions listed bellow:
 
 **🎯Miniconda**
 
-Managing python environments with miniconda is always a good choice. Check [the official website](https://docs.anaconda.com/miniconda/install/#quick-command-line-install) for an installation guide.
+Managing python environments with miniconda is always a good choice. Check [the official website](https://www.anaconda.com/docs/getting-started/miniconda/install#quickstart-install-instructions) for an installation guide.
 
 After installation, if you do not intend to install all the packages in `base` environment, create a new conda environment named `PMPP` (or whatever you like) and activate it:
 

configs/lib-tests.yml

@@ -1,4 +1,8 @@
 OpTest:
+  VecAdd:
+    - nElems: 64
+    - nElems: 1024
+    - nElems: 2048
   Conv2D:
     - inputHeight: 32
       inputWidth: 32
@@ -13,4 +17,8 @@ OpTest:
     # # [NOTE] The following calculation results would be wrong?
     # - nInputs: 64
     # - nInputs: 128
-    # - nInputs: 512
+    # - nInputs: 512
+  PrefixSum:
+    - nInputs: 32
+    - nInputs: 64
+    - nInputs: 128

csrc/cmake/compilers/cuda-compiler-configs.cmake

@@ -7,7 +7,7 @@
 #         - `CMAKE_CUDA_STANDARD`: CUDA Standard. Default: 20.
 # ==================================================================================================
 
-include(${PROJECT_SOURCE_DIR}/cmake/utils/logging.cmake)
+include(${CMAKE_CURRENT_LIST_DIR}/../utils/logging.cmake)
 
 enable_language(CUDA)
 
@@ -24,4 +24,5 @@ log_info("CMAKE_CUDA_STANDARD: ${CMAKE_CUDA_STANDARD}")
 
 string(APPEND CMAKE_CUDA_FLAGS " --expt-relaxed-constexpr")
 string(APPEND CMAKE_CUDA_FLAGS_RELEASE " -O3")
+string(APPEND CMAKE_CUDA_FLAGS_RELWITHDEBINFO " -O3 -lineinfo")
 string(APPEND CMAKE_CUDA_FLAGS_DEBUG " -lineinfo")

csrc/cmake/compilers/cxx-compiler-configs.cmake

@@ -11,14 +11,15 @@
 #         - `STACK_SIZE`: Stack size for the executable. Default: 1048576 (1MB).
 # ==================================================================================================
 
-include(${PROJECT_SOURCE_DIR}/cmake/utils/common.cmake)
+include(${CMAKE_CURRENT_LIST_DIR}/../utils/common.cmake)
 
 enable_language(CXX)
 
+set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
+set(CMAKE_CXX_STANDARD_REQUIRED   ON)
+
 set_default_values(
-    CMAKE_EXPORT_COMPILE_COMMANDS ON
     CMAKE_CXX_STANDARD            20
-    CMAKE_CXX_STANDARD_REQUIRED   ON
     CMAKE_CXX_SCAN_FOR_MODULES    OFF
     STACK_SIZE                    1048576
 )
@@ -31,13 +32,15 @@ log_info("STACK_SIZE: ${STACK_SIZE}")
 if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
     string(APPEND CMAKE_CXX_FLAGS " /permissive- /Zc:forScope /openmp /Zc:__cplusplus")
     string(APPEND CMAKE_CXX_FLAGS_RELEASE " /O2")
+    string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " /O2 /Zi")
     string(APPEND CMAKE_CXX_FLAGS_DEBUG " /Zi")
     # Set stack size
     string(APPEND CMAKE_EXE_LINKER_FLAGS " /STACK:${STACK_SIZE}")
 # Compiler flags for Clang
 elseif(CMAKE_CXX_COMPILER_ID STREQUAL "Clang")
     string(APPEND CMAKE_CXX_FLAGS " -fopenmp -Wall -Wextra -Werror")
     string(APPEND CMAKE_CXX_FLAGS_RELEASE " -O3")
+    string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " -O3 -g")
     string(APPEND CMAKE_CXX_FLAGS_DEBUG " -g")
     # Set stack size
     if (WIN32)
@@ -49,6 +52,7 @@ elseif(CMAKE_CXX_COMPILER_ID STREQUAL "Clang")
 elseif(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
     string(APPEND CMAKE_CXX_FLAGS " -fopenmp -Wall -Wextra -Werror")
     string(APPEND CMAKE_CXX_FLAGS_RELEASE " -O3")
+    string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " -O3 -g")
     string(APPEND CMAKE_CXX_FLAGS_DEBUG " -g")
     # Set stack size
     if (WIN32)

csrc/lib/ops/prefixSum/op.cuh

@@ -0,0 +1,73 @@
+#pragma once
+
+#include "pmpp/pch.hpp"
+
+#include "pmpp/utils/math.hpp"
+
+namespace pmpp::ops::cuda
+{
+
+template <typename ScalarT>
+__global__ void koggeStonePrefixSumKernel(const ScalarT* input,
+                                          ScalarT* output, uint32_t n)
+{
+    extern __shared__ ScalarT shmem[];
+
+    uint32_t btid = threadIdx.x;                            // Block Thread ID
+    uint32_t gtid = blockIdx.x * blockDim.x + threadIdx.x;  // Global Thread ID
+
+    if (gtid < n) {
+        shmem[btid] = input[gtid];
+    } else {
+        shmem[btid] = 0;
+    }
+
+    __syncthreads();
+
+    for (uint32_t stride = 1; stride < blockDim.x; stride *= 2) {
+        ScalarT tmp = 0;
+        if (btid >= stride) {
+            tmp = shmem[btid] + shmem[btid - stride];
+        }
+        __syncthreads();
+        if (btid >= stride) {
+            shmem[btid] = tmp;
+        }
+        __syncthreads();
+    }
+
+    if (gtid < n) {
+        output[gtid] = shmem[btid];
+    }
+}
+
+template <typename ScalarT>
+void launchPrefixSum(const ScalarT* d_input, ScalarT* d_output, uint32_t n)
+{
+    constexpr uint32_t blockSize = 256;
+    uint32_t gridSize = ceilDiv(n, blockSize);
+    koggeStonePrefixSumKernel<<<gridSize, blockSize,
+                                blockSize * sizeof(ScalarT)>>>(d_input,
+                                                               d_output, n);
+}
+
+namespace torch_impl
+{
+inline auto prefixSum(const torch::Tensor& A) -> torch::Tensor
+{
+    torch::Tensor B = torch::empty_like(A);
+    switch (A.scalar_type()) {
+    case torch::kFloat32: {
+        pmpp::ops::cuda::launchPrefixSum(A.data_ptr<fp32_t>(),
+                                         B.data_ptr<fp32_t>(), A.size(0));
+        break;
+    }
+    default:
+        AT_ERROR("Unsupported dtype: ", A.dtype());
+    }
+
+    return B;
+}
+}  // namespace torch_impl
+
+}  // namespace pmpp::ops::cuda

csrc/lib/ops/prefixSum/op.hpp

@@ -0,0 +1,39 @@
+#pragma once
+
+#include "pmpp/pch.hpp"
+
+namespace pmpp::ops::cpu
+{
+template <typename ScalarT>
+void launchPrefixSum(ScalarT* input, ScalarT* output, size_t n)
+{
+    output[0] = input[0];
+    for (size_t i = 1; i < n; ++i) {
+        output[i] = output[i - 1] + input[i];
+    }
+}
+
+namespace torch_impl
+{
+inline auto prefixSum(const torch::Tensor& A) -> torch::Tensor
+{
+    torch::Tensor B = torch::empty_like(A);
+    switch (A.scalar_type()) {
+    case torch::kFloat32: {
+        pmpp::ops::cpu::launchPrefixSum(A.data_ptr<fp32_t>(),
+                                        B.data_ptr<fp32_t>(), A.size(0));
+        break;
+    }
+    case torch::kInt32: {
+        pmpp::ops::cpu::launchPrefixSum(A.data_ptr<int32_t>(),
+                                        B.data_ptr<int32_t>(), A.size(0));
+        break;
+    }
+    default:
+        AT_ERROR("Unsupported dtype: ", A.dtype());
+    }
+
+    return B;
+} 
+}  // namespace torch_impl
+}  // namespace pmpp::ops::cpu

csrc/lib/ops/torch_bind.cu

@@ -13,6 +13,7 @@ TORCH_LIBRARY(pmpp, m)
     m.def("conv2d(Tensor input, Tensor kernel) -> Tensor");
     m.def("alphabet_histogram(Tensor input, int divider) -> Tensor");
     m.def("mul_reduction(Tensor input) -> Tensor");
+    m.def("prefix_sum(Tensor input) -> Tensor");
 }
 
 // Register the implementations.
@@ -27,6 +28,7 @@ TORCH_LIBRARY_IMPL(pmpp, CPU, m)
     m.impl("alphabet_histogram",
            &pmpp::ops::cpu::torch_impl::alphabetHistogram);
     m.impl("mul_reduction", &pmpp::ops::cpu::torch_impl::mulReduction);
+    m.impl("prefix_sum", &pmpp::ops::cpu::torch_impl::prefixSum);
 }
 
 TORCH_LIBRARY_IMPL(pmpp, CUDA, m)
@@ -38,4 +40,5 @@ TORCH_LIBRARY_IMPL(pmpp, CUDA, m)
     m.impl("alphabet_histogram",
            &pmpp::ops::cuda::torch_impl::alphabetHistogram);
     m.impl("mul_reduction", &pmpp::ops::cuda::torch_impl::mulReduction);
+    m.impl("prefix_sum", &pmpp::ops::cuda::torch_impl::prefixSum);
 }

csrc/lib/ops/torch_impl.hpp

@@ -8,6 +8,8 @@
 #include "./cvtRGBtoGray/op.hpp"
 #include "./matmul/op.cuh"
 #include "./matmul/op.hpp"
+#include "./prefixSum/op.cuh"
+#include "./prefixSum/op.hpp"
 #include "./reduction/op.cuh"
 #include "./reduction/op.hpp"
 #include "./vecAdd/op.cuh"