hw08/cuda

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)
 

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,20 +31,22 @@ 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);
         return -1;
     }
     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掘机”的，笑不活了:)
         }
     }