Complete saxpy

Author: mialana

cuda-introduction/.vscode/c_cpp_properties.json

@@ -0,0 +1,17 @@
+{
+    "configurations": [
+        {
+            "name": "Linux",
+            "includePath": [
+                "${workspaceFolder}/source/**"
+            ],
+            "defines": [],
+            "compilerPath": "/usr/local/cuda/bin/nvcc",
+            "cStandard": "gnu17",
+            "cppStandard": "gnu++17",
+            "intelliSenseMode": "linux-gcc-x64",
+            "configurationProvider": "ms-vscode.makefile-tools"
+        }
+    ],
+    "version": 4
+}

cuda-introduction/.vscode/extension.json

@@ -0,0 +1,7 @@
+{
+    "recommendations": [
+        "nvidia.nsight-vscode-edition",
+        "ms-vscode.cpptools",
+        "ms-vscode.makefile-tools"
+    ]
+}

cuda-introduction/.vscode/launch.json

@@ -3,9 +3,10 @@
     "configurations": [
         {
             "name": "CUDA C++: Launch",
+            "preLaunchTask": "CMake: build",
             "type": "cuda-gdb",
             "request": "launch",
-            "program": "${workspaceFolder}/build/bin/saxpy",
+            "program": "${command:cmake.launchTargetPath}",
         }
     ]
 }

cuda-introduction/.vscode/settings.json

@@ -1,6 +1,63 @@
 {
     "files.associations": {
         "cpp": "cuda-cpp",
-        "hpp": "cuda-cpp"
+        "hpp": "cuda-cpp",
+        "array": "cpp",
+        "atomic": "cpp",
+        "bit": "cpp",
+        "cctype": "cpp",
+        "charconv": "cpp",
+        "chrono": "cpp",
+        "clocale": "cpp",
+        "cmath": "cpp",
+        "compare": "cpp",
+        "concepts": "cpp",
+        "cstdarg": "cpp",
+        "cstddef": "cpp",
+        "cstdint": "cpp",
+        "cstdio": "cpp",
+        "cstdlib": "cpp",
+        "cstring": "cpp",
+        "ctime": "cpp",
+        "cwchar": "cpp",
+        "cwctype": "cpp",
+        "deque": "cpp",
+        "string": "cpp",
+        "unordered_map": "cpp",
+        "vector": "cpp",
+        "exception": "cpp",
+        "algorithm": "cpp",
+        "functional": "cpp",
+        "iterator": "cpp",
+        "memory": "cpp",
+        "memory_resource": "cpp",
+        "numeric": "cpp",
+        "optional": "cpp",
+        "random": "cpp",
+        "ratio": "cpp",
+        "string_view": "cpp",
+        "system_error": "cpp",
+        "tuple": "cpp",
+        "type_traits": "cpp",
+        "utility": "cpp",
+        "format": "cpp",
+        "initializer_list": "cpp",
+        "iomanip": "cpp",
+        "iosfwd": "cpp",
+        "iostream": "cpp",
+        "istream": "cpp",
+        "limits": "cpp",
+        "new": "cpp",
+        "numbers": "cpp",
+        "ostream": "cpp",
+        "queue": "cpp",
+        "ranges": "cpp",
+        "span": "cpp",
+        "sstream": "cpp",
+        "stdexcept": "cpp",
+        "streambuf": "cpp",
+        "text_encoding": "cpp",
+        "typeinfo": "cpp",
+        "variant": "cpp"
     }
 }

cuda-introduction/.vscode/tasks.json

@@ -0,0 +1,16 @@
+{
+    "version": "2.0.0",
+    "tasks": [
+        {
+            "type": "cmake",
+            "label": "CMake: build",
+            "command": "build",
+            "targets": [
+                "${command:cmake.buildTargetName}"
+            ],
+            "group": "build",
+            "problemMatcher": [],
+            "detail": "CMake template build task"
+        },
+    ]
+}

cuda-introduction/source/common.cu

@@ -9,7 +9,7 @@ unsigned divup(unsigned size, unsigned div)
 {
     // TODO: implement a 1 line function to return the divup operation.
     // Note: You only need to use addition, subtraction, and division operations.
-    return 0;
+    return (size + div - 1) / div;
 }
 
 void clearHostAndDeviceArray(float *res, float *dev_res, unsigned size, const int value)
@@ -26,6 +26,11 @@ bool compareReferenceAndResult(const float *ref, const float *res, unsigned size
     bool passed = true;
     for (unsigned i = 0; i < size; i++)
     {
+#if 0
+        // print thread info regardless of fail or not
+        std::cout << "ID: " << i << " \t Res: " << res[i] << " \t Ref: " << ref[i] << std::endl;
+#endif
+
         // LOOK: Check if floating point values are equal within an epsilon as returns can vary slightly between CPU and GPU
         if (std::fabs(res[i] - ref[i]) > epsilon)
         {

cuda-introduction/source/saxpy.cu

@@ -9,20 +9,22 @@
 __global__ void saxpy(float* const z, const float* const x, const float* const y, const float a, const unsigned size)
 {
     // TODO 9: Compute the global index for each thread.
-    unsigned idx = 0;
+    unsigned idx = threadIdx.x;
 
     // TODO 10: Check if idx is out of bounds. If yes, return.
-    if (idx >= 0)
-        return;
+    if (idx >= size)
+        {return;}
 
     // TODO 11: Perform the SAXPY operation: z = a * x + y.
+
+    z[idx] = a * x[idx] + y[idx];
 }
 
 int main(int argc, char *argv[])
 {
     // TODO 1: Set the size. Start with something simple like 64.
     // TODO Optional: Try out these sizes: 256, 1024, 2048, 14, 103, 1025, 3127
-    const unsigned size = 0;
+    const unsigned size = 64;
 
     // Host arrays.
     float* x = new float[size];
@@ -51,11 +53,16 @@ int main(int argc, char *argv[])
     // Device arrays
     float *d_x, *d_y, *d_z;
 
+    int sizeInBytes = size * sizeof(float);
+
     // TODO 2: Allocate memory on the device. Fill in the blanks for d_x, then do the same commands for d_y and d_z.
-    // CUDA(cudaMalloc((void **)& pointer, size in bytes)));
+    CUDA(cudaMalloc((void **)& d_x, sizeInBytes));
+    CUDA(cudaMalloc((void **)& d_y, sizeInBytes));
+    CUDA(cudaMalloc((void **)& d_z, sizeInBytes));
 
     // TODO 3: Copy array contents of X and Y from the host (CPU) to the device (GPU). Follow what you did for 2,
-    // CUDA(cudaMemcpy(dest ptr, source ptr, size in bytes, direction enum));
+    CUDA(cudaMemcpy(d_x, x, sizeInBytes, cudaMemcpyHostToDevice));
+    CUDA(cudaMemcpy(d_y, y, sizeInBytes, cudaMemcpyHostToDevice));
 
     CUDA(cudaDeviceSynchronize());
 
@@ -69,17 +76,20 @@ int main(int argc, char *argv[])
     // TODO 4: Setup threads and blocks.
     // Start threadPerBlock as 128, then try out differnt configurations: 32, 64, 256, 512, 1024
     // Use divup to get the number of blocks to launch.
-    const unsigned threadsPerBlock = 0;
+    const unsigned threadsPerBlock = 128;
 
     // TODO 5: Implement the divup function in common.cpp
     const unsigned blocks = divup(size, threadsPerBlock);
 
     // TODO 6: Launch the GPU kernel with blocks and threadPerBlock as launch configuration
-    // saxpy<<< >>> (....);
+    saxpy<<<blocks, threadsPerBlock>>> (d_z, d_x, d_y, a, size);
 
     // TODO 7: Copy the answer back to the host (CPU) from the device (GPU).
     // Copy what you did in 3, except for d_z -> z.
 
+    cudaMemcpy(z, d_z, sizeInBytes, cudaMemcpyDeviceToHost);
+
+
     // LOOK: Use postprocess to check the result
     compareReferenceAndResult(z_gold, z, size, 1e-6);
     std::cout << "****************************************************" << std::endl << std::endl;
@@ -88,6 +98,10 @@ int main(int argc, char *argv[])
     // TODO 8: free device memory using cudaFree
     // CUDA(cudaFree(device pointer));
 
+    CUDA(cudaFree(d_x));
+    CUDA(cudaFree(d_y));
+    CUDA(cudaFree(d_z));
+
     // free host memory
     delete[] x;
     delete[] y;