Update NCU profile

Author: jamesnulliu

.gitignore

@@ -161,5 +161,5 @@ cython_debug/
 /pmpp/_torch_ops
 vcpkg_installed
 /tmp
-/output/*
+/outputs/*
 !.gitkeep

configs/lib-tests.yml

@@ -1,8 +1,10 @@
 OpTest:
   VecAdd:
-    - nElems: 64
+    - nElems: 32
     - nElems: 1024
     - nElems: 2048
+    - nElems: 1048576  # 1024*1024
+    - nElems: 8388608  # 1024*1024
   Conv2D:
     - inputHeight: 32
       inputWidth: 32
@@ -14,10 +16,6 @@ OpTest:
       divider: 4
   MulReduction:
     - nInputs: 32
-    # # [NOTE] The following calculation results would be wrong?
-    # - nInputs: 64
-    # - nInputs: 128
-    # - nInputs: 512
   PrefixSum:
     - nInputs: 32
     - nInputs: 64

csrc/CMakeLists.txt

@@ -9,9 +9,12 @@ include(${PROJECT_SOURCE_DIR}/cmake/utils/common.cmake)
 
 log_info("CMake Tookchain File Path: ${CMAKE_TOOLCHAIN_FILE}")
 
+set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
+
 set_default_values(
     PROJECT_NAMESPACE "pmpp::"
     STDOUT_IS_TERMINAL OFF
+    CMAKE_EXECUTABLE_SUFFIX ""
 )
 
 # Utility functions for installing libraries:

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

@@ -15,7 +15,6 @@ 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(

csrc/lib/ops/vecAdd/op.cuh

@@ -12,15 +12,15 @@ __global__ void vecAddKernel(const fp32_t* a, const fp32_t* b, fp32_t* c,
 
     int gtid = threadIdx.x + blockDim.x * blockIdx.x;
     if (gtid < n) {
-        // [GM] 2 load, 1 store
+        // [GM] 2 load, 1 store, 3 inst
         c[gtid] = a[gtid] + b[gtid];
     }
 }
 
 void launchVecAdd(const fp32_t* d_A, const fp32_t* d_B, fp32_t* d_C, size_t n)
 {
-    dim3 blockSize = 256;
-    dim3 gridSize = ceilDiv(n, 256);
+    dim3 blockSize = {std::min<uint32_t>(n, 1024), 1, 1};
+    dim3 gridSize = {ceilDiv<uint32_t>(n, blockSize.x), 1, 1};
 
     vecAddKernel<<<gridSize, blockSize>>>(d_A, d_B, d_C, int32_t(n));
 

output/.gitkeep outputs/.gitkeepoutput/.gitkeep renamed to outputs/.gitkeep

@@ -1,4 +1,4 @@
-# Env Variables: CC, CXX, NVCC_CCBIN
+# Env Variables: CC, CXX, NVCC_CCBIN, VCPKG_HOME
 
 set -e  # Exit on error
 
@@ -50,8 +50,8 @@ done
 
 cmake -S $SOURCE_DIR -B $BUILD_DIR -G Ninja \
     -DCMAKE_TOOLCHAIN_FILE="$VCPKG_HOME/scripts/buildsystems/vcpkg.cmake" \
-    -DVCPKG_TARGET_TRIPLET="x64-linux" \
     -DVCPKG_OVERLAY_TRIPLETS="csrc/cmake/vcpkg-triplets" \
+    -DVCPKG_TARGET_TRIPLET="x64-linux" \
     -DSTDOUT_IS_TERMINAL=$STDOUT_IS_TERMINAL \
     -DCMAKE_BUILD_TYPE=$BUILD_TYPE \
     -DCMAKE_CXX_STANDARD=$CXX_STANDARD \

scripts/build.sh

@@ -1,6 +1,6 @@
 TEST_FILE="./build/test/pmpp_test"
 GTEST_FILTER="OpTest.VecAdd"
-OUTPUT_FILE="output/nsight_profile.ncu-rep"
+OUTPUT_FILE="outputs/nsight_profile.ncu-rep"
 
 while [[ $# -gt 0 ]]; do
     case $1 in
@@ -17,5 +17,6 @@ while [[ $# -gt 0 ]]; do
 done
 
 
-ncu --export $OUTPUT_FILE --force-overwrite $TEST_FILE \
-    --gtest_filter=$GTEST_FILTER
+ncu --export $OUTPUT_FILE --force-overwrite \
+    --set "full" \
+    $TEST_FILE --gtest_filter=$GTEST_FILTER