[hw08] done

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 52)
 
 project(hellocmake LANGUAGES CXX CUDA)
 

main.cu

@@ -8,19 +8,27 @@
 
 // 这是基于“边角料法”的，请把他改成基于“网格跨步循环”的：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);
+    // int i = blockIdx.x * blockDim.x + threadIdx.x;
+    for(int i=blockIdx.x * blockDim.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;
+        // ans: 多线程会同时访问和修改counter，需要使用原子操作。
+        // int loc = *counter;
+        // atomicAdd(&)
+        // *counter += 1;
+        // res[loc] = n;
+
+        int loc = atomicAdd(counter, 1);
         res[loc] = n;
+        // printf("%d\n", loc);
     }
 }
 
@@ -31,20 +39,28 @@ int main() {
     std::vector<int, CudaAllocator<int>> counter(1);
 
     // fill_sin 改成“网格跨步循环”以后，这里三重尖括号里的参数如何调整？10 分
-    fill_sin<<<n / 1024, 1024>>>(arr.data(), n);
+    // ans: 将block数调整为32即可。
+    fill_sin<<<32, 1024>>>(arr.data(), n);
 
     // 这里的“边角料法”对于不是 1024 整数倍的 n 会出错，为什么？请修复：10 分
-    filter_positive<<<n / 1024, 1024>>>(counter.data(), res.data(), arr.data(), n);
+    // ans: 会缺少数据，应该向上取整，让总线程数大于n。
+    filter_positive<<<(n+1024-1) / 1024, 1024>>>(counter.data(), res.data(), arr.data(), n);
+    // cudaDeviceSynchronize();
 
     // 这里 CPU 访问数据前漏了一步什么操作？请补上：10 分
+    // ans: 要将数据从cuda拷贝到内存上。
+    int counter_cpu;
+    std::vector<int> res_cpu(n);
+    checkCudaErrors(cudaMemcpy(&counter_cpu, counter.data(), sizeof(int), cudaMemcpyDeviceToHost));
+    checkCudaErrors(cudaMemcpy(res_cpu.data(), res.data(), sizeof(int)*res.size(), cudaMemcpyDeviceToHost));
 
-    if (counter[0] <= n / 50) {
-        printf("Result too short! %d <= %d\n", counter[0], n / 50);
+    if (counter_cpu <= n / 50) {
+        printf("Result too short! %d <= %d\n", counter_cpu, 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]);
+    for (int i = 0; i < counter_cpu; i++) {
+        if (res[i] != n) {
+            printf("Wrong At %d: %d < 0\n", i, res[i]);
             return -1;  // 突然想起了ICPC有一年队名叫“蓝翔WA掘机”的，笑不活了:)
         }
     }