CUDA_C_编程权指南 professional CUDA C++ programming chapter two CUDA编程模型 Programming Model_cuda c编程权威指南用的是什么软件

CUDA C 编程权指南-professional CUDA C++ programming chapter-two CUDA编程模型Programming Model

1. Introduction — CUDA C Programming Guide (nvidia.com)

CUDA Runtime API :: CUDA Toolkit Documentation (nvidia.com)

CUDA C编程权指南 professional CUDA C programming​product.dangdang.com/25089854.html

以下的内容主要来自这个页面 ：

1. Introduction — CUDA C Programming Guide (nvidia.com)

2.1. Kernels 核函数

CUDA C++通过允许编程者定义C++函数来扩展C++，也叫做核函数，当核函数被调用的时候，会在N个不同的CUDA线程上被执行N次，而不像C++函数只执行一次。

核函数使用标识符 global 来标志，执行的CUDA线程数量通过扩展的C++执行配置三重括号<<<>>>来指定，见

C++ Language Extensions 。每个执行核函数的线程都给定了一个独特的线程ID，可在核函数内通过内置变量 threadIdx.x、threadIdx.y、threadIdx.z来访问。

下面的program，使用了内置的变量threadIdx，对A和B的N个数值相加，最后存放到数组C内。<<<1, N>>>执行配置代表了blocksPerGrid=每个grid网格的线程块block数量是1，threadsPerBlock=每个线程块block的线程thread数量是N。执行核函数VecAdd的N个线程每个都执行了一次 加和+运算。

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;  ## 某个线程块block内的该线程的标号index
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);  ## VecAdd<<<blocksPerGrid, threadsPerBlock>>>
    ...
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

2.2. Thread Hierarchy 线程的层次结构

是结构体包含了三个变量x, y, z。所以线程能被组织到1dim, 2dim, 3dim的线程标号index。 也就形成了1dim, 2dim, 3dim的线程块block。提供了便捷的方式来计算 向量、矩阵 或者 3dim数据。某个线程的index或者它的线程ID可以这么来看：对1dim的线程块，只需要用到threadIdx.x来index，它的线程ID也是threadIdx.x；对2dim的线程块，假设线程块大小是(Dx, Dy)，则需要用到(threadIdx.x, threadIdx.y)来index，它的线程ID是threadIdx.x + threadIdx.y * Dx；对3dim的线程块，假设线程块大小是(Dx, Dy，Dz)，则需要用到(threadIdx.x, threadIdx.y, threadIdx.z)来index，它的线程ID是threadIdx.x + threadIdx.y * Dx + threadIdx.z * Dx * Dy 。

下面的program，累加了NxN大小的矩阵A和B，并保存到矩阵C，dim3 threadsPerBlock(N, N);每个线程块block的线程数量是NxN的2dim，所以实际的index需要threadIdx.x和threadIdx.y。

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
                       float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

每个线程块block的线程数量是有限的，因一个block内所有的threads都应该在相同的流多处理器核内，共享流多处理器核的有限内存。当前的限制是一个block最多仅可包含1026-2个线程。

具有相同形状的线程块block能同时执行同一个核函数，执行核函数的总线程数量=每个block的线程数量 x block的个数。线程块能被组织到1dim, 2dim, 3dim的网格形状，网格内线程块block的个数通常由数据量决定。

每个block的线程数量，以及每个网格的block线程块数量，通过执行配置三重括号来指定<<<blocksPerGrid, threadsPerBlock>>>，可以是int类型的，也可以就是dim3结构体类型的。上面的program给出了2dim的配置。

blockDim内置变量给出了线程块的dim，也就是包含的线程数量；网格grid内的线程块block的标号在内置变量blockIdx内。上面的program改到多个线程块block，现在网格grid内的线程块block也是2dim的了，线程块内的线程也是2dim的。blockIdx.x 网格grid内x方向的线程块标号 index，blockDim.x 某一个线程块x方向包含的线程数量，threadIdx.x 该线程在所在线程块内x方向的标号 index。

矩阵的每个element都分配了一个线程

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{   // blockDim.x 某一个线程块x方向包含的线程数量
    // blockIdx.x 网格grid内x方向的线程块标号 index
    // threadIdx.x 该线程在所在线程块内x方向的标号 index
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16); // 每个线程块block的线程数量，x方向16个线程，y方向16个线程
    // 每个网格Grid内线程块block的数量，x方向、y方向
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22

线程块内的线程能通过共享内存来合作，通过同步操作。指定同步点的内置函数是 __syncthreads()，作用在线程块内部。会阻塞起来直到线程块内所有线程都完成了__syncthreads语句之前的操作。协作组提供了更加丰富的线程同步机制。

Cooperative Groups API

要高效的执行，共享内存需要低延迟，在每个处理器核附近，更像L1缓存，__syncthreads也是轻量语句的。全局内存速度慢很多。

2.2.1. Thread Block Clusters 线程块簇

自从运算能力9.0开始，CUDA编程模型引入了层次结构-线程块簇，由很多线程块组成。和流多处理器上的线程块内线程组织方式类似，GPU上的GPU处理簇GPC组织了一个簇内的线程块。

和线程块类似，线程块簇也存在1dim, 2dim, 3dim，一个簇内的线程块数量可以由用户定义，CUDA内当前最多允许一个簇内有eight个线程块的。支持的数量可通过API来查询 。

当使用簇的时候，可通过

Cluster Group API 来查询，考虑到兼容性，之前的内置变量含义都不变。gridDim还是每个grid内的block数量。

线程块簇的使用方式：1、编译时的核函数属性__cluster_dims__(X,Y,Z)，放在函数声明那，2、或者用CUDA核函数执行API cudaLaunchKernelEx。下面的program编译时确定了簇的大小，所以运行时就不能修改了。

// Kernel definition
// Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{

}

int main()
{
    float *input, *output;
    // Kernel invocation with compile time cluster size
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // The grid dimension is not affected by cluster launch, and is still enumerated
    // using number of blocks.
    // The grid dimension must be a multiple of cluster size.
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

要在运行时指定簇的大小，就需要用到API cudaLaunchKernelEx，

// Kernel definition
// No compile time attribute attached to the kernel
__global__ void cluster_kernel(float *input, float* output)
{

}

int main()
{
    float *input, *output;
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // Kernel invocation with runtime cluster size
    {
        cudaLaunchConfig_t config = {0};
        // The grid dimension is not affected by cluster launch, and is still enumerated
        // using number of blocks.
        // The grid dimension should be a multiple of cluster size.
        config.gridDim = numBlocks;
        config.blockDim = threadsPerBlock;

        cudaLaunchAttribute attribute[1];
        attribute[0].id = cudaLaunchAttributeClusterDimension;
        attribute[0].val.clusterDim.x = 2; // Cluster size in X-dimension
        attribute[0].val.clusterDim.y = 1;
        attribute[0].val.clusterDim.z = 1;
        config.attrs = attribute;
        config.numAttrs = 1;

        cudaLaunchKernelEx(&config, cluster_kernel, input, output);
    }
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34

在运算能力9.0的GPU内，簇内的所有线程块都组织在一个单个的GPC也就是GPU处理簇内的。允许簇内的线程块使用

Cluster Group API来执行硬件层级的同步操作。簇组也提供了成员函数来查询簇组的大小，包括线程的数量 num_threads() API、线程块的数量 num_blocks() API。簇组内的线程或者线程块标号index，可以通过这两个API来查询 andAPI。同一个簇内的线程块能访问分布式共享内存Distributed Shared Memory，在DSM内，同一个簇的所有线程能在分布式共享内存上读、写、执行atomic操作。

2.3. Memory Hierarchy 内存的层级

CUDA在执行期间会访问很多内存空间来获取数据，每个线程都有私有的本地内存和寄存器，每个线程块都有共享内存，线程块内的所有线程都可见共享内存，共享内存的生命周期和线程块相同。在同一个线程簇内的所有线程块都能在每个的共享内存上读、写、atomic操作。所有的线程都能访问同一个全局内存。

两个额外的只读内存，所有线程都能访问，也就是常量内存和texture内存。全局内存、常量内存和texture内存分别$不同目的做了优化操作，texture内存提供了不同的寻址模式以及data filtering。详见：

Texture and Surface Memory。全局、常量和texture内存在核函数执行期间都是可见的。线程块簇的分布式共享内存是由线程块的共享内存组成的。

2.4. Heterogeneous Programming

像下图那样，CUDA编程模型假设CUDA线程在GPU device上执行C++ program，是主机的协处理器。核函数是在GPU上执行，但是剩下的C++ program在CPU上执行。

CUDA编程模型还假设了主机和设备都有独立的内存，也就是主机内存和设备内存，通过CUDA运行时API可以管理全局内存、常量内存和texture内存的，包括内存分配、释放，主机和设备之间的内存数据传输。

统一内存管理，提供了管理主机和设备内存的桥梁，管理内存是将主机和设备的所有内存看作一个单个的、协作的内存块有着同一个统一的地址空间。这个能力允许分配>设备内存的空间，通过消除在主机和设备之间传输数据能极大的简化应用移植。详见：

Unified Memory Programming

2.5. Asynchronous SIMT Programming Model

异步单指令多数据编程模型，CUDA编程模型内，一个线程是内存操作或者运算的最低抽象层次。自从nvidia ampere 架构开始，CUDA编程模型通过异步编程提供了内存操作的加速。异步编程模型是指CUDA线程的异步操作。

异步编程模型定义了CUDA线程同步时的异步阻塞，异步编程模型能解释 API

cuda::memcpy_async　是怎么在GPU运算的同时从全局内存异步移动数据的。

2.5.1. Asynchronous Operations 异步操作的

一个异步操作是指通过一个CUDA线程来初始化一个操作，由另一个线程来异步执行。一个或者多个CUDA线程使用异步操作来同步，初始化异步操作的CUDA线程不要求在同步线程内。

一个异步线程总是和初始化异步操作的这个CUDA线程相关，一个异步操作使用一个同步对象来完成当前操作，这样的异步对象能被用户管理，像：cuda::memcpy_async，或者间接的在库内管理 cooperative_groups::memcpy_async。

一个同步对象可以是or a，同步对象能在不同的命名空间scope内使用，一个命名空间是指可能使用同步对象去同步异步操作的线程集合。

2.6. Compute Capability 运算能力

设备的运算能力可见版本号，有时候也叫做SM版本，版本号给出了GPU硬件的性能，被应用用来决定当前GPU的特性。运算能力有一个主版本号和一个小版本号。主版本号相同的设备有相同的核心架构，主版本号是9代表了architecture，主版本号是8代表了architecture。

小版本号代表了核心架构的持续提升，可能包括了new feature。

CUDA-Enabled GPUs 这个页面列出了所有的CUDA设备和它们的运算能力，

Compute Capabilities 给出了所有的技术指标。特定运算能力的GPU不应该和CUDA版本混淆起来，CUDA一般会兼容前几代GPU，但是运算能力不相同的GPU的核心架构是不相同的。CUDA是应用开发者用来开发可以运行在好几代GPU上的应用程序。GPU是硬件设备的。

以下内容来自书籍

《CUDA C编程权威指南》([美]程润伟（John Cheng）)【简介_书评_在线阅读】 - 当当图书 (dangdang.com)

CUDA是一种通用的并行计算平台和编程模型，在C语言基础上扩展的。使用CUDA可以像编写C语言一样实现并行的算法。这个chapter内向量相加的例子都可以在GitHub的repository内找到

vectorAdd CUDA sample.

cuda-samples/Samples/0_Introduction/vectorAdd/vectorAdd.cu at master · NVIDIA/cuda-samples (github.com)​github.com/NVIDIA/cuda-samples/blob/master/Samples/0_Introduction/vectorAdd/vectorAdd.cu

2.1 CUDA编程模型概述

CUDA编程模型提供了一个计算机架构抽象来作$应用程序和其可用的硬件之间的桥梁。GPU编程模型根据GPU架构的计算能力，提供了以下几个特有的功能：1、一种通过层次结构在GPU中组织线程的方式，2、一种通过层次结构在GPU中访问内存的方式。

2.1.1 CUDA编程结构

“统一寻址”（Unified Memory）的编程模型的改进，主机内存和设备内存看作统一的同一个内存，不需要在host和device之间复制数据。核函数是在GPU上运行。内核被启动以后，管理权返回给主机，设备上运行核函数，CPU可以做其他的事情，CUDA编程模型是异步的，GPU执行并行运算，CPU执行串行运算，

2.1.2 内存管理

CUDA编程模型假设主机和设备拥有独立的内存，CUDA运行时API负责分配和释放内存，并在两者之间传输数据。对应的函数是这些，cudaMalloc和Malloc几乎是相同的，只是分配的内存地方不相同。cudaMemcpy函数负责主机和设备之间的数据传输，cudaMemcpy的第一个参数是目的地，第二个参数是source，第三个是方向的呢。cudaMemcpy是同步执行的，数据传输完成以前，主机的程序会被阻塞起来，直到传输完成。

可以使用运行时API函数cuda::cudaGetErrorString获取返回的错误信息，除去核函数之外，其他的CUDA调用都会返回错误的枚举类型cudaError_t。

cudaError_t err = cudaSuccess;
string message = cudaGetErrorString(err);
1
2

内存层次结构

CUDA编程模型给出了内存层次结构，GPU设备有用于不同用途的存储类型。最主要的两类内存是全局内存和共享内存，全局内存对所有线程可见，共享内存存在每个线程块内，线程块内的每个线程使用同一块共享内存。

下面的program来自nvidia的example。

https://github.com/ZouJiu1/cuda-samples/blob/master/Samples/0_Introduction/vectorAdd ，数据传输是在主机内存和GPU全局内存之间进行的。

编译方式就是：nvcc -Xcompiler -std=c99

vectorAdd.cu -o sum，运行是：./sum

编译器的选项可以从这个网页查询的：

1. Introduction — cuda-compiler-driver-nvcc 12.2 documentation (nvidia.com)

主机调用核函数以后，控制权就交给了CPU，CPU可以执行其他函数，所以核函数是异步的。cudaMemcpy函数将结果从设备复制到主机，这个函数是同步的，会阻塞起来直到数据传输完成。

/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions
 * are met:
 *  * Redistributions of source code must retain the above copyright
 *    notice, this list of conditions and the following disclaimer.
 *  * Redistributions in binary form must reproduce the above copyright
 *    notice, this list of conditions and the following disclaimer in the
 *    documentation and/or other materials provided with the distribution.
 *  * Neither the name of NVIDIA CORPORATION nor the names of its
 *    contributors may be used to endorse or promote products derived
 *    from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
 * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
 * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
 * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
 * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
 * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
 * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */

/**
 * Vector addition: C = A + B.
 *
 * This sample is a very basic sample that implements element by element
 * vector addition. It is the same as the sample illustrating Chapter 2
 * of the programming guide with some additions like error checking.
 */

#include <stdio.h>

// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>

#include <helper_cuda.h>
/**
 * CUDA Kernel Device code
 *
 * Computes the vector addition of A and B into C. The 3 vectors have the same
 * number of elements numElements.
 */
__global__ void vectorAdd(const float *A, const float *B, float *C,
                          int numElements) {
  // blockDim.x 一个线程块block内的线程数量
  // blockIdx.x 网格grid内该线程块block的标号 index
  // threadIdx.x 线程块block内该线程的标号 index
  int i = blockDim.x * blockIdx.x + threadIdx.x;

  if (i < numElements) {  // 防止线程越界的
    C[i] = A[i] + B[i] + 0.0f;   // 累加
  }
}

/**
 * Host main routine
 */
int main(void) {
  // Error code to check return values for CUDA calls
  cudaError_t err = cudaSuccess;

  // Print the vector length to be used, and compute its size
  int numElements = 50000;
  size_t size = numElements * sizeof(float);
  printf("[Vector addition of %d elements]\n", numElements);

  // Allocate the host input vector A
  float *h_A = (float *)malloc(size);  //分配主机的内存

  // Allocate the host input vector B
  float *h_B = (float *)malloc(size);

  // Allocate the host output vector C
  float *h_C = (float *)malloc(size);

  // Verify that allocations succeeded
  if (h_A == NULL || h_B == NULL || h_C == NULL) {   // 分配的结果
    fprintf(stderr, "Failed to allocate host vectors!\n");
    exit(EXIT_FAILURE);
  }

  // Initialize the host input vectors
  for (int i = 0; i < numElements; ++i) {    // 给定随机数值的
    h_A[i] = rand() / (float)RAND_MAX;
    h_B[i] = rand() / (float)RAND_MAX;
  }

  // Allocate the device input vector A
  float *d_A = NULL; 
  err = cudaMalloc((void **)&d_A, size);   // 分配设备内存

  if (err != cudaSuccess) {    // 检查结果的
    fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Allocate the device input vector B
  float *d_B = NULL;      
  err = cudaMalloc((void **)&d_B, size);    // 分配设备内存

  if (err != cudaSuccess) {     // 检查结果的
    fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Allocate the device output vector C
  float *d_C = NULL;
  err = cudaMalloc((void **)&d_C, size);    // 分配设备内存

  if (err != cudaSuccess) {     // 检查结果的
    fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Copy the host input vectors A and B in host memory to the device input
  // vectors in
  // device memory
  printf("Copy input data from the host memory to the CUDA device\n");
  err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); // 同步函数的，从主机端复制到设备端

  if (err != cudaSuccess) {     // 检查结果的
    fprintf(stderr,
            "Failed to copy vector A from host to device (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

  if (err != cudaSuccess) {
    fprintf(stderr,
            "Failed to copy vector B from host to device (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Launch the Vector Add CUDA Kernel
  int threadsPerBlock = 256; // 每个block的线程数量
  int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; // 每个grid的block数量
  printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,
         threadsPerBlock);
  // 三重括号内是执行配置的
  vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
  err = cudaGetLastError();

  if (err != cudaSuccess) {
    fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Copy the device result vector in device memory to the host result vector
  // in host memory.
  printf("Copy output data from the CUDA device to the host memory\n");
  err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

  if (err != cudaSuccess) {
    fprintf(stderr,
            "Failed to copy vector C from device to host (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Verify that the result vector is correct
  for (int i = 0; i < numElements; ++i) {
    if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {
      fprintf(stderr, "Result verification failed at element %d!\n", i);
      exit(EXIT_FAILURE);
    }
  }

  printf("Test PASSED\n");

  // Free device global memory
  err = cudaFree(d_A);   // 释放分配的显存

  if (err != cudaSuccess) {
    fprintf(stderr, "Failed to free device vector A (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  err = cudaFree(d_B);   // 释放分配的显存

  if (err != cudaSuccess) {
    fprintf(stderr, "Failed to free device vector B (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  err = cudaFree(d_C);   // 释放分配的显存

  if (err != cudaSuccess) {
    fprintf(stderr, "Failed to free device vector C (error code %s)!\n",
            cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // Free host memory
  free(h_A);            // 释放主机分配的内存
  free(h_B);
  free(h_C);

  printf("Done\n");
  return 0;
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215

2.1.3 线程管理

在host启动device以后，device会产生很多线程来运行kernel函数，CUDA有着明确的抽象层次概念，也就是两部分，线程块block和网格grid组成的。grid包含了很多的block，每个block包含了很多thread，grid内所有threads使用全局内存，每一个线程块内部都有一块共享内存，共享内存由线程块内的线程访问，其他线程块不能访问，不过GPU9.0以后，多了线程块簇，共享内存组成了分布式共享内存，此时不同线程块可以访问其他的共享内存。

https://zhuanlan.zhihu.com/p/659588909