add capter6

capter6/ReadMe.md

@@ -32,7 +32,14 @@
 
 ## CUDA 中不同类型的内存
 
-1. **全局内存（global memory）**，即核函数中所有线程都可以访问的内存，可读可写，由主机端分配和释放；  
+CUDA 中的内存类型有：全局内存、常量内存、纹理内存、寄存器、局部内存、共享内存。
+CUDA 的内存，即设备内存，主机无法直接访问。
+
+------
+
+### 全局内存
+
+**全局内存（global memory）**，即核函数中所有线程都可以访问的内存，可读可写，由主机端分配和释放；  
 如 cudaMalloc() 的设备内存 d_x, d_y, d_z。
 
 全局内存由于没有放到 GPU 芯片上，所以具有较高的延迟和较低的访问速度，但是容量大（显存）。  
@@ -51,6 +58,92 @@
 对于静态全局内存变量，其访问权限：  
 
 1. 核函数中可以直接访问静态全局内存变量，不必以参数形式传给核函数；
-2. 主机中不可以直接访问静态全局内存变量，可以通过 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 调用。
+2. 主机中不可以直接访问静态全局内存变量，可以通过 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 调用。  
 
+------
 
+### 常量内存
+
+**常量内存（constant memory）**，仅有 64 kb，可见范围和生命周期与全局内存一样；具有缓存，从而高速；
+常量内存仅可读、不可写。  
+
+使用常量内存的方法：一是在核函数外定义常量内存变量；二是向核函数传递常量参数，默认存放在常量内存：  
+
+1. 核函数中可以直接访问常量全局内存变量，不必以参数形式传给核函数，但不可更改（只读）；
+2. 主机中不可以直接访问常量全局内存变量，可以通过 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 调用。
+
+------
+
+### 纹理内存
+
+**纹理内存(texture memory)**，类似常量内存，也是一种具有缓存的全局内存，具有相同可见范围和生命周期。
+
+可以将某些只读的全局内存数据用 `__ldg()` 函数通过只读数据缓存（read-only data cache）读取，  
+既可以达到使用纹理内存的加速效果，又可使代码简洁：  
+    
+    int __ldg(const int* ptr);  // 函数原型。
+
+全局内存的读取在默认情况下就利用了 `__ldg()` 函数，所以不需要显式地使用。
+
+------
+
+### 寄存器
+
+在核函数中定义的、不加任何限定符的变量一般存放在寄存器（register）；核函数中不加任何限定符的数组可能放在  
+寄存器，也可能放在局部内存中。寄存器可读可写。
+
+各种内建变量，如 gridDim、blockDim 等都保存在特殊的寄存器中。
+
+寄存器变量仅被一个线程看见，寄存器的生命周期也和所属线程相同。
+
+寄存器内存在芯片上，是所有内存中访问速度最高的。一个寄存器占 32b（4字节），一个双精度浮点数占 2个寄存器。
+
+------
+
+### 局部内存
+
+局部内存(local memory)也是全局内存的一部分，每个线程最多可以使用 512 kb 的局部内存，但过多使用会降低性能。
+局部内存的用法类似寄存器。
+
+------
+
+### 共享内存
+
+共享内存（shared memory）与寄存器类似，都是位于芯片上，读写速度较快。
+
+共享内存对整个线程块可见，一个线程块上的所有线程都可以访问共享内存上的数据；共享内存的生命  
+周期也与所属线程块一致。
+
+共享内存的主要作用是减少对全局内存的访问，或者改善对全局内存的访问模式。
+
+------
+
+### L1 和 L2 缓存
+
+SM 层次的 L1 缓存（一级缓存）和设备层次 L2 缓存（二级缓存）。它们主要用来缓存全局内存和设备内存的访问。
+
+------
+
+## SM 及其占有率
+
+一个 GPU 由多个 SM（流多处理器）构成，一个 SM 包含如下资源：
+
+1. 一定数量的寄存器；
+2. 一定数量的共享内存；
+3. 常量内存的缓存；
+4. 纹理内存的缓存；
+5. L1 缓存；
+6. 两个或四个线程束调度器，用于在不同线程上下文间迅速切换，及为准备就绪的线程束发出执行指令；
+7. 执行核心。
+
+一般来说，要尽量让 SM 的占有率不小于某值（如 25%），才有可能获得较高的性能。
+
+SM 中线程的执行是以线程束为单位的，所以最好将线程块大小取为线程束大小（32个线程）的整数倍（如 128）.
+
+------
+
+## CUDA 运行时 API 函数查询设备
+
+使用 CUDA 运行时 API 函数查询所用GPU 规格。
+
+------

capter6/query.cu

@@ -0,0 +1,31 @@
+#include "error.cuh"
+#include <stdlib.h>
+
+
+int main(int argc, char *argv[])
+{
+    int device_id = 0;
+    if (argc > 1) device_id = atoi(argv[1]);
+
+    CHECK(cudaSetDevice(device_id));
+
+    cudaDeviceProp prop;
+    CHECK(cudaGetDeviceProperties(&prop, device_id));
+
+    printf("Device id: %d\n", device_id);
+    printf("Device name: %s\n", prop.name);
+    printf("Compute capability: %d.%d\n", prop.major, prop.minor);
+    printf("Amount of global memory: %g GB\n", prop.totalGlobalMem/(1024.0*1024*1024));
+    printf("Amount of constant memory: %g KB\n", prop.totalConstMem/1024.0);
+    printf("Maximum grid size: %d, %d, %d\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
+    printf("Maximum block size: %d, %d, %d\n", prop.maxThreadsDim[0],prop.maxThreadsDim[1],prop.maxThreadsDim[2]);
+    printf("Number of SMs: %d\n", prop.multiProcessorCount);
+    printf("Maximum amount of shared memory per block: %g KB\n", prop.sharedMemPerBlock/1024.0);
+    printf("Maximum amount of shared memory per SM: %g KB\n", prop.sharedMemPerMultiprocessor/1024.0);
+    printf("Maximum number of registers per block: %d K\n", prop.regsPerBlock/1024);
+    printf("Maximum number of registers per SM: %d K\n", prop.regsPerMultiprocessor/1024);
+    printf("Maximum number of threads per block: %d\n", prop.maxThreadsPerBlock);
+    printf("Maximum number of threads per SM: %d\n", prop.maxThreadsPerMultiProcessor);
+
+    return 0;
+}

capter6/static.cu

@@ -1,11 +1,16 @@
 #include "error.cuh"
 
 
-// 静态全局内存变量。  
+// 静态全局内存变量(设备内存)。  
 __device__ int d_x = 1;
 __device__ int d_y[2] = {2, 3};
 
 
+// 常量内存变量（设备内存）。
+__constant__ double d_m = 23.33;
+__constant__ double d_n[] = {12.2, 34.1, 14.3};
+
+
 // 核函数。
 __global__ void add_array()
 {
@@ -24,6 +29,12 @@ __global__ void display()
 {
     printf("d_x: %d, d_y: {%d, %d}\n", d_x, d_y[0], d_y[1]);
 }
+__global__ void show()
+{
+    // 常量内存变量在核函数中不可更改（只读）。
+
+    printf("d_m: %f, d_n: {%f, %f, %f}\n", d_m, d_n[0], d_n[1], d_n[2]);
+}
 
 
 int main()
@@ -32,20 +43,38 @@ int main()
     add_array<<<1, 1>>>();
     add_var<<<1, 1>>>();
     CHECK(cudaDeviceSynchronize());
-    
+
+    show<<<1, 1>>>();
+    CHECK(cudaDeviceSynchronize());
+
     int h_y[2] = {10, 20};
+    int h_x = 7;
+
+    double h_m = 22.23;
+    double h_n[3] = {1.1, 2.2, 3.3};
 
     // 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)));
+    CHECK(cudaMemcpyToSymbol(d_x, &h_x, sizeof(int)));  // 对于全局内存变量，非数组对象不用取址。
     display<<<1, 1>>>();
 
+    CHECK(cudaMemcpyToSymbol(d_m, &h_m, sizeof(double)));  // 对于常量内存变量，非数组对象不用取址。 
+    CHECK(cudaMemcpyToSymbol(d_n, h_n, sizeof(double)*3));
+    show<<<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]);
+    CHECK(cudaMemcpyFromSymbol(&h_x, d_x, sizeof(int)));  // 对于全局内存变量，非数组对象不用取址。
+    printf("host, h_y: %d, %d, h_x: %d\n", h_y[0], h_y[1], h_x);
     display<<<1, 1>>>();
 
+    CHECK(cudaMemcpyFromSymbol(h_n, d_n, sizeof(double)*3));
+    CHECK(cudaMemcpyFromSymbol(&h_m, d_m, sizeof(double)));  // 对于常量内存变量，非数组对象不用取址。  
+    printf("host, h_n: %f, %f, h_m: %f\n", h_n[0], h_n[1], h_m);
+    show<<<1, 1>>>();
+
     return 0;
 }