working on capter6

capter6/ReadMe.md

@@ -0,0 +1,56 @@
+# CUDA 的内存组织
+
+< CPU >   
+// ----------------------------------------------------------------------------------  
+// 内存  
+// ----------------------------------------------------------------------------------  
+        ||  
+        ||  
+< GPU > ||  
+// ----------------------------------------------------------------------------------  
+// 全局内存  
+// ----------------------------------------------------------------------------------  
+        ||  
+// ----------------------------------------------------------------------------------   
+// 纹理内存  
+// ----------------------------------------------------------------------------------  
+        ||  
+// ----------------------------------------------------------------------------------  
+// 常量内存  
+// ----------------------------------------------------------------------------------  
+        ||                                                 ||  
+// ---------------------------------------         // -------------------------------  
+// 共享内存 [block0]                                // 共享内存 [block1]                  
+// ---------------------------------------         // -------------------------------  
+        ||                           ||  
+// -----------------------    // -----------------------  
+// 局部内存 [thread00]         // 局部内存 [thread01]            ......  
+// 寄存器                      // 寄存器  
+// -----------------------    // -----------------------  
+
+------
+
+## CUDA 中不同类型的内存
+
+1. **全局内存（global memory）**，即核函数中所有线程都可以访问的内存，可读可写，由主机端分配和释放；  
+如 cudaMalloc() 的设备内存 d_x, d_y, d_z。
+
+全局内存由于没有放到 GPU 芯片上，所以具有较高的延迟和较低的访问速度，但是容量大（显存）。  
+全局内存主要为核函数提供数据，并在主机和设备、设备和设备之间传递数据。
+
+全局内存的生命周期由主机端维护，期间不同的核函数可以多次访问全局内存。
+
+除以上动态分配的全局内存变量外，还可以使用 **静态全局内存变量**，其所占内存数量在编译器确定；  
+这样的静态全局内存变量必须在 所有主机和设备函数外部定义，例如：  
+
+```cuda  
+    __device__ real epsilon;  // 单个静态全局内存变量， `__device` 表示是设备中的变量。
+    __device__ real arr[10];  // 固定长度的静态全局内存数组变量。
+```
+
+对于静态全局内存变量，其访问权限：  
+
+1. 核函数中可以直接访问静态全局内存变量，不必以参数形式传给核函数；
+2. 主机中不可以直接访问静态全局内存变量，可以通过 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 调用。
+
+

capter6/error.cuh

@@ -0,0 +1,26 @@
+#pragma once
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include <cuda_runtime_api.h>
+#include <device_launch_parameters.h>
+
+#define CHECK(call)                                                     \
+do {                                                                    \
+    const cudaError_t error_code = call;                                \
+    if (error_code != cudaSuccess)                                      \
+    {                                                                   \
+        printf("CUDA ERROR: \n");                                       \
+        printf("    FILE: %s\n", __FILE__);                             \
+        printf("    LINE: %d\n", __LINE__);                             \
+        printf("    ERROR CODE: %d\n", error_code);                     \
+        printf("    ERROR TEXT: %s\n", cudaGetErrorString(error_code)); \
+        exit(1);                                                        \
+    }                                                                   \
+}while(0);                                                              \
+
+
+
+
+
+
+

capter6/static.cu

@@ -0,0 +1,54 @@
+#include "error.cuh"
+
+
+// 静态全局内存变量。  
+__device__ int d_x = 1;
+__device__ int d_y[2] = {2, 3};
+
+
+// 核函数。
+__global__ void add_array()
+{
+    d_y[0] += d_x;
+    d_y[1] += d_x;
+
+    printf("d_y: {%d, %d}\n", d_y[0], d_y[1]);
+}
+__global__ void add_var()
+{
+    d_x += 2;
+
+    printf("d_x: %d\n", d_x);
+}
+__global__ void display()
+{
+    printf("d_x: %d, d_y: {%d, %d}\n", d_x, d_y[0], d_y[1]);
+}
+
+
+int main()
+{
+    display<<<1, 1>>>();
+    add_array<<<1, 1>>>();
+    add_var<<<1, 1>>>();
+    CHECK(cudaDeviceSynchronize());
+    
+    int h_y[2] = {10, 20};
+
+    // cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, 
+    //     size_t offset=0, cudaMemcpyKind kind=cudaMemcpyHostToDevice);
+    CHECK(cudaMemcpyToSymbol(d_y, h_y, sizeof(int)));
+    display<<<1, 1>>>();
+
+    // cudaError_t cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, 
+    //     size_t offset=0, cudaMemcpyKind kind=cudaMemcpyDeviceToHost);
+    CHECK(cudaMemcpyFromSymbol(h_y, d_y, sizeof(int) * 2));
+    printf("host, d_y: %d, %d\n", h_y[0], h_y[1]);
+    display<<<1, 1>>>();
+
+    return 0;
+}
+
+
+
+