From 49fe367e6d5b91c99f013e87a3463b6622d91bee Mon Sep 17 00:00:00 2001
From: Chen Hu <86994637+Tigerrr07@users.noreply.github.com>
Date: Wed, 30 Mar 2022 17:28:31 +0800
Subject: [PATCH] hw08/cuda

---
 CMakeLists.txt |  2 +-
 main.cu        | 20 +++++++++++---------
 2 files changed, 12 insertions(+), 10 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 07d8dc9..6cc7b9d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -6,7 +6,7 @@ if (NOT CMAKE_BUILD_TYPE)
     set(CMAKE_BUILD_TYPE Release)
 endif()
 # 如果需要指定显卡版本号的话：
-# set(CMAKE_CUDA_ARCHITECTURES 52)
+set(CMAKE_CUDA_ARCHITECTURES 70)
 
 project(hellocmake LANGUAGES CXX CUDA)
 
diff --git a/main.cu b/main.cu
index 018bc33..ae87097 100644
--- a/main.cu
+++ b/main.cu
@@ -8,18 +8,18 @@
 
 // 这是基于“边角料法”的，请把他改成基于“网格跨步循环”的：10 分
 __global__ void fill_sin(int *arr, int n) {
-    int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < n) return;
-    arr[i] = sinf(i);
+    // 基于网格跨步循环，自动根据n指定blockDim和gridDim
+    for (size_t i = blockDim.x * blockIdx.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x)
+        arr[i] = sinf(i);
 }
 
 __global__ void filter_positive(int *counter, int *res, int const *arr, int n) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < n) return;
+    if (i >= n) return;
     if (arr[i] >= 0) {
         // 这里有什么问题？请改正：10 分
-        int loc = *counter;
-        *counter += 1;
+        // 会有其他线程干扰对counter+1的过程，应替代为原子操作
+        int loc = atomicAdd(counter, 1);
         res[loc] = n;
     }
 }
@@ -31,12 +31,14 @@ int main() {
     std::vector<int, CudaAllocator<int>> counter(1);
 
     // fill_sin 改成“网格跨步循环”以后，这里三重尖括号里的参数如何调整？10 分
-    fill_sin<<<n / 1024, 1024>>>(arr.data(), n);
+    fill_sin<<<128, 1024>>>(arr.data(), n);
 
     // 这里的“边角料法”对于不是 1024 整数倍的 n 会出错，为什么？请修复：10 分
-    filter_positive<<<n / 1024, 1024>>>(counter.data(), res.data(), arr.data(), n);
+    // 出错：若n = 1025，n / 1024 = 1，总共会启动1*1024=1024个线程，会漏掉1个元素
+    filter_positive<<<(n + 1024 - 1)/ 1024, 1024>>>(counter.data(), res.data(), arr.data(), n);
 
     // 这里 CPU 访问数据前漏了一步什么操作？请补上：10 分
+    cudaDeviceSynchronize();
 
     if (counter[0] <= n / 50) {
         printf("Result too short! %d <= %d\n", counter[0], n / 50);
@@ -44,7 +46,7 @@ int main() {
     }
     for (int i = 0; i < counter[0]; i++) {
         if (res[i] < 0) {
-            printf("Wrong At %d: %f < 0\n", i, res[i]);
+            printf("Wrong At %d: %d < 0\n", i, res[i]);
             return -1;  // 突然想起了ICPC有一年队名叫“蓝翔WA掘机”的，笑不活了:)
         }
     }