add cuda codes

Author: archibate

cudaguide/1_moderncuda/README.md

@@ -320,13 +320,23 @@ CUDA 中的函数分为三大类：
 
 但是，CPU 上的每个进程只会调用一次 `main` 入口点，而 GPU 上的 `__global__` 会被调用 n 次（n 的大小在 host 函数中指定），所有启动的 n 个 `__global__` 函数互相之间是并行执行的，每个线程的入口点都是 `__global__`，因此一个“网格”含有多个“线程”。
 
+通过在函数的定义前加上关键字前缀，可以改变函数的类型。
+
+其中 `__host__` 可以省略，不带任何前缀时默认就算 `__host__`。因此 C++ 的代码可以无缝植入到 CUDA 代码中，C++ 的函数默认就会被视为 `__host__`，放在 CPU 侧编译。
+
+以下是一个案例：
+
 ```cuda
 #include <cuda_runtime.h>
 
-/*__host__*/ void host_func() {
+/*__host__*/ void host_func() {  // 仅为此函数编译 CPU 版本
+    printf("hello from host!\n");
+    std::cout << "cout from host!" << std::endl;
 }
 
-__device__ void device_func() {
+__device__ void device_func() {  // 仅为此函数编译 GPU 版本
+    printf("hello from device!\n"); // OK！CUDA 为 printf 做了个特殊的 device 版重载，可以在 GPU 中调用的一个特供版本
+    std::cout << "cout from host!" << std::endl; // 编译错误，cout 是 C++ 的华丽胡哨垃圾，CUDA 官方没有适配，无法在 GPU 调用
 }
 
 __host__ __device__ void host_device_func() {
@@ -374,7 +384,9 @@ int main() {
     // kernel<<<blockDim, gridDim>>>(...)
     kernel<<<3, 4>>>();
 
-    // 强制同步：等待此前启动过的所有内核执行完成
+    // CUDA 内核的启动都是异步的，类似于 std::thread(kernel).detach() 的效果
+    // 因此需要强制同步，等待此前启动过的所有内核执行完成，才退出程序
+    // 否则可能内核还没来得及启动，CPU 程序就被你退出了，导致收不到 printf 的信息
     cudaDeviceSynchronize();
 
     return 0;
@@ -398,7 +410,7 @@ int main() {
 | 启动内核 `kernel<<<3, 4>>>()` | `std::async` 或 `std::thread` |
 | `cudaDeviceSynchronize` | `future.wait()` 或 `thread.join()` |
 
-### 小彭老师为你准备的 CUDA 框架
+### 小彭老师赋能 CUDA 现代化
 
 ```cuda
 #include <cuda_runtime.h>
@@ -433,18 +445,14 @@ int main() {
     cfg.numAttrs = 0;
     CHECK_CUDA(cudaLaunchKernelEx(&cfg, kernel, x));
 
-    const char *name;
-    CHECK_CUDA(cudaFuncGetName(&name, kernel));
-    printf("内核名字：%s\n", name);
-
     // 1. 强制同步：等待此前启动过的所有内核执行完成
     CHECK_CUDA(cudaDeviceSynchronize());
 
     // 2. 仅同步 0 号流（null-stream）
     CHECK_CUDA(cudaStreamSynchronize(0));
 
     // 3. 仅同步 0 号流，但使用小彭老师现代 CUDA 框架
-    CudaStream::nullStream().join();
+    CudaStream::defaultStream().join();
 
     return 0;
 }

cudaguide/1_moderncuda/cudapp.cuh

@@ -1,12 +1,13 @@
 #pragma once
 
-#include "debug.hpp"
 #include <cuda_runtime.h>
+#include <nvfunctional>
 #include <version>
 #include <cstddef>
 #include <cstdio>
 #include <cstdlib>
 #include <cstdarg>
+#include <cstdarg>
 #include <memory>
 #include <new>
 #include <string>
@@ -187,6 +188,8 @@ private:
 
 public:
     CudaMemPool(std::nullptr_t) noexcept {}
+    CudaMemPool(CudaMemPool &&) = default;
+    CudaMemPool &operator=(CudaMemPool &&) = default;
 
     struct Builder {
     private:
@@ -259,12 +262,17 @@ private:
 
 public:
     CudaEvent(std::nullptr_t) noexcept {}
+    CudaEvent(CudaEvent &&) = default;
+    CudaEvent &operator=(CudaEvent &&) = default;
 
     struct Builder {
     private:
         int flags = cudaEventDefault;
 
     public:
+        Builder() = default;
+        explicit Builder(int flags) noexcept : flags(flags) {}
+
         Builder &withBlockingSync(bool blockingSync = true) noexcept {
             if (blockingSync) {
                 flags |= cudaEventBlockingSync;
@@ -303,24 +311,28 @@ public:
         CHECK_CUDA(cudaEventSynchronize(*this));
     }
 
-    bool joinReady() const {
+    bool poll() const {
         cudaError_t res = cudaEventQuery(*this);
         if (res == cudaSuccess) {
             return true;
         }
         if (res == cudaErrorNotReady) {
             return false;
         }
-        CHECK_CUDA(res);
+        CHECK_CUDA(res /* cudaEventQuery */);
         return false;
     }
 
     float elapsedMillis(CudaEvent const &event) const {
         float result;
-        CHECK_CUDA(cudaEventElapsedTime(&result, *this, event));
+        CHECK_CUDA(cudaEventElapsedTime(&result, event, *this));
         return result;
     }
 
+    float operator-(CudaEvent const &event) const {
+        return elapsedMillis(event);
+    }
+
     ~CudaEvent() {
         if (*this) {
             CHECK_CUDA(cudaEventDestroy(*this));
@@ -335,12 +347,17 @@ private:
 
 public:
     CudaStream(std::nullptr_t) noexcept {}
+    CudaStream(CudaStream &&) = default;
+    CudaStream &operator=(CudaStream &&) = default;
 
     struct Builder {
     private:
         int flags = cudaStreamDefault;
 
     public:
+        Builder() = default;
+        explicit Builder(int flags) noexcept : flags(flags) {}
+
         Builder &withNonBlocking(bool nonBlocking = true) noexcept {
             if (nonBlocking) {
                 flags |= cudaStreamNonBlocking;
@@ -357,10 +374,14 @@ public:
         }
     };
 
-    static CudaStream nullStream() noexcept {
+    static CudaStream defaultStream() noexcept {
         return CudaStream(nullptr);
     }
 
+    static CudaStream perThreadStream() noexcept {
+        return CudaStream(cudaStreamPerThread);
+    }
+
     void copy(void *dst, void *src, size_t size, cudaMemcpyKind kind) const {
         CHECK_CUDA(cudaMemcpyAsync(dst, src, size, kind, *this));
     }
@@ -381,11 +402,17 @@ public:
         copy(dst, src, size, cudaMemcpyHostToHost);
     }
 
-    void record(CudaEvent const &event) const {
+    void recordEvent(CudaEvent const &event) const {
         CHECK_CUDA(cudaEventRecord(event, *this));
     }
 
-    void wait(CudaEvent const &event,
+    CudaEvent recordEvent() const {
+        CudaEvent event = CudaEvent::Builder().build();
+        recordEvent(event);
+        return event;
+    }
+
+    void waitEvent(CudaEvent const &event,
               unsigned int flags = cudaEventWaitDefault) const {
         CHECK_CUDA(cudaStreamWaitEvent(*this, event, flags));
     }
@@ -403,22 +430,23 @@ public:
         auto userData = std::make_unique<Func>();
         cudaStreamCallback_t callback = [](cudaStream_t stream,
                                            cudaError_t status, void *userData) {
+            CHECK_CUDA(status /* joinAsync cudaStreamCallback */);
             std::unique_ptr<Func> func(static_cast<Func *>(userData));
-            (*func)(stream, status);
+            (*func)();
         };
         joinAsync(callback, userData.get());
         userData.release();
     }
 
-    bool joinReady() const {
+    bool poll() const {
         cudaError_t res = cudaStreamQuery(*this);
         if (res == cudaSuccess) {
             return true;
         }
         if (res == cudaErrorNotReady) {
             return false;
         }
-        CHECK_CUDA(res);
+        CHECK_CUDA(res /* cudaStreamQuery */);
         return false;
     }
 
@@ -428,7 +456,7 @@ public:
     }
 
     ~CudaStream() {
-        if (*this) {
+        if (*this && *this != cudaStreamPerThread) {
             CHECK_CUDA(cudaStreamDestroy(*this));
         }
     }
@@ -522,8 +550,8 @@ struct CudaAllocator : private Arena {
     };
 };
 
-template <class T>
-using CudaVector = std::vector<T, CudaAllocator<T>>;
+template <class T, class Arena = CudaManagedArena>
+using CudaVector = std::vector<T, CudaAllocator<T, Arena>>;
 
 #if defined(__clang__) && defined(__CUDACC__) && defined(__GLIBCXX__)
 __host__ __device__ static void printf(const char *fmt, ...) {

cudaguide/1_moderncuda/main.cu

@@ -22,26 +22,27 @@ int main() {
 
     // 3. cudaLaunchKernelEx
     cudaLaunchConfig_t cfg{};
-    cfg.blockDim = dim3(3);
-    cfg.gridDim = dim3(4);
-    cfg.dynamicSmemBytes = 0;
-    cfg.stream = 0;
+    cfg.blockDim = dim3(3); // threadIdx 的变化范围
+    cfg.gridDim = dim3(4); // blockIdx 的变化范围
+    cfg.dynamicSmemBytes = 0; // shared-memory 大小（暂不使用）
+    cfg.stream = 0; // 在 0 号流（默认流）上启动
     cfg.attrs = nullptr;
     cfg.numAttrs = 0;
     CHECK_CUDA(cudaLaunchKernelEx(&cfg, kernel, x));
 
+    // 还能查询内核名字
     const char *name;
     CHECK_CUDA(cudaFuncGetName(&name, kernel));
     printf("内核名字：%s\n", name);
 
     // 1. 强制同步：等待此前启动过的所有内核执行完成
     CHECK_CUDA(cudaDeviceSynchronize());
 
-    // 2. 仅同步 0 号流（null-stream）
+    // 2. 仅同步 0 号流（默认流）：等待此前在 0 号流上启动过的所有内核执行完成
     CHECK_CUDA(cudaStreamSynchronize(0));
 
     // 3. 仅同步 0 号流，但使用小彭老师现代 CUDA 框架
-    CudaStream::nullStream().join();
+    CudaStream::defaultStream().join();
 
     return 0;
 }

cudaguide/2_cudastream/CMakeLists.txt

@@ -0,0 +1,33 @@
+cmake_minimum_required(VERSION 3.12)
+
+set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CUDA_STANDARD 20)
+set(CMAKE_CUDA_SEPARABLE_COMPILATION OFF)
+set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --expt-extended-lambda")
+
+if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+    # try to set CMAKE_CUDA_ARCHITECTURES to native card version
+    if (CMAKE_VERSION VERSION_GREATER_EQUAL 3.24)
+        set(CMAKE_CUDA_ARCHITECTURES native)
+    else()
+        message(STATUS "Running __nvcc_device_query for CMAKE_CUDA_ARCHITECTURES")
+        execute_process(COMMAND "__nvcc_device_query"
+            RESULT_VARIABLE NVCC_DEVICE_QUERY_RESULT
+            OUTPUT_VARIABLE CMAKE_CUDA_ARCHITECTURES
+            OUTPUT_STRIP_TRAILING_WHITESPACE)
+        if (NVCC_DEVICE_QUERY_RESULT EQUAL 0)
+            message(STATUS "CMAKE_CUDA_ARCHITECTURES set to ${CMAKE_CUDA_ARCHITECTURES}")
+        else()
+            message(WARNING "CMAKE_CUDA_ARCHITECTURES not defined and __nvcc_device_query failed.")
+            unset(CMAKE_CUDA_ARCHITECTURES)
+        endif()
+    endif()
+endif()
+
+project(cudaradixsort LANGUAGES CXX CUDA)
+
+file(GLOB sources "*.cpp" "*.cu")
+add_executable(${PROJECT_NAME} ${sources})
+target_link_libraries(${PROJECT_NAME} PRIVATE cusparse cublas)
+find_package(fmt REQUIRED)
+target_link_libraries(${PROJECT_NAME} PRIVATE fmt::fmt)