跳到主要内容

3. MUSA 软件架构

MUSA(Meta-computing Unified System Architecture)，中文名为元计算统一系统架构，是显卡厂商摩尔线程(Moore Threads)推出的运算平台。元计算一词涵盖人工智能计算，图形计算，物理仿真计算等重要领域。MUSA 是一种通用并行计算架构，该架构使 GPU 能够解决复杂的计算问题。 它包含了 MUSA 指令集架构（ISA）以及 GPU 内部的并行计算引擎。

MUSA 带有一个软件环境，允许开发人员使用 C/C++ 语言为 MUSA 架构编写程序，编写出的程序可以在支持 MUSA 的 GPU 上以超高性能运行。同时 MUSA 软件栈还提供了各个领域的加速库，供开发者直接使用，便于快速适配各种高性能计算场景。具体如图 1 所示。

img1

图1 MUSA 软件栈

3.1. MUSA 开发套件​

MUSA 开发套件（MUSA Toolkit）是一套用于开发、优化及部署高性能 GPU 加速应用程序的工具包，它包含了 GPU 加速库、调试优化工具、C/C++ 编译器以及用于部署应用程序的运行时库。

借助 MUSA Toolkit，开发者可以将自己的应用程序快速部署到嵌入式系统、桌面工作站、企业数据中心以及超级计算机上。并且 MUSA Toolkit 具有跨多 GPU 的分布式计算能力，支持将应用程序从单个 GPU 工作站扩展到具有数千个 GPU 的云平台上。

3.2. MUSA 驱动与运行时库​

3.2.1. 初始化函数​

MUSA 运行时库（MUSA Runtime）没有特定的初始化函数（Initialization），在程序第一次调用 Runtime 库函数时会自动完成初始化。因此在记录 Runtime 函数调用时间和解释程序中第一个 Runtime 调用返回的错误代码时，需要将初始化考虑在内。

在初始化期间，Runtime 将会为系统中每一个设备创建一个 MUSA context，这个 context 是设备的基础 context(primary context)，它被程序中所有的主机线程所共享。创建过程在后台运行，并且 Runtime 将隐藏 primary context，使之对 Runtime API 这一层的程序员不可见。

当一个主机线程调用 musaDeviceReset() 函数时，它将会销毁线程当前控制设备的 primary context。即当线程下一次调用 Runtime 函数时将会重启初始化，一个新的 MUSA primary context 将被创建出来。

3.2.2. 设备内存​

MUSA 编程模型假定系统由主机侧(host) 和设备侧(device) 构成，它们分别具有独立的内存空间。Runtime 负责设备内存（Device Memory）的分配，回收，拷贝以及在主机和设备之间传输数据的工作。

设备内存可以有两种分配方式：线性内存 (Linear Memory) 或者 MUSA 数组 (MUSA Arrays)。

MUSA 数组是一块不透明的内存空间，它主要用于纹理存取优化。

线性内存空间与平时我们访问的内存类似，对于计算能力 1.x 的设备来说，它存在于一个 32 位的地址空间。对于更高计算能力的设备而言，它存在于一个 40 位的地址空间中。因此，单独分配的实体可以使用指针来相互应用。

通常，我们使用 musaMalloc() 函数分配线性内存空间，释放线性内存空间时调用 musaFree() 函数，musaMemcpy() 函数用于主机和设备之间传输数据。以下是 musa Vector Add 代码示例的一些片段：

  #include <iostream>

  __global__ void axpy(float *x, float *y, float a) {
    y[threadIdx.x] = a * x[threadIdx.x];
  }

  int main(int argc, char *argv[]) {
    const int kDataLen = 4;

    float a = 2.0f;
    float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
    float host_y[kDataLen];

    // Copy input data to device.
    float *device_x;
    float *device_y;
    musaMalloc(&device_x, kDataLen * sizeof(float));
    musaMalloc(&device_y, kDataLen * sizeof(float));
    musaMemcpy(device_x, host_x, kDataLen * sizeof(float),
              musaMemcpyHostToDevice);

    // Launch the kernel.
    axpy<<<1, kDataLen>>>(device_x, device_y, a);

    // Copy output data to host.
    musaDeviceSynchronize();
    musaMemcpy(host_y, device_y, kDataLen * sizeof(float),
              musaMemcpyDeviceToHost);

    // Print the results.
    for (int i = 0; i < kDataLen; ++i) {
      std::cout << "y[" << i << "] = " << host_y[i] << "\n";
    }

    musaFree(device_x);
    musaFree(device_y);
    return 0;
  }

上述代码展示了设备内存的分配，传输以及回收过程。

除了上面展示的方法，我们还可以使用 musaMallocPitch() 和 musaMalloc3D() 函数来分配线性内存。这些函数能够确保分配的内存满足设备内存访问的对齐要求，对行地址的访问以及多维数组间的数据传输提供高性能保证，因此非常适用于对二维和三维数组内存空间的分配。下面的代码片段展示了分配和使用尺寸为 width x height 的二维数组的技术：

  // Host code
  int width = 64, height = 64;
  float* devPtr;
  size_t pitch;
  musaMallocPitch(&devPtr, &pitch,
                  width * sizeof(float), height);
  MyKernel<<<100, 512>>>(devPtr, pitch, width, height);

  // Device code
  __global__ void MyKernel(float* devPtr,
                          size_t pitch, int width, int height)
  {
      for (int r = 0; r < height; ++r) {
          float* row = (float*)((char*)devPtr + r * pitch);
          for (int c = 0; c < width; ++c) {
              float element = row[c];
          }
      }
  }

更多详细的内容请参考 MUSA Runtime API Reference。
下面的代码示例展示了多种使用 Runtime API 访问全局变量的技术：

  __constant__ float constData[256];
  float data[256];
  musaMemcpyToSymbol(constData, data, sizeof(data));
  musaMemcpyFromSymbol(data, constData, sizeof(data));

  __device__ float devData;
  float value = 3.14f;
  musaMemcpyToSymbol(devData, &value, sizeof(float));

  __device__ float* devPointer;
  float* ptr;
  musaMalloc(&ptr, 256 * sizeof(float));
  musaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

使用 musaGetSymbolAddress() 函数可以获得被声明存储在全局内存中的变量地址。如果需要获取分配内存的大小，可以使用musaGetSymbolSize() 函数。

3.2.3. Device Memory L2 Access Management​

当内核反复访问全局内存中的数据区域时，可以视为 Persisting accesses; 反过来，如果数据预计只被访问一次，则为 Streaming accesses。为了提高缓存利用率并减少 DRAM 回写，我们可以使数据在 L2 中驻留更长时间，从而提供更高的带宽和对全局内存访问时更低的延迟。

对全局内存的数据访问可分为 3 类：

  • 流式访问 (Streaming accesses)：数据仅访问一次，因此 cache lines 不太可能保留在 L2 中。
  • 持久访问 (Persisting accesses)：重复访问数据，因此 cache lines 更有可能持久保持在 L2 中。
  • 正常访问 (Normal access)：正常访问。

根据访问的类别，我们可以获得缓存替换策略的提示，以实现更好的缓存使用。例如，我们可以用 Persisting accesses 替换 Streaming accesses，因为我们知道后者很可能会在不久的将来缓存命中。

许多 AI 或通用计算工作负载具有可以频繁访问的共享数据，并且本地内存不够大，无法容纳它们。例如，50 层 ResNet 网络有 2600 万个权重参数，并在正向传递中计算 1600 万次激活。内存访问量相当大，因此 L2 缓存效率非常低。如果我们能够正确控制 L2 驻留，可以减少全局内存访问时的延迟，带来一定的性能提升。MUSA 提供了一套 L2 set-aside APIs 来进行驻留和替换控制。

3.2.3.1. L2 cache Set-Aside for Persisting Accesses​

可以留出一部分L2高速缓存以用于持久存储对全局存储器的数据访问，而且相比于非持久数据访问，持久性数据访问在使用L2高速缓冲时具有更高的优先权。为持久化数据保留的 LLC 大小可以在硬件限制基础上进行调整。大小需要在程序执行开始前通过musa api进行设置。

  musaGetDeviceProperties(&prop, device_id);
  size_t size = min(int(prop.l2CacheSize * 0.75), prop.persistingL2CacheMaxSize);
  musaDeviceSetLimit(musaLimitPersistingL2CacheSize, size); /* set-aside 3/4 of L2 cache for persisting accesses or the max allowed*/

3.2.3.2. L2 Policy for Persisting Accesses​

访问策略窗口（access policy window）指定了全局存储器的连续区域以及用于访问该区域的L2高速缓存中的持久性属性。开发者需要指定access policy window 相关字段的属性：

  musaStreamAttrValue stream_attribute;                                         // Stream level attributes data structure
  stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast<void*>(ptr); // Global Memory data pointer
  stream_attribute.accessPolicyWindow.num_bytes = num_bytes;                    // Number of bytes for persistence access.
                                                                                // (Must be less than musaDeviceProp::accessPolicyMaxWindowSize)
  stream_attribute.accessPolicyWindow.hitRatio  = 0.6;                          // Hint for cache hit ratio
  stream_attribute.accessPolicyWindow.hitProp   = musaAccessPropertyPersisting; // Type of access property on cache hit
  stream_attribute.accessPolicyWindow.missProp  = musaAccessPropertyStreaming;  // Type of access property on cache miss.//Set the attributes to a MUSA stream of type musaStream_t
  musaStreamSetAttribute(stream, musaStreamAttributeAccessPolicyWindow, &stream_attribute);

一般在设置访问策略窗口时，会在全局内存中指定一块区域，这里假设A，然后使用 hitRatio 指定这块区域中的持久性访问的比例作为 hitProp 区域，但是 L2 set-aside cache size 是有限的，最终作为 hitProp 区域的是 min(A*hitPrtio, L2 set-aside cache size)。

计算 access policy window 起止地址的方法：

  if hit hint Ratio is 1.0:
      start = the base VA from base_ptr
      end = start + num_bytes
  else:
      start = the base VA from base_ptr
      end = start + num_bytes * Ratio
      // Line 5 is a basic implementation.
      // A better solution: start = a random value between
      // base_ptr and (base_ptr + num_bytes - num_bytes * Ratio)

此外如果两个 kernel 共同访问 L2 set-aside cache，如果两者设置的 size 相交不超过 L2 set-aside cache size，那么没有问题，如果超过，那么会出现冲突。

3.2.3.3. L2 Access Properties​

不同的全局内存数据访问定义了三种类型的访问属性：
（1）musaAccessPropertyStreaming：使用流属性进行的内存访问不太可能保留在 L2 高速缓存中，因为这些访问被优先驱逐。
（2）musaAccessPropertyPersisting：由于具有持久属性而发生的内存访问更可能保留在 L2 高速缓存中，因为这些访问优先保留在 L2 高速缓存的预留部分中。
（3）musaAccessPropertyNormal：此访问属性将以前应用的持久访问属性强制重置为正常状态。 先前的MUSA内核具有持久属性的内存访问可能在其预期用途后很长时间被保留在 L2 缓存中。 这种使用后的持久性减少了不使用持久性的后续内核可使用的 L2 缓存数量。 使用 musaAccessPropertyNormal 属性重置访问属性窗口会删除先前访问的持久（优先保留）状态，就好像先前访问没有访问属性一样。

3.2.3.4. L2 Persistence Example​

下面是一个例子：

  musaStream_t stream;
  musaStreamCreate(&stream);                                                                  // Create MUSA stream

  musaDeviceProp prop;                                                                        // MUSA device properties variable
  musaGetDeviceProperties( &prop, device_id);                                                 // Query GPU properties
  size_t size = min( int(prop.l2CacheSize * 0.75) , prop.persistingL2CacheMaxSize );
  musaDeviceSetLimit( musaLimitPersistingL2CacheSize, size);                                  // set-aside 3/4 of L2 cache for persisting accesses or the max allowed

  size_t window_size = min(prop.accessPolicyMaxWindowSize, num_bytes);                        // Select minimum of user defined num_bytes and max window size.

  musaStreamAttrValue stream_attribute;                                                       // Stream level attributes data structure
  stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast<void*>(data1);               // Global Memory data pointer
  stream_attribute.accessPolicyWindow.num_bytes = window_size;                                // Number of bytes for persistence access
  stream_attribute.accessPolicyWindow.hitRatio  = 0.6;                                        // Hint for cache hit ratio
  stream_attribute.accessPolicyWindow.hitProp   = musaAccessPropertyPersisting;               // Persistence Property
  stream_attribute.accessPolicyWindow.missProp  = musaAccessPropertyStreaming;                // Type of access property on cache miss

  musaStreamSetAttribute(stream, musaStreamAttributeAccessPolicyWindow, &stream_attribute);   // Set the attributes to a MUSA Stream

  for(int i = 0; i < 10; i++) {
      musa_kernelA<<<grid_size,block_size,0,stream>>>(data1);                                 // This data1 is used by a kernel multiple times
  }                                                                                           // [data1 + num_bytes) benefits from L2 persistence
  musa_kernelB<<<grid_size,block_size,0,stream>>>(data1);                                     // A different kernel in the same stream can also benefit
                                                                                              // from the persistence of data1

  stream_attribute.accessPolicyWindow.num_bytes = 0;                                          // Setting the window size to 0 disable it
  musaStreamSetAttribute(stream, musaStreamAttributeAccessPolicyWindow, &stream_attribute);   // Overwrite the access policy attribute to a MUSA Stream
  musaCtxResetPersistingL2Cache();                                                            // Remove any persistent lines in L2

  musa_kernelC<<<grid_size,block_size,0,stream>>>(data2);

3.2.3.5. Reset L2 Access to Normal​

来自先前 MUSA 内核的持久化 L2 高速缓存行可能在使用后很长一段时间内一直保持在 L2 中。因此，将 L2 高速缓存重置为正常对于使用具有正常优先级的 L2 高速缓存的流式传输或正常存储器访问是重要的。有三种方法可以将持久访问重置为正常状态。

  1. 使用访问属性 musaAccessPropertyNormal 重置以前的持久内存区域。
  2. 通过调用 musaCtxResetPersistingL2Cache() 将所有持久化的二级缓存行重置为正常。
  3. 最终，没有使用的的 cache line 将自动重置为正常，强烈建议不要依赖自动复位，因为自动复位所需的时间长度不确定。

3.2.3.6. Manage Utilization of L2 set-aside cache​

在不同流中的核函数有可能同时使用访问策略窗口（access policy window），但是 L2 set-aside cache 是全局的，被共同使用，当使用的资源超过 L2 set-aside cache size 是，会出现性能下降，所以在使用之前要考虑如下事情：

  1. L2 预留缓存的大小。
  2. 可以同时执行的 MUSA 内核。
  3. 可能同时执行的所有 MUSA 内核的访问策略窗口。
  4. 需要何时以及如何重置 L2 才能允许普通访问或流访问使用具有相同优先级的先前预留的 L2 缓存。

3.2.3.7. Query L2 cache Properties​

与L2缓存相关的属性是 musaDeviceProp 结构的一部分，可以使用 MUSA 运行时 API musaGetDeviceProperties 查询。

MUSA设备属性包括：

  1. l2CacheSize：GPU 上的可用 L2 缓存大小。
  2. persistenceingL2CacheMaxSize：可以为持久内存访问而保留的 L2 高速缓存的最大数量。
  3. accessPolicyMaxWindowSize：访问策略窗口的最大尺寸。

3.2.3.8. Control L2 Cache Set-Aside Size for Persisting Memory Access​

使用 MUSA 运行时 musaDeviceGetLimit 查询用于持久存储访问的 L2 预留缓存大小，并使用 MUSA 运行时 musaDeviceSetLimit 作为 musaLimit 进行设置。设置此限制的最大值是 musaDeviceProp ::persistingL2CacheMaxSize。

enum musaLimit {
    /* other fields not shown */
    musaLimitPersistingL2CacheSize
};

3.2.4. 共享内存​

MUSA 使用 __shared__ 空间标识符来分配共享内存（Shared Memory）。

共享内存的速度大多数时候比全局内存快得多，它可以用作临时内存（或缓存）来减少来自 MUSA 对全局内存访问。以下的矩阵乘法示例就验证了这一点。

以下代码示例是矩阵乘法的直接实现，不利用共享内存。每个线程读取 A 的一行和 B 的一列，并计算 C 的相应元素，如图所示。从全局内存中读取 A B.width 次，读取 B A.height 次。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
    int width;
    int height;
    float* elements;
} Matrix;

// Thread block size
#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
    // Load A and B to device memory
    Matrix d_A;
    d_A.width = A.width; d_A.height = A.height;
    size_t size = A.width * A.height * sizeof(float);
    musaMalloc(&d_A.elements, size);
    musaMemcpy(d_A.elements, A.elements, size,
               musaMemcpyHostToDevice);
    Matrix d_B;
    d_B.width = B.width; d_B.height = B.height;
    size = B.width * B.height * sizeof(float);
    musaMalloc(&d_B.elements, size);
    musaMemcpy(d_B.elements, B.elements, size,
               musaMemcpyHostToDevice);

    // Allocate C in device memory
    Matrix d_C;
    d_C.width = C.width; d_C.height = C.height;
    size = C.width * C.height * sizeof(float);
    musaMalloc(&d_C.elements, size);

    // Invoke kernel
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
    MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);

    // Read C from device memory
    musaMemcpy(C.elements, d_C.elements, size,
               musaMemcpyDeviceToHost);

    // Free device memory
    musaFree(d_A.elements);
    musaFree(d_B.elements);
    musaFree(d_C.elements);
}

// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
    // Each thread computes one element of C
    // by accumulating results into Cvalue
    float Cvalue = 0;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    for (int e = 0; e < A.width; ++e)
        Cvalue += A.elements[row * A.width + e]
                * B.elements[e * B.width + col];
    C.elements[row * C.width + col] = Cvalue;
}

3.2.4-1

以下代码示例是利用共享内存的矩阵乘法实现。在此实现中，每个线程块负责计算 C 的一个正方形子矩阵 C-sub，而块内的每个线程负责计算 C-sub 的一个元素。如图所示，C-sub 等于两个矩形矩阵的乘积：A 的维度为（A.width，block_size）的子矩阵，其行索引与 C-sub相同；B 的维度为（block_size，A.width）的子矩阵，其列索引与 C-sub 相同。为了适应设备的资源，这两个矩形矩阵被划分为大小为 block_size 的正方形矩阵，并且 C-sub 是这些正方形矩阵的乘积之和。这些乘积通过使用一个线程从全局内存加载两个对应的正方形矩阵到共享内存中来执行，其中一个线程加载每个矩阵的一个元素，然后每个线程计算一个元素的乘积。每个线程将这些乘积的结果累加到一个寄存器中，计算完成后，将结果写入全局内存。 通过这种分块计算的方式，我们利用快速的共享内存节省了大量的全局内存带宽，因为 A 仅从全局内存中被读取（B.width / block_size）次，而 B 则被读取（A.height / block_size）次。 在下面的代码示例中，Matrix 类型增加了一个 stride 字段，以便可以使用相同类型来有效地表示子矩阵。__device__ 函数用于获取和设置元素，并从矩阵构建子矩阵。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {
    int width;
    int height;
    int stride;
    float* elements;
} Matrix;
// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{
    return A.elements[row * A.stride + col];
}
// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col,
                           float value)
{
    A.elements[row * A.stride + col] = value;
}
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A
 __device__ Matrix GetSubMatrix(Matrix A, int row, int col)
{
    Matrix Asub;
    Asub.width    = BLOCK_SIZE;
    Asub.height   = BLOCK_SIZE;
    Asub.stride   = A.stride;
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row
                                         + BLOCK_SIZE * col];
    return Asub;
}
// Thread block size
#define BLOCK_SIZE 16
// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
    // Load A and B to device memory
    Matrix d_A;
    d_A.width = d_A.stride = A.width; d_A.height = A.height;
    size_t size = A.width * A.height * sizeof(float);
    musaMalloc(&d_A.elements, size);
    musaMemcpy(d_A.elements, A.elements, size,
               musaMemcpyHostToDevice);
    Matrix d_B;
    d_B.width = d_B.stride = B.width; d_B.height = B.height;
    size = B.width * B.height * sizeof(float);
    musaMalloc(&d_B.elements, size);
    musaMemcpy(d_B.elements, B.elements, size,
    musaMemcpyHostToDevice);
    // Allocate C in device memory
    Matrix d_C;
    d_C.width = d_C.stride = C.width; d_C.height = C.height;
    size = C.width * C.height * sizeof(float);
    musaMalloc(&d_C.elements, size);
    // Invoke kernel
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
    MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
    // Read C from device memory
    musaMemcpy(C.elements, d_C.elements, size,
               musaMemcpyDeviceToHost);
    // Free device memory
    musaFree(d_A.elements);
    musaFree(d_B.elements);
    musaFree(d_C.elements);
}
// Matrix multiplication kernel called by MatMul()
 __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
    // Block row and column
    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;
    // Each thread block computes one sub-matrix Csub of C
    Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
    // Each thread computes one element of Csub
    // by accumulating results into Cvalue
    float Cvalue = 0;
    // Thread row and column within Csub
    int row = threadIdx.y;
    int col = threadIdx.x;
    // Loop over all the sub-matrices of A and B that are
    // required to compute Csub
    // Multiply each pair of sub-matrices together
    // and accumulate the results
    for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
        // Get sub-matrix Asub of A
        Matrix Asub = GetSubMatrix(A, blockRow, m);
        // Get sub-matrix Bsub of B
        Matrix Bsub = GetSubMatrix(B, m, blockCol);
        // Shared memory used to store Asub and Bsub respectively
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
        // Load Asub and Bsub from device memory to shared memory
        // Each thread loads one element of each sub-matrix
        As[row][col] = GetElement(Asub, row, col);
        Bs[row][col] = GetElement(Bsub, row, col);
        // Synchronize to make sure the sub-matrices are loaded
        // before starting the computation
        __syncthreads();
        // Multiply Asub and Bsub together
        for (int e = 0; e < BLOCK_SIZE; ++e)
            Cvalue += As[row][e] * Bs[e][col];
        // Synchronize to make sure that the preceding
        // computation is done before loading two new
        // sub-matrices of A and B in the next iteration
        __syncthreads();
    }
    // Write Csub to device memory
    // Each thread writes one element
    SetElement(Csub, row, col, Cvalue);
}

3.2.4-2

3.2.5. 分布式共享内存​

分布式共享内存（Distributed Shared Memory）为线程块集群中的线程提供了访问集群中其他线程块的共享内存的能力，其对应的地址空间称之为分布式共享内存地址空间。属于线程块集群的线程可以在分布式地址空间中读取、写入或执行原子操作，无论该地址是本地线程块还是远程线程块。不管内核是否使用分布式共享内存，静态或动态的共享内存大小规格都是每个线程块。分布式共享内存的大小是每个集群的线程块数乘以每个线程块的共享内存大小。

访问分布式共享内存中的数据需要所有线程块都存在。用户可以使用 Cluster Group API 中的 cluster.sync() 来保证所有线程块都已开始执行。另外，用户还需要确保在线程块退出之前完成所有的分布式共享内存操作。

MUSA 提供了一种访问分布式共享内存的机制。我们可以从简单的直方图计算，以及如何使用线程块集群在 GPU 上对其进行优化入手。计算直方图的标准方法是在每个线程块的共享内存中进行计算，然后执行全局内存原子操作。这种方法的局制是共享内存容量。一旦共享内存无法容纳直方图，用户就需要直接计算直方图，从而执行全局内存的原子操作。对于分布式共享内存，MUSA 提供了一个中间步骤，根据直方图元素量的大小，可以直接在共享内存、分布式共享内存或全局内存中计算直方图。

以下代码段演示了如何根据直方图元素的数量在共享内存或分布式共享内存中计算直方图。

#include <cooperative_groups.h>

// Distributed Shared memory histogram kernel
__global__ void clusterHist_kernel(int *bins, const int nbins, const int bins_per_block, const int *__restrict__ input,
                                   size_t array_size)
{
  extern __shared__ int smem[];
  namespace cg = cooperative_groups;
  int tid = cg::this_grid().thread_rank();

  // Cluster initialization, size and calculating local bin offsets.
  cg::cluster_group cluster = cg::this_cluster();
  unsigned int clusterBlockRank = cluster.block_rank();
  int cluster_size = cluster.dim_blocks().x;

  for (int i = threadIdx.x; i < bins_per_block; i += blockDim.x)
  {
    smem[i] = 0; //Initialize shared memory histogram to zeros
  }

  // cluster synchronization ensures that shared memory is initialized to zero in
  // all thread blocks in the cluster. It also ensures that all thread blocks
  // have started executing and they exist concurrently.
  cluster.sync();

  for (int i = tid; i < array_size; i += blockDim.x * gridDim.x)
  {
    int ldata = input[i];

    //Find the right histogram bin.
    int binid = ldata;
    if (ldata < 0)
      binid = 0;
    else if (ldata >= nbins)
      binid = nbins - 1;

    //Find destination block rank and offset for computing
    //distributed shared memory histogram
    int dst_block_rank = (int)(binid / bins_per_block);
    int dst_offset = binid % bins_per_block;

    //Pointer to target block shared memory
    int *dst_smem = cluster.map_shared_rank(smem, dst_block_rank);

    //Perform atomic update of the histogram bin
    atomicAdd(dst_smem + dst_offset, 1);
  }

  // cluster synchronization is required to ensure all distributed shared
  // memory operations are completed and no thread block exits while
  // other thread blocks are still accessing distributed shared memory
  cluster.sync();

  // Perform global memory histogram, using the local distributed memory histogram
  int *lbins = bins + cluster.block_rank() * bins_per_block;
  for (int i = threadIdx.x; i < bins_per_block; i += blockDim.x)
  {
    atomicAdd(&lbins[i], smem[i]);
  }
}

上面的程序可以在运行时以集群大小启动，具体取决于分布式共享内存的数量。如果直方图比较小，适合一块共享内存，用户则可以启动集群大小为 1 的内核。

以下是根据共享内存要求动态的启动集群内核的代码示例。

// Launch via extensible launch
{
  musaLaunchConfig_t config = {0};
  config.gridDim = array_size / threads_per_block;
  config.blockDim = threads_per_block;

  // cluster_size depends on the histogram size.
  // ( cluster_size == 1 ) implies no distributed shared memory, just thread block local shared memory
  int cluster_size = 2; // size 2 is an example here
  int nbins_per_block = nbins / cluster_size;

  //dynamic shared memory size is per block.
  //Distributed shared memory size =  cluster_size * nbins_per_block * sizeof(int)
  config.dynamicSmemBytes = nbins_per_block * sizeof(int);

  musa_CHECK(::musaFuncSetAttribute((void *)clusterHist_kernel, musaFuncAttributeMaxDynamicSharedMemorySize, config.dynamicSmemBytes));

  musaLaunchAttribute attribute[1];
  attribute[0].id = musaLaunchAttributeClusterDimension;
  attribute[0].val.clusterDim.x = cluster_size;
  attribute[0].val.clusterDim.y = 1;
  attribute[0].val.clusterDim.z = 1;

  config.numAttrs = 1;
  config.attrs = attribute;

  musaLaunchKernelEx(&config, clusterHist_kernel, bins, nbins, nbins_per_block, input, array_size);
}

3.2.6. 锁页内存​

MUSA-runtime 提供了一些接口，让用户实现对 host memory 的 page-locked（pinned mem）。

  • musaHostAlloc() 和 musaFreeHost() 申请和释放锁页内存；
  • musaHostRegister() 指定范围的 host memory 进行 page-lock。

使用 page-locked 内存有以下好处：

  • 锁页内存和 GPU 设备之间的拷贝可以与某些内核同时执行。
  • 锁页内存可以直接映射到 GPU 设备的内存地址空间上，无需再次执行 host 到 device 的拷贝，就能实现让设备读取到数据。
  • 锁页内存和 GPU 之间的读写带宽会更高，如果分配的时候，采用的 write-combining，还可以获得更大的带宽。

> 注意：
>
> 为了获取最大的拷贝带宽，MUSA-Runtime 会为 musaHostAlloc() 分配对齐到 128 字节的主机内存。另一方面，musaHostRegister 会 page-lock 指定范围的 host memory。如需要获得最大带宽，请尝试将 host memory 的地址对齐到 128 字节。

Linux 系统上，通过 aligned_alloc(需要支持 c11 标准)或者 posix_memalign 来申请地址对齐的 host memory，Windows 上则使用 _aligned_malloc 命令。以下为 Linux 申请地址对齐到 128 字节的 host memory 的示例：

#include <stdio.h>
#include <stdlib.h>

/* Check if size is a multiple of alignment and adjust size */
#define ALIGNMENT_CHECK_AND_ADJUST(size, alignment) \
    (((size) + (alignment) - 1) & ~((alignment) - 1))

 int main() {

    size_t alignment = 128;
    size_t size = 129;
    
    void *hostptr = aligned_alloc(alignment, ALIGNMENT_CHECK_AND_ADJUST(size, alignment));
    if (hostptr == NULL) {
        printf("Failed to allocate alignment memory.\n");
        return 1;
    }

    /* Or implement it like this
    if (posix_memalign(&hostptr, alignment, ALIGNMENT_CHECK_AND_ADJUST(size, alignment)) != 0) {
        printf("Failed to allocate alignment memory.\n");
        return 1;
    }
    */

    free(hostptr);

    return 0;

3.2.6.1. Portable memory​

在多卡环境中，锁页内存创建后，所有 GPU 设备都可以使用。为了保证所有 GPU 设备都能受益，需要将标志 musaHostAllocPortable 传递给 musaHostAlloc() 来分配，或者通过将标志 musaHostRegisterPortable 传递给 musaHostRegister() 进行 page-locked。

3.2.6.2. Write-Combining memory​

默认情况下，锁页内存可以分配为 cacheable。您可以传参分配为 write-combining，无需将标志 musaHostAllocWriteCombined 传递给 musaHostAlloc() 来实现。使用 write-combining，可以减少对 L1/L2 cache 的使用。此外，通过 PCI Express 总线传输期间不会侦听 write-combining 的内存，传输速度将提高 40%。

从主机读取 write-combining 内存非常慢，因此 write-combining 内存一般用于主机只写入的内存。

建议避免在 WC 内存上使用 CPU 原子指令，因为并非所有 CPU 都支持该功能。

3.2.6.3. Mapped memory​

通过将标志 musaHostAllocMapped 传递给 musaHostAlloc() ，或者将标志 musaHostRegisterMapped 传递给 musaHostRegister()，都可以将锁页内存映射到设备的地址空间。 因此，这样的内存块通常有两个地址：一个在主机内存中，由 musaHostAlloc() 或 malloc() 返回；另一个在设备内存中，可以使用 musaHostGetDevicePointer() 检索，然后用于从 kernel 访问块。唯一的例外是使用 musaHostAlloc() 分配的指针，以及统一虚拟地址空间中提到的主机和设备使用统一地址空间。

虽然 kernel 访问主机内存不能提供与设备内存相同的带宽，但还是存在一些优势：

  • 无需在 GPU 设备内存中进行分块 ，且不需要执行该块与主机内存的数据拷贝；数据传输根据 kernel 执行的需要，隐式的执行。
  • 无需使用流来控制数据传输，来达到与 kernel 执行重叠。内核发起的数据传输，自动与 kernel 执行重叠。

然而，由于主机和设备共享映射页锁定内存，因此应用程序必须使用流或事件同步内存访问，以避免任何潜在的先写后读（read-after-write）、先读后写（write-after-read）或先写后写（write-after-write）危险。

为了能够检索指向任何映射的页面锁定内存的设备指针，必须在执行任何其他 musa 调用之前通过使用 musaDeviceMapHost 标志调用 musaSetDeviceFlags() 来启用页面锁定内存映射。 否则，musaHostGetDevicePointer() 将返回一个错误。

如果设备不支持映射的页面锁定主机内存，musaHostGetDevicePointer() 也会返回错误。 应用程序可以通过检查 canMapHostMemory 设备属性来查询此功能，对于支持映射页面锁定主机内存的设备，该属性等于 1。

> 注意：
>
> 从主机或其他设备的角度来看，在映射的页面锁定内存上运行的原子函数不是原子的。
>
> MUSA 运行时要求将从设备发起的对主机内存的 1 字节、2 字节、4 字节和 8 字节常规对齐（naturally aligned ）的加载和存储从主机和其他设备的视角保留为单独访问设备。在某些平台上，内存的原子操作可能会被硬件分解为单独的加载和存储操作。这些组件加载和存储操作对保存常规对齐的访问有相同的要求。 例如，MUSA 运行时不支持 PCI Express 总线拓扑，在这种情况下，PCI Express 桥将 8 字节常规对齐的写入拆分为设备和主机之间的两个 4 字节写入。

3.2.7 Memory Synchronization Domains​

3.2.7.1 Memory Fence Interference​

执行 memory fence 或 memory flush 等操作可能会导致 MUSA 应用程序性能下降，因为上述操作等待的事务比 MUSA 内存一致性模型所需事务更多。

__managed__ int x = 0;
__device__ musa::atomic<int, musa::thread_scope_device> a(0);
__managed__ musa::atomic<int, musa::thread_scope_system> b(0);

Thread 1 (MP)

x = 1;
a = 1;

Thread 2 (MP)

while (a != 1);
assert(x == 1);
b = 1;

Thread 3 (CPU)

while (b != 1);
assert(x == 1);

参考以上示例，MUSA 内存一致性模型保证断言条件为真，因此线程 1 对 x 的写入必须在线程 2 写入 b 之前对线程 3 可见。

在设备范围操作中，释放-获取 a 造成的内存事件序只能使 x 对线程 2 可见，而非线程 3。另一方面，释放-获取 b 提供的系统范围的内存事件序需要确保：线程 2 的 b 写入操作对线程 3 可见，同时其他线程对线程 2 可见的写操作也对线程 3 可见。这称为累积性。由于 GPU 执行时无法知道哪些写操作在源码逻辑上可见，哪些仅在执行时偶然可见，因此，必须对正在进行的内存操作映射一个保守、宽泛的网络。

以上操做也存在弊端，由于 GPU 需要等待源码逻辑存在但非必须的内存操作，所以 fence/flush 操作可能比必要的时间长。

> 注意：
>
> fence 在代码中显式地作为内嵌函数或者原子函数出现，如__threadfence_system()；fence 也可能隐式地在任务执行边界处的synchronizes-with 逻辑中实现，如上述示例。

一个常见的场景：

kernel 在本地 GPU 中执行计算，而与它并行的 kernel（例如来自 MCCL）正与对端 GPU 进行通信。完成后，本地 GPU 的 kernel 将隐式地 flush 写操作以便满足任何 synchronizes-with 关系。这可能会在执行低速 PCIe 写操作的通信 kernel 上造成全部/部分不必要的等待。

3.2.8 异步并发执行​

MUSA 将以下操作作为独立任务，它们可以相互并发地执行：

  • 主机侧计算
  • 设备侧计算
  • 主机到设备的数据传输
  • 设备到主机的数据传输
  • 在给定设备的内存中进行数据传输
  • 设备间的数据传输

这些操作之间的并发级别取决于设备的计算能力和实现的特性，如下所述。

3.2.8.1 主机和设备间的并发执行​

异步的 MUSA API 在完成任务之前，主机线程可以执行其他任务，这可以理解为并行主机执行的特性。使用异步接口，许多设备侧操作可以一起排队，资源可用时将由 MUSA 驱动下发到硬件。这减轻了主机侧线程管理设备的许多成本，用户可以自由地处理其他任务。以下设备操作会异步执行：

  • kernel launches
  • 单个设备内存中的数据传输
  • 设备间的数据传输
  • 使用带 async 后缀的函数执行的数据传输
  • 写入设备内存的函数调用

用户可以通过将 MUSA_LAUNCH_BLOCKING 环境变量设置为 1 来全局禁用异步 kernel launches，以便为 MUSA 应用程序进行调试。

3.2.8.2 Kernel 并发执行​

目前硬件支持在一个设备，更细粒度地，在 1 个 MP 上同时执行多个 kernel。通过在 MUSA 应用上创建流，可以同时启动多个 kernel，以最大化减少调度侧负载。

3.2.8.3 数据传输和 Kernel 执行的重叠​

设备可以在 kernel 执行时同时进行数据传输，应用程序可以通过检查 asyncEngineCount 属性值来查询此特性。异步的数据传输在有一方是非锁页的主机内存时，会强制转换为同步操作。

3.2.8.4 数据传输并发​

设备可以并行执行设备和主机间的数据传输操作。应用程序可以通过检查 asyncEngineCount 属性值来查询此特性。为了实现数据传输的 overlap，涉及的任何主机内存都必须是页锁定的。当前的数据传输 overlap 特性较为复杂，因为 MUSA 使用了不同的引擎来执行传输任务。当前，以下不同流上的传输任务是并行执行的：

  • 主机到设备的数据传输与设备到主机的数据传输

  • 设备间的数据传输与主机到设备的数据传输

3.2.8.5 流​

应用程序通过 MUSA 流来管理上述并发操作。流表征了可能由不同主机线程发出的一系列命令，它们将按顺序执行。另一方面，不同的流可能会按照不同的顺序或者同时执行命令。因此，不同流上的操作无法保证也不应依赖共享数据的正确性。流上发出的命令可能会在满足所有前置依赖关系时才被执行，依赖项可以是同一流上先前的命令，或来自于其他流。MUSA 提供的多个 synchronize 后缀的调用可以保证所有已启动的命令在调用处完成。

3.2.8.5.1 创建和销毁​

创建一个流对象，并将其指定为 kernel launches 以及主机 <-> 设备数据传输接口所使用的流来定义它。以下的实例代码创建了两个流并在锁页内存中分配了一个 float 数组 hostPtr。

musaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    musaStreamCreate(&stream[i]);
float* hostPtr;
musaMallocHost(&hostPtr, 2 * size);

每个流都由以下示例代码定义成一个命令序列，包含了从一次主机到设备的数据传输，一次 kernel launch 以及一次设备到主机的数据传输。

for (int i = 0; i < 2; ++i) {
    musaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, musaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
    musaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, musaMemcpyDeviceToHost, stream[i]);
}

这些流每一个都将输入数组 hostPtr 的一部分复制到设备内存中的 inputDevPtr 数组，然后调用 MyKernel() 在设备上处理 inputDevPtr，最后将结果 outputDevPtr 复制回 hostPtr 相应部分。需要注意的是，为了实现不同流的并行，hostPtr 必须是锁页内存。

调用 musaStreamDestroy() 可以释放流，运行如下命令：

for (int i = 0; i < 2; ++i)
    musaStreamDestroy(stream[i]);

如果调用 musaStreamDestroy() 时设备仍在流中执行任务，则函数将立即返回并得到错误码。反之如果设备完成了流中的所有任务，与流相关的资源将得到释放。

3.2.8.5.2 默认流​

如果未指定任何流的 kernel launches 以及主机 <-> 设备数据传输操作，或者将流参数设置为 nullptr，将被分配到默认流上执行，它们也将顺序执行。

对于使用 default-stream legacy 标志编译或在 include MUSA 头文件之前定义了 MUSA_API_PER_THREAD_DEFAULT_STREAM 宏的代码，默认流是一个常规流，且每个主机线程都有自己的默认流。

3.2.8.5.3 显式流同步​

有多种方法可以显式地将流与其他流同步。

  • musaDeviceSynchronize() 等待所有主机线程的所有流中的所有先前命令完成。
  • musaStreamSynchronize() 将流作为参数，并等待给定流中的所有先前命令完成。它可用于将主机与特定流同步，以允许其他流在设备上继续执行。
  • musaStreamWaitEvent() 将流和事件作为参数（有关事件的描述，请参见事件章节），并使调用 musaStreamWaitEvent() 后添加到给定流中的所有命令延迟它们的执行，直到给定事件完成。
  • musaStreamQuery() 提供了应用程序检查流中所有先前命令是否已完成的方法。
3.2.8.5.4 隐式流同步​

如果主机线程在不同流中的两个命令之间发出以下命令，则它们无法同时执行：

  • 调用写入设备内存的 MUSA 接口
  • 任何任何针对空流的 MUSA 接口

需要检查依赖项的操作包括在同一流中启动的任何其他命令以及对该流上的任何 musaStreamQuery() 的调用。因此，应用程序应遵循以下准则，以提高它们可能的并发 kernel 执行能力：

  • 应该在依赖的操作之前发出所有独立操作
  • 任何类型的同步应尽可能推迟
3.2.8.5.5. 重叠行为​

两个流之间的执行重叠（overlap）量取决于向每个流发出命令的顺序以及设备是否支持数据传输和 kernel 执行的重叠（请参见数据传输和 Kernel 执行的重叠），并发 kernel 执行（请参见 Kernel 并发执行）以及并发数据传输（请参见数据并发传输）。

例如，在不支持并发数据传输的设备上，创建和销毁代码示例中的两个流根本不重叠，因为从主机到设备的数据搬移是在设备到主机的数据搬移发往 stream[0] 之后发往 stream[1] 的，因此只能在发向 stream[0] 的从设备到主机的数据搬移完成后才能启动。如果代码按以下方式重写（假设设备支持数据传输和 kernel 执行的重叠），则发往 stream[1] 的从主机到设备的数据搬移重叠于发往 stream[0] 的 kernel。

for (int i = 0; i < 2; ++i)
    musaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, musaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel<<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    musaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, musaMemcpyDeviceToHost, stream[i]);

在支持并发数据传输的设备上，创建和销毁代码示例中的两个流确实重叠：发往 stream[1] 的从主机到设备的内存复制重叠于发往 stream[0] 的从设备到主机的内存复制，甚至重叠于发往 stream[0] 的 kernel（假设设备支持数据传输和 kernel 执行的重叠）。

3.2.8.5.6. 流优先级​

可以使用 musaStreamCreateWithPriority() 在创建流时指定流的相对优先级。也可以使用 musaDeviceGetStreamPriorityRange() 函数获取有序的允许优先级范围[最高优先级，最低优先级]。在运行时，高优先级流中的待处理工作优先于低优先级流中的待处理工作。

以下代码示例获取当前设备的允许优先级范围，并创建具有最高和最低可用优先级的流。

// 获得当前设备的流优先级范围
int priority_high, priority_low;
musaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// 创建具有最高优先级以及最低优先级的流
musaStream_t st_high, st_low;
musaStreamCreateWithPriority(&st_high, musaStreamNonBlocking, priority_high);
musaStreamCreateWithPriority(&st_low, musaStreamNonBlocking, priority_low);

3.2.8.6. 事件​

MUSA 运行时还提供了一种紧密监视设备进度以及进行准确计时的方法，即允许应用程序在程序的任何时刻异步记录事件（events），并查询这些事件何时完成。当事件之前的所有任务（或可选地，给定流中的所有命令）完成时，事件就完成了。默认流中的事件在所有流中之前的所有任务和命令完成后才会完成。

3.2.8.6.1. 创建和销毁​

以下示例代码创建了两个事件：

musaEvent_t start, stop;
musaEventCreate(&start);
musaEventCreate(&stop);

它们可以通过如下命令释放：

musaEventDestroy(start);
musaEventDestroy(stop);
3.2.8.6.2. 计时（Elapsed time）​

上文代码示例中创建的事件可以用以下方式计时：

musaEventRecord(start, 0);
for (int i = 0; i < 2; ++i) {
    musaMemcpyAsync(inputDev + i * size, inputHost + i * size,
                    size, musaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>
               (outputDev + i * size, inputDev + i * size, size);
    musaMemcpyAsync(outputHost + i * size, outputDev + i * size,
                    size, musaMemcpyDeviceToHost, stream[i]);
}
musaEventRecord(stop, 0);
musaEventSynchronize(stop);
float elapsedTime;
musaEventElapsedTime(&elapsedTime, start, stop);

3.2.8.7. 同步调用​

当调用同步函数时，在设备完成所请求的任务之前，控制不会返回到主机线程。主机线程是否将让出、阻塞或自旋可以通过在主机线程执行任何其他 MUSA 调用之前使用一些特定标志（详见参考手册）调用 musaSetDeviceFlags() 来指定。

3.2.9. 多设备系统​

3.2.9.1. 设备枚举​

主机系统可以拥有多个设备。以下代码示例展示了如何枚举这些设备，查询其属性，并确定 MUSA 可用设备的数量。

int deviceCount;
musaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    musaDeviceProp deviceProp;
    musaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n",
           device, deviceProp.major, deviceProp.minor);
}

3.2.9.2. 设备选择​

主机线程可以随时调用 musaSetDevice() 设置其所操作的设备。设备内存分配和 kernel launches 在当前设置的设备上进行，流和事件是与当前设置的设备相关联创建的。如果没有调用 musaSetDevice()，当前设备则为设备 0。

以下代码示例演示了如何设置当前设备的内存和 kernel 执行。

size_t size = 1024 * sizeof(float);
musaSetDevice(0);            // 设置 0 号设备为当前设备
float* p0;
musaMalloc(&p0, size);       // 在 0 号设备上分配显存
MyKernel<<<1000, 128>>>(p0); // 在 0 号设备上 kernel launches
musaSetDevice(1);            // 设置 1 号设备为当前设备
float* p1;
musaMalloc(&p1, size);       // 在 1 号设备上分配显存
MyKernel<<<1000, 128>>>(p1); // 在 1 号设备上 kernel launches

3.2.9.3. 流和事件行为​

MUSA 支持改变当前设备为 1 号设备后，在 0 号设备的 s0 流上 kernel launches，此时 kernel 会在 0 号设备上执行，如下所示。

musaSetDevice(0);               // 设置 0 号设备为当前设备
musaStream_t s0;
musaStreamCreate(&s0);          // 在 0 号设备上创建流 0
MyKernel<<<100, 64, 0, s0>>>(); // 在 0 号设备的流 0 中 kernel launches
musaSetDevice(1);               // 设置 1 号设备为当前设备
musaStream_t s1;
musaStreamCreate(&s1);          // 在 1 号设备上创建流 1
MyKernel<<<100, 64, 0, s1>>>(); // 在 1 号设备的流 1 中 kernel launches

// 此 kernel launches 是允许的:
MyKernel<<<100, 64, 0, s0>>>(); // 在 0 号设备的流 0 中 kernel launches

  • 即使在与当前设备不相关联的流上搬移数据，也会成功。

  • 如果输入事件绑定的设备和输入流所在的设备不同，则 musaEventRecord() 将失败。

  • 如果两个输入事件与不同的设备相关联，则 musaEventElapsedTime() 将失败。

  • 即使输入事件与当前设备不同的设备相关联，musaEventSynchronize() 和 musaEventQuery() 也会成功。

  • 即使输入流和输入事件与不同的设备相关联，musaStreamWaitEvent()也会成功。因此，musaStreamWaitEvent()可用于使多个设备相互同步。

每个设备都有自己的默认流（请参见默认流），因此发送到设备的默认流的命令可能与发送到任何其他设备的默认流的命令乱序或并发执行。

3.2.9.4. 点对点内存访问​

根据系统属性，特别是 PCIe 拓扑结构，多个设备可以访问彼此的内存（即，在一个设备上执行的 kernel 可以引用另一个设备的内存指针，不限于显存和锁页内存）。如果使用 musaDeviceCanAccessPeer() 查看两个设备，返回 true，则这两个设备之间支持点对点内存访问功能。

在多个 MUSA 设备共享上游 PCIe 交换器端口的拓扑结构下，musaDeviceCanAccessPeer 将返回 true。如果多个设备都直接连接根端口（一般来说，指 CPU），那么 musaDeviceCanAccessPeer 将在 AMD ZEN 及后续架构，以及 Intel 至强架构的 CPU 下返回 true（无论 MUSA 设备是否跨越 NUMA 节点以及 QPI/UPI 总线）。

两个设备使用统一的地址空间（请参见统一虚拟地址空间），因此可以使用相同的指针对这两个设备的内存寻址，如下代码所示。

musaSetDevice(0);                   // 设置 0 号设备为当前设备
float* p0;
size_t size = 1024 * sizeof(float);
musaMalloc(&p0, size);              // 在 0 号设备上分配显存
MyKernel<<<1000, 128>>>(p0);        // 在 0 号设备上 kernel launches
musaSetDevice(1);                   // 设置 1 号设备为当前设备
musaDeviceEnablePeerAccess(0, 0);   // 使能点对点访问 0 号设备内存

// 在1号设备上kernel launches
// 此 kernel 访问在设备 0 上 p0 指向的显存，是允许的
MyKernel<<<1000, 128>>>(p0);
3.2.9.4.1. Linux 上开启 IOMMU​

在 Linux 上，MUSA 当前支持在启用 IOMMU 的裸机 PCIe 点对点内存访问。当开启 IOMMU 时，点对点通信性能将可能下降，但 IOMMU 会使得更多的 MTGPU 设备支持 peer-to-peer 访问。

3.2.9.5. 点对点数据搬移​

MUSA 支持在两个设备的内存之间搬移数据（Memory Copy）。

目前 MUSA 支持统一设备虚拟地址空间，因此可以使用在设备内存 一节中提到的常规常规数据搬移函数进行操作，也可以使用 musaMemcpyPeer，musaMemcpyPeerAsync，musaMemcpy3DPeer() 或者 musaMemcpy3DPeerAsync，如下所示：

musaSetDevice(0);                   // 设置 0 号设备为当前设备
float* p0;
size_t size = 1024 * sizeof(float);
musaMalloc(&p0, size);              // 在 0 号设备上分配显存
musaSetDevice(1);                   // 设置 1 号设备为当前设备
float* p1;
musaMalloc(&p1, size);              // 在 1 号设备上分配显存
musaSetDevice(0);                   // 在 0 号设备上分配显存
MyKernel<<<1000, 128>>>(p0);        // 在 0 号设备上 kernel launches
musaSetDevice(1);                   // 在 1 号设备上分配显存
musaMemcpyPeer(p1, 1, p0, 0, size); // 拷贝 p0 指向的数据到 p1
MyKernel<<<1000, 128>>>(p1);        // 在1号设备上 kernel launches

需要注意的是，在支持统一设备虚拟地址空间的前提下，进行数据搬移不要求使用 musaSetDevice() 切换到源内存所在的设备上。

此外，在两个设备之间的数据搬移（在隐式默认流中）：

  • 在两个设备中先前发出的所有命令都完成之前不会开始；

  • 在复制到任一设备后发出的任何命令（请参见异步并发执行）开始之前，数据搬移会完成。

与流的正常行为一致，两个设备之间的异步数据搬移可能会与另一个流中的数据搬移或 kernel 重叠。

3.2.10. 统一虚拟地址空间​

当前，所有设备使用单一地址空间，通过 MUSA API 调用进行的所有锁页内存分配，以及所有设备上的显存分配都在此虚拟地址的范围内，因此：

  • 可以使用 musaPointerGetAttributes() 函数从指针的值确定通过 MUSA 分配在主机上的任何锁页内存或任何设备上的任何显存的位置。

  • 通过 musaHostAlloc() 分配的锁页内存可以在所有设备上自动移植（请参见 Portable Memory ），并且 musaHostAlloc() 返回的指针可以直接在这些设备上运行的 kernel 中使用（即无需通过 musaHostGetDevicePointer() 获取设备指针，如 Mapped Memory 中所述）。

应用程序可以通过检查 unifiedAddressing 设备属性（请参见设备枚举 ）是否等于 1 来查询某个特定设备是否使用了统一设备地址空间。

3.2.11. 进程间通信​

任何由主机线程创建的显存指针或事件句柄都可以直接被同一进程内的任何其他线程引用。然而，它在该进程之外无效，因此不能直接被属于不同进程的线程引用。

为了跨进程共享设备内存指针和事件，应用程序必须使用在参考手册中详细描述的进程间通信（Inter-process Communication, IPC）API。IPC API 在支持 MUSA 的所有设备上支持。注意，IPC API 当前不支持 musaMallocManaged，musaMallocHost，musaHostAlloc 分配的锁页内存。

使用此 API，应用程序可以使用 musaIpcGetMemHandle() 获取给定设备内存指针的 IPC 句柄，使用标准 IPC 机制（例如进程间共享内存或文件）将其传递给另一个进程，并使用 musaIpcOpenMemHandle() 从 IPC 句柄中检索设备指针，该指针是该其他进程内的有效指针。事件句柄可以使用类似的入口点共享。

注意，由 musaMalloc() 分配的显存可以基于性能原因从较大的内存块中进行子分配。在这种情况下，MUSA IPC API 将共享整个底层内存块，这可能导致其他子分配被共享，从而可能导致进程之间的信息泄漏。为了防止这种行为，建议仅共享大小对齐为 4 KiB 的 MUSA 内存。

使用 IPC API 通信的应用程序应使用相同的 MUSA 驱动程序和运行时进行编译、链接和运行。

3.2.12. 错误检查​

所有运行时函数都会返回一个错误代码，但对于异步函数（请参见异步并发执行），因为函数在设备完成任务之前就返回了，所以该错误代码可能无法报告设备上发生的任何异步错误。函数错误码仅仅可以报告执行任务之前主机上发生的错误，通常与参数校验等相关； 如果发生了异步错误，则该错误将在之后调用的运行时函数中报出。

检查异步错误的唯一方法是在调用异步函数之后立即调用同步函数，例如 musaDeviceSynchronize()（或其他任何同步机制，请参见异步并发执行），并检查 musaDeviceSynchronize() 返回的错误码。

运行时为每个主机线程维护一个错误变量，该变量初始化为 musaSuccess，并在每次发生错误时（无论是参数验证错误还是异步错误）均被错误码覆盖。

程序运行时，驱动会为每一个主机线程维护一个错误码变量，一旦发生错误（无论是参数验证还是异步错误），错误码变量都会赋值成该错误码。可以通过调用 musaPeekAtLastError() 获取此变量值，也可以通过调用 musaGetLastError() 获取此变量值并将其重置为 musaSuccess。

Kernel launches 不会返回任何错误码，因此必须在 kernel launch 后立即调用 musaPeekAtLastError() 或 musaGetLastError() 来捕捉任何启动前的错误。为了确保 musaPeekAtLastError() 或 musaGetLastError() 返回的任何错误的确是因 kernel launch 导致的，而不是 kernel launch 之前调用的其他函数引起的，必须确保在 kernel launch 之前将运行时错误码变量设置为 musaSuccess。例如，在 kernel launch 之前调用 musaGetLastError()。因为 kernel launch 是异步的，因此想要检查异步错误，应用程序必须在 kernel launch 之后，先进行同步，然后才可以调用 musaPeekAtLastError() 或 musaGetLastError() 来捕捉 kernel launch 的错误。

值得注意的是，调用 musaStreamQuery() 或 musaEventQuery() 返回的 musaErrorNotReady 不会被当作错误，也不会在 musaPeekAtLastError() 或 musaGetLastError() 被返回。

3.2.13. 调用堆栈​

使用 musaDeviceGetLimit() 可以查询调用堆栈的大小，并使用 muaDeviceSetLimit() 进行设置。

当调用堆栈溢出时，如果通过 MUSA 调试器（如 Msight 等）运行应用程序或者程序产生一个未定义的启动错误时，kernel 调用失败，并出现堆栈溢出的错误。当编译器无法确定堆栈大小时，它会报一个无法静态确定堆栈大小的警告。这种情况通常发生在应用程序使用了递归函数的时候。一旦编译器发生这种警告提示堆栈大小不够，用户需要手动设置堆栈大小。

3.2.14. Texture Memory​

MUSA 支持 GPU 用于进行图形操作的纹理内存（Texture Memory）。纹理内存相比全局内存，在读取数据方面具有一些性能优势。

有两种不同的 API 来访问纹理和曲面内存：

使用 Texture Object API 中描述的设备函数从核函数中读取 Texture 内存。使用设备函数读取 Texture 内存的过程被叫做 texture fetch。每次 texture fetch 需指定一个参数，该参数是由 Texture 对象 API 生成 Texture 对象或者由 Texture 引用 API 生成 Texture 引用。

Texture对象详细说明如下：

  • Texture，即 Texture 内存。如 Texture Object API 所述，Texture 对象在程序运行阶段创建，并且在创建 Texture 对象时指定 Texture 内容。如 Texture Reference API 所述，Texture 引用在编译阶段创建，然后通过运行时函数将 Texture 引用绑定到运行时确定的 Texture 内存上。不同的 Texture 引用可能会绑定到同一块 Texture 内存或者内存区域有重叠的不同 Texture 内存。Texture 内存可以是线性内存或者 MUSA 数组（参见 MUSA Arrays）。
  • Texture 的维数指定 Texture 是使用一维坐标寻址为一维数组，还是使用二维坐标寻址为二维数组，或者三维坐标寻址为三维数组。数组中元素称为 Texels，是 texture elements 的缩写。Texture 的宽度、高度和深度确定数组在每一个方向上的元素个数。
  • Texel 的类型，仅限于基本整数和单精度浮点类型，以及内置向量类型中定义的任何1、2和4分量向量类型，这些向量类型是从基本整数和单精度浮点类型派生而来的。
  • 读取模式，支持 musaReadModeNormalizedFloat 和 musaReadModeElementType 两种模式。如果应用程序采用 musaReadModeNormalizedFloat，并且 texel 的类型是 16 位或 8 位整数类型，则 texture fetch 返回的值实际上将返回为浮点类型，无符号整数类型映射到 [0.0 ，1.0]，有符号整形映射到 [-1.0，1.0]。例如值为0xff的无符号 8 位 Texel 读取为1。应用程序采用 musaReadModeElementType，则不执行任何转换。
  • Texture采样坐标是否归一化。默认情况下，使用范围为 [0，N-1] 的浮点坐标采样Texture，其中N是Texture在采样坐标对应维度的尺寸。例如，大小为 64x32 的 Texture 将被分别在 x 和 y 维度的 [0，63] 和 [0，31] 范围内的坐标进行采样。Texture坐标归一化将导致采样坐标在 [0.0，1.0-1/N] 范围内指定，而不是在 [0，N-1] 范围内。在采样过程中如果想不去关心 Texture 具体尺寸大小，则可以采用采样坐标归一化，这天然适合满足某些应用程序的要求。
  • 寻址模式。使用超出范围的坐标调用B.8节的设备函数是有效的。寻址模式定义了在这种情况下会发生什么。默认寻址模式是将采样有效范围的元素值作为超出边界坐标的元素值：[0，N）用于非标准化坐标，[0.0，1.0）用于归一化坐标。如果指定边界模式，则使用超出范围的Texture采样坐标进行 texture fetch 时会返回零。对于归一化坐标，还可以使用环绕模式和镜像模式。使用环绕模式时，每个坐标x都会转换为 frac（x）= x – floor（x），其中 floor（x）是不大于 x 的最大整数。当使用镜像模式时，如果 floor（x）为偶数，则每个坐标x都转换为 frac（x），如果 flor（x）是奇数，则将每个坐标 x 转换为 1-frac（x）。寻址模式被指定为大小为三的阵列，其第一、第二和第三元素分别指定第一、第二、第三纹理坐标的寻址模式；寻址模式为 musaAddressModeBorder、musaAddressModeClamp、musaAddressModeWrap 和musaAddressModeMirror；musaAddressModeWrap 和 musaAddressModeMirror 仅支持归一化采样坐标。
  • 滤波模式，指定获取 Texture 时返回的值是如何基于输入采样坐标计算的。只能对配置为返回浮点数据的 Texture 进行线性过滤。它在相邻的纹素之间进行低精度的线性插值。启用时，将读取采样坐标周围的 texel值，并根据采样坐标位于 texel 之间的位置对 texels 值进行加权插值。对一维 Texture 执行简单线性插值，对二维 Texture 执行双线性插值，对三维 Texture 执行三线性插值。滤波模式支持 musaFilterModePoint 和 musaFilterMode Linear 两种模式。如果是 musaFilterModePoint，则返回的值是坐标最接近采样坐标的texel值。如果是 musaFilterModeLinear，则返回的值是坐标最接近采样坐标的两个（一维纹理）、四个（二维纹理）或八个（三维纹理）texel 值的线性插值。musaFilterModeLinear 仅对浮点类型的返回值有效。

3.2.14.1. Texture Object API​

使用 musaCreateTextureObject(), 根据描述 Texture 资源的结构体 musaResourceDesc 来创建Texture 对象。Texture 资源描述符结构如下：

    struct musaTextureDesc
    {
        enum musaTextureAddressMode addressMode[3];
        enum musaTextureFilterMode  filterMode;
        enum musaTextureReadMode    readMode;
        int                         sRGB;
        int                         normalizedCoords;
        unsigned int                maxAnisotropy;
        enum musaTextureFilterMode  mipmapFilterMode;
        float                       mipmapLevelBias;
        float                       minMipmapLevelClamp;
        float                       maxMipmapLevelClamp;
    };
  • addressMode 指定寻址模式；
  • filterMode 指定滤波模式；
  • readMode 指定读取模式；
  • normalizedCoords 执行是否进行采样坐标归一化；
  • 查看参考手册了解 sRGB，maxAnisotropy，mipmapFilterMode，mipmapLevelBias， minMipmapLevelClamp，和 maxMipmapLevelClamp。

下面的代码示例对 texture 进行简单的转换操作。

    // Simple transformation kernel
    __global__ void transformKernel(float* output,
                                       musaTextureObject_t texObj,
                                       int width, int height,
                                       float theta)
    {
        // Calculate normalized texture coordinates
        unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    
        float u = x / (float)width;
        float v = y / (float)height;
    
        // Transform coordinates
        u -= 0.5f;
        v -= 0.5f;
        float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
        float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;
    
        // Read from texture and write to global memory
        output[y * width + x] = tex2D<float>(texObj, tu, tv);
    }
    // Host code
    int main()
    {
        const int height = 1024;
        const int width = 1024;
        float angle = 0.5;
    
        // Allocate and set some host data
        float *h_data = (float *)std::malloc(sizeof(float) * width * height);
        for (int i = 0; i < height * width; ++i)
            h_data[i] = i;
    
        // Allocate MUSA array in device memory
        musaChannelFormatDesc channelDesc =
            musaCreateChannelDesc(32, 0, 0, 0, musaChannelFormatKindFloat);
        musaArray_t muArray;
        musaMallocArray(&muArray, &channelDesc, width, height);
    
        // Set pitch of the source (the width in memory in bytes of the 2D array pointed
        // to by src, including padding), we dont have any padding
        const size_t spitch = width * sizeof(float);
        // Copy data located at address h_data in host memory to device memory
        musaMemcpy2DToArray(muArray, 0, 0, h_data, spitch, width * sizeof(float),
                              height, musaMemcpyHostToDevice);
    
        // Specify texture
        struct musaResourceDesc resDesc;
        memset(&resDesc, 0, sizeof(resDesc));
        resDesc.resType = musaResourceTypeArray;
        resDesc.res.array.array = muArray;
    
        // Specify texture object parameters
        struct musaTextureDesc texDesc;
        memset(&texDesc, 0, sizeof(texDesc));
        texDesc.addressMode[0] = musaAddressModeWrap;
        texDesc.addressMode[1] = musaAddressModeWrap;
        texDesc.filterMode = musaFilterModeLinear;
        texDesc.readMode = musaReadModeElementType;
        texDesc.normalizedCoords = 1;
    
        // Create texture object
        musaTextureObject_t texObj = 0;
        musaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);
    
        // Allocate result of transformation in device memory
        float *output;
        musaMalloc(&output, width * height * sizeof(float));
    
        // Invoke kernel
        dim3 threadsperBlock(16, 16);
        dim3 numBlocks((width + threadsperBlock.x - 1) / threadsperBlock.x,
                        (height + threadsperBlock.y - 1) / threadsperBlock.y);
        transformKernel<<<numBlocks, threadsperBlock>>>(output, texObj, width, height, angle);
        // Copy data from device back to host
        musaMemcpy(h_data, output, width * height * sizeof(float),
                    musaMemcpyDeviceToHost);
    
        // Destroy texture object
        musaDestroyTextureObject(texObj);
    
        // Free device memory
        musaFreeArray(muArray);
        musaFree(output);
    
        // Free host memory
        free(h_data);
    
        return 0;
    }

3.2.14.2. [[DEPRECATED]] Texture Reference API​

Texture 引用 API 已被弃用。

Texture 引用的一些属性是不可变的，并且必须在编译时已知；它们是在声明 Texture 引用时指定的。Textrue 引用在文件范围内声明为 texture 类型的变量：

    texture<DataType, Type, ReadMode> texRef;

参数释义:

  • DataType 指定 texel 类型;
  • Type 指定 Texture引用的类型，对于一维、二维或三维纹理，分别等于 musaTextureType1D、musaTexturesType2D 或 musaTextureType3D，对于一维或二维分层纹理，则分别等于musaTextureType1DLayered 或 musaTextureType2DLayered；Type 是一个可选参数，默认为musaTextureType1D；
  • ReadMode 指定读取模式，是一个可选参数，默认类型是 musaReadModeElementType。

Texture 引用只能声明为静态全局变量，不能作为参数传递给函数。

Texture 引用的其他属性是可变的，可以在运行时通过主机运行时进行更改。正如参考手册中所解释的，运行时 API 有一个低级 C 风格接口和一个高级C++风格接口。texture 类型在高级 API 中定义为从低级 API 中定义的 textureReference 类型公开派生的结构，如下所示：

    struct textureReference {
        int                          normalized;
        enum musaTextureFilterMode   filterMode;
        enum musaTextureAddressMode  addressMode[3];
        struct musaChannelFormatDesc channelDesc;
        int                          sRGB;
        unsigned int                 maxAnisotropy;
        enum musaTextureFilterMode   mipmapFilterMode;
        float                        mipmapLevelBias;
        float                        minMipmapLevelClamp;
        float                        maxMipmapLevelClamp;
    }
  • normalized 指定是否进行坐标归一化；
  • filterMode 指定滤波模式；
  • addressMode 指定寻址模式；
  • channelDesc 描述 texel 的数据格式，必须与 texture reference 申明中的 DataType 匹配。channelDesc类型定义如下：
    struct musaChannelFormatDesc {
      int x, y, z, w;
      enum musaChannelFormatKind f;
    };

其中 x、y、z 和 w 等于返回值的每个分量的位数，f 为

  • musaChannelFormatKindSigned 有符号整数，
  • musaChannelFormatKindUnsigned 无符号整数，
  • musaChannelFormatKindFloat 浮点数。
  • 查看参考手册了解 sRGB, maxAnisotropy, mipmapFilterMode, mipmapLevelBias, minMipmapLevelClamp, 和 maxMipmapLevelClamp.

normalized, addressMode 和 filterMode 可以在主机代码中直接修改。

在核函数可以使用 Texture 引用从 Texture 内存中读取之前，应用程序必须先使用 musaBindTexture() 或 musaBindTexture2D() 将Texture引用绑定到 Texture（用于线性内存），或使用musaBindTextureToArray（用于MUSA数组）。musaUnbindTexture() 用于取消绑定Texture引用。一旦Texture引用被解除绑定，它就可以安全地重新绑定到另一个MUSA数组，即使使用先前绑定Texture的核函数尚未结束。建议使用 musaMallocPitch() 在线性内存中分配二维纹理，并使用由 musaMallocPitch() 返回的pitch值作为 musaBindTexture2D()的输入参数。

以下代码示例将 2D 纹理引用绑定到 devPtr 指向的线性内存：

  • 使用低级 API：
    texture<float, musaTextureType2D, musaReadModeElementType> texRef;
    textureReference* texRefPtr;
    musaGetTextureReference(&texRefPtr, &texRef);
    musaChannelFormatDesc channelDesc = musaCreateChannelDesc<float>();
    size_t offset;
    musaBindTexture2D(&offset, texRefPtr, devPtr, &channelDesc, width, height, pitch);
  • 使用高级 API：
    texture<float, musaTextureType2D, musaReadModeElementType> texRef;
    musaChannelFormatDesc channelDesc = musaCreateChannelDesc<float>();
    size_t offset;
    musaBindTexture2D(&offset, texRef, devPtr, channelDesc, width, height, pitch);

以下代码示例将 2D 纹理引用绑定到 MUSA 数组 muArray:

  • 使用低级 API：
    texture<float, musaTextureType2D, musaReadModeElementType> texRef;
    textureReference* texRefPtr;
    musaGetTextureReference(&texRefPtr, &texRef);
    musaChannelFormatDesc channelDesc;
    musaGetChannelDesc(&channelDesc, muArray);
    musaBindTextureToArray(texRef, muArray, &channelDesc);
  • 使用高级 API：
    texture<float, musaTextureType2D, musaReadModeElementType> texRef;
    musaBindTextureToArray(texRef, muArray);

将 Texture 绑定到 Texture 引用时指定的格式必须与声明 Texture 引用时所指定的参数相匹配；否则，texture fetch 的结果是未定义的。

如表 15 所示，可以绑定到核函数的 Texture 数量是有限制的。

下面的代码示例将一些简单的转换内核应用于纹理

    // 2D float texture
    texture<float, musaTextureType2D, musaReadModeElementType> texRef;
    
    // Simple transformation kernel
    __global__ void transformKernel(float* output,
                                    int width, int height,
                                    float theta)
    {
        // Calculate normalized texture coordinates
        unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    
        float u = x / (float)width;
        float v = y / (float)height;
    
        // Transform coordinates
        u -= 0.5f;
        v -= 0.5f;
        float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
        float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;


        // Read from texture and write to global memory
        output[y * width + x] = tex2D(texRef, tu, tv);
    }
    
    // Host code
    int main()
    {
        // Allocate MUSA array in device memory
        musaChannelFormatDesc channelDesc =
            musaCreateChannelDesc(32, 0, 0, 0, musaChannelFormatKindFloat);
        musaArray* muArray;
        musaMallocArray(&muArray, &channelDesc, width, height);
    
        // Copy to device memory some data located at address h_data
        // in host memory
        musaMemcpyToArray(muArray, 0, 0, h_data, size,
                          musaMemcpyHostToDevice);
    
        // Set texture reference parameters
        texRef.addressMode[0] = musaAddressModeWrap;
        texRef.addressMode[1] = musaAddressModeWrap;
        texRef.filterMode     = musaFilterModeLinear;
        texRef.normalized     = true;
    
        // Bind the array to the texture reference
        musaBindTextureToArray(texRef, muArray, channelDesc);
    
        // Allocate result of transformation in device memory
        float* output;
        musaMalloc(&output, width * height * sizeof(float));
    
        // Invoke kernel
        dim3 dimBlock(16, 16);
        dim3 dimGrid((width  + dimBlock.x - 1) / dimBlock.x,
                     (height + dimBlock.y - 1) / dimBlock.y);
        transformKernel<<<dimGrid, dimBlock>>>(output, width, height, angle);
    
        // Free device memory
        musaFreeArray(muArray);
        musaFree(output);
    
        return 0;
    }

3.2.14.3. MUSA Arrays​

MUSA 数组是为 texture fetch 而优化的不透明内存布局。它们是一维、二维或三维的，由元素组成，每个元素都有 1、2 或 4 个分量，这些分量可以是有符号或无符号的 8 位、16 位或 32 位整数、16 位浮点数或32位浮点数。MUSA 数组只能由核函数通过 Texture fetch 函数进行读取和写入。

3.2.14.4. Read/Write Coherency​

Texture 内存被会缓存，并且在同一核函数调用中，该缓存在全局内存写入方面不能保持一致性，因此任何在同一核函数访问由先前该核通过全局写入生成的地址将返回未定义的数据。 换句话说，仅当该存储位置已由先前的核函数调用或内存拷贝更新时，线程才能安全地读取某些 Texture 位置，但如果被来自同一个核函数的相同线程或其他线程在先前更新过，则该线程不能安全地读取该内存位置。

3.2.15. Graphics Interoperability​

在图形渲染管线中可能需要处理一些计算任务，一些图形标准如 OpenGL、Direct3D 受制于自身 compute shader 相对较为有限的功能，无法高效处理一些复杂的计算任务。MUSA 为以上问题提供了一种解决方案，它支持将其他图形标准的资源映射到自己的地址空间中，以使 MUSA 能够读取这些资源中已经计算好的数据；或将计算结果写入这些资源，以供其他图形标准继续使用。这种能力让用户避免了进行耗时的数据拷贝，也可以发挥 MUSA 丰富灵活的编程模型和功能的优势，更好地满足了用户对图形渲染管线中复杂计算任务的需求。

在上一节提到的资源映射操作之前，需要先把图形资源注册（register）到 MUSA 中。MUSA 从设计上为不同的图形标准提供不同的 API 来注册资源，但它们统一返回一个指向 MUSA 内部定义的 musaGraphicsResource 结构体的句柄，之后对该句柄的使用将不再关心资源的来源。资源使用结束后，需要使用 musaGraphicsUnregisterResource 取消注册（unregister）。图形资源的注册及取消都是 MUSA 上下文范围中（context scope）的操作。

一旦资源注册到 MUSA，就可以使用 musaGraphicsMapResources 和 musaGraphicsUnmapResources 多次映射及取消映射资源。可以使用 musaGraphicsResourceSetMapFlags 来提示 MUSA 驱动针对资源的使用方式进行优化，如只写、只读等。如果资源已经被映射，可以通过 musaGraphicsResourceGetMappedPointer 获取该资源使用的设备内存地址，该地址可以使用 musaMemcpy 进行读写，也可以作为内核参数传入，并在内核中读写。资源处于被映射的状态时，通过其他图形 API 或另一个 MUSA 上下文访问它将会得到未定义的结果。

目前 MUSA 仅支持与 OpenGL 的互操作。

3.2.15.1. OpenGL Interoperability​

MUSA 支持使用 musaGraphicsGLRegisterBuffer 将 OpenGL 缓冲区对象（OpenGL buffer object）映射至自己的地址空间，暂不支持将 OpenGL 的纹理对象（OpenGL texture object）及渲染缓冲区对象（OpenGL renderbuffer object）映射至自己的地址空间。

用户需要保证调用 OpenGL Interoperability 相关 API 的当前主机线程所绑定的 OpenGL 上下文，与想要共享的资源所属的 OpenGL 上下文一致。

以下示例代码实现了使用 MUSA 读取 OpenGL 缓冲区对象的数据，并在计算后写入另一个 OpenGL 缓冲区对象（为了简明，代码中未处理错误码）：

#include <array>
#include <vector>
#include <execution>
#include <random>
#include <algorithm>
#include <iostream>

__device__ int coef = 4;
__global__ void mul(int* src, int* dst) {
    dst[threadIdx.x * blockDim.x + blockIdx.x] = src[threadIdx.x * blockDim.x + blockIdx.x] * coef / 128;
}

int main() {
    // Initialize OpenGL for device 0
    // and make the OpenGL context current
    ...

    // Explicitly set device 0
    musaSetDevice(0);

    // Dispatch parameters and buffer size
    constexpr size_t blockDim = 32;
    constexpr size_t gridDim = 32;
    constexpr size_t bufferSize = blockDim * gridDim;

    // Create two OpenGL buffer objects representing src and dst
    std::array<GLuint, 2> buffers;
    glGenBuffers(2, buffers.data());

    // Create a random integer generator for src
    std::random_device rd;
    std::mt19937 engine(rd());
    std::uniform_int_distribution<> dist(0, 127);
    auto integerGenerator = [&engine, &dist] { return dist(engine); };

    // Generate random data, allocate device memory for src, and copy data to it
    std::vector<int> srcData(bufferSize);
    std::generate(std::execution::seq, srcData.begin(), srcData.end(), integerGenerator);
    glBindBuffer(GL_ARRAY_BUFFER, buffers[0]);
    glBufferData(GL_ARRAY_BUFFER, bufferSize * sizeof(int), srcData.data(), GL_STATIC_READ);
    glBindBuffer(GL_ARRAY_BUFFER, 0);

    // Allocate device memory with uninitialized data for dst
    glBindBuffer(GL_ARRAY_BUFFER, buffers[1]);
    glBufferData(GL_ARRAY_BUFFER, bufferSize * sizeof(int), nullptr, GL_STATIC_DRAW);
    glBindBuffer(GL_ARRAY_BUFFER, 0);

    // ======= Begin of interoperation =======
    // Step1: register OpenGL buffer objects to MUSA
    std::array<musaGraphicsResource*, 2> resources;
    musaGraphicsGLRegisterBuffer(&resources[0], buffers[0], musaGraphicsRegisterFlagsReadOnly);
    musaGraphicsGLRegisterBuffer(&resources[1], buffers[1], musaGraphicsRegisterFlagsWriteDiscard);

    // Step2: then map the registered resources to MUSA address space, and get the mapped device addresses
    musaGraphicsMapResources(2, resources.data(), 0);
    void* srcDevPtr;
    void* dstDevPtr;
    size_t size;
    musaGraphicsResourceGetMappedPointer(&srcDevPtr, &size, resources[0]);
    musaGraphicsResourceGetMappedPointer(&dstDevPtr, &size, resources[1]);

    // Step3: so that we can direct access these resources in MUSA kernel
    mul<<<gridDim, blockDim>>>(srcDevPtr, dstDevPtr);
    musaDeviceSynchronize();

    // Step4: remember to unmap the resources after use
    musaGraphicsUnmapResources(2, resources.data(), 0);

    // Step5: and finally unregister the resources, after when OpenGL could continue to use
    musaGraphicsUnregisterResource(resources[0]);
    musaGraphicsUnregisterResource(resources[1]);
    // ======= End of interoperation =======

    // Map the dst OpenGL buffer object to the host and check the answers
    glBindBuffer(GL_ARRAY_BUFFER, buffers[1]);
    int* dstHostPtr = static_cast<int*>(glMapBuffer(GL_ARRAY_BUFFER, MAP_READ_BIT));
    for (size_t i = 0; i < bufferSize; i++) {
        if (dstHostPtr[i] != srcData[i] * 4 / 128) {
            std::cerr << "Got wrong answer!" << std::endl;
        }
    }
    glUnmapBuffer(GL_ARRAY_BUFFER);
    glBindBuffer(GL_ARRAY_BUFFER, 0);

    // Destroy OpenGL buffer objects and their device memories
    glDeleteBuffers(2, buffers.data());

    // Other resource destructions
    ...

    std::cout << "OpenGL Interoperability test finished!" << std::endl;
}

3.2.17 驱动 API​

3.2.17.1 上下文​

在 driver API 中执行的所有资源和操作都封装在MUSA 上下文（context）中。当 context 被销毁时，MUSA驱动会自动清理这些资源。

一个主机线程一次可能只有一个当前的设备上下文。使用 muCtxCreate() 创建 context 时，这个 context 将成为当前线程的默认context，每个线程可以有不同的默认 context。在 context 中调用大多数不涉及设备枚举或 context 管理的 MUSA 函数，如果不是当前线程的有效的 context，函数将返回 MUSA_ERROR_INVALID_CONTEXT。

每个主机线程都可能绑定了一堆 context，但只能有一个默认 context，musa 用栈来管理这些绑定的 context。muCtxCreate() 会将新的 context 推入栈顶；muCtxPopCurrent() 会重置当前主机线程默认 context，所以你可以调用 muCtxPopCurrent() 以从主机线程剥离当个 context。所有 context 都可能被推到任何主机线程的 context 栈中。

context 管理模块还为每个 context 维护了一个使用计数。muCtxCreate() 会创建一个使用计数为 1 的 context。muCtxAttach() 会增加使用计数，而 muCtxDetach() 将减少计数。当调用 muCtxDetach() 或 muCtxDestroy() 导致使用计数变为 0 时，context 将被销毁。
driver API 可与 runtime API 进行互操作，并且可以通过 muDevicePrimaryCtxRetain() 从 driver API 访问 runtime 管理的 primary context。

使用计数促进了在同一 context 中运行的第三方编写的代码之间的互操作性。例如，如果同时加载三个库且让它们使用相同的 context，每个库将调用 muCtxAttach() 来增加使用计数，并调用 muCtxDetach() 来减少使用计数。对于大多数库，预计应用程序会在加载或初始化库之前创建 context；这样，应用程序就可以使用自己的启发式方法创建 context，而库只对传递给它的 context 进行操作。对于希望创建自己的 context 的库，它们的用户不知道它们是否已经创建了自己的 context ，可以配合使用 muCtxPushCurrent() 和muCtxPopCurrent() 来获取可用的 context。

3.2.17.2 模块​

模块（Module）是设备端代码和数据的动态可加载包，类似于 Linux 的动态库.so，可由 mcc 编译生成；所有 including 函数、全局变量、texture 或接口引用等符号均在 module 中维护，以便独立第三方编写的 module 可以在相同的 MUSA context 中进行互操作。

以下是加载一个 module 并检索 'kernel_test' 命名的 kernel 句柄的完整步骤：

    status = muInit(0);
    if (status != MUSA_SUCCESS) {
        std::cout << "muInit failed!" << std::endl;
        exit(EXIT_FAILURE);
    }
    
    MUcontext ctx = nullptr;
    status = muCtxCreate(&ctx, 0, 0);
    if (status != MUSA_SUCCESS) {
        std::cout << "muCtxCreate failed!" << std::endl;
        exit(EXIT_FAILURE);
    }
    
    MUmodule module;
    MUfunction function;
    
    status = muModuleLoad(&module, "kernel_test.elf");
    if (status != MUSA_SUCCESS) {
        std::cout << "muModuleLoadData failed!" << std::endl;
        exit(EXIT_FAILURE);
    }
    std::cout << "muModuleLoadData Success!" << std::endl;
    
    status = muModuleGetFunction(&function, module, "kernel_test");
    if (status != MUSA_SUCCESS) {
        std::cout << "muModuleGetFunction failed!" << std::endl;
        exit(EXIT_FAILURE);
    }
    
    int threadsPerBlock = 1024;
    int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
    void* extra[] = {};
    
    status = muLaunchKernel(function,
                            blocksPerGrid, 1, 1,   /* grid dim */
                            threadsPerBlock, 1, 1, /* block dim */
                            0,                     /* shared mem */
                            0,                     /* stream */
                            nullptr,               /* params */
                            extra                  /* extra */
                            );
    
    if (status != MUSA_SUCCESS) {
        std::cout << "muLaunchKernel failed!" << std::endl;
        exit(EXIT_FAILURE);
    }

3.2.17.3 Kernel 执行​

muLaunchKernel() 可以用来启动具有给定参数配置的 kernel。

muLaunchKernel() 的倒数第二个参数可以用来传递参数构成的指针数组，其中第 n 个指针对应于第 n 个参数，并指向从中复制参数的内存区域。muLaunchKernel() 的最后一个参数是用来传递额外选项的，您也可以选择将参数数组指针作为额外选项之一来传参。当参数作为额外选项（MU_LAUNCH_PARAM_BUFFER_POINTER 选项）传递时，它们作为指向单个缓冲区的指针传递，但要求参数的偏移匹配设备代码中每个参数类型的对齐要求。
Memory Synchronization Domains 中列出了内置向量类型的设备端代码中的对齐要求。对于所有其他基本类型，设备代码中的对齐要求与主机代码中的对齐要求相匹配，因此可以使用 alignof() 获得。唯一的例外是当主机编译器将 double、long long 以及 64 位系统上的 long 类型对齐到单 Word 边界而不是双 Word 边界时（例如使用 gcc 的编译标志 -mno-align-double），因为在设备代码中，这些类型总是按双 Word 对齐；MUdeviceptr 是一个整数，但用于代表一个指针，所以它的对齐要求是 alignof(void*)。

以下代码示例使用宏 (ALIGN_UP()) 调整每个参数的偏移量以满足其对齐要求，并使用另一个宏 (ADD_TO_PARAM_BUFFER()) 将每个参数添加到传递给 MU_LAUNCH_PARAM_BUFFER_POINTER 选项的参数缓冲区。

    #define ALIGN_UP(offset, alignment) \
        (offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)

    char paramBuffer[1024];
    size_t paramBufferSize = 0;
    
    #define ADD_TO_PARAM_BUFFER(value, alignment)                   \
        do {                                                        \
            paramBufferSize = ALIGN_UP(paramBufferSize, alignment); \
            memcpy(paramBuffer + paramBufferSize,                   \
                &(value), sizeof(value));                        \
            paramBufferSize += sizeof(value);                       \
        } while (0)
    
    int i;
    ADD_TO_PARAM_BUFFER(i, __alignof(i));
    float4 f4;
    ADD_TO_PARAM_BUFFER(f4, 16); // float4's alignment is 16
    char c;
    ADD_TO_PARAM_BUFFER(c, __alignof(c));
    float f;
    ADD_TO_PARAM_BUFFER(f, __alignof(f));
    MUdeviceptr devPtr;
    ADD_TO_PARAM_BUFFER(devPtr, __alignof(devPtr));
    float2 f2;
    ADD_TO_PARAM_BUFFER(f2, 8); // float2's alignment is 8
    
    void* extra[] = {
        MU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
        MU_LAUNCH_PARAM_BUFFER_SIZE,    &paramBufferSize,
        MU_LAUNCH_PARAM_END
    };
    muLaunchKernel(muFunction,
                blockWidth, blockHeight, blockDepth,
                gridWidth, gridHeight, gridDepth,
                0, 0, 0, extra);

注意：主机端和设备端对于字节对齐的要求不一致，结构体的对齐要求等于其字段对齐要求的最大值。因此，包含内置向量类型、MUdeviceptr 或非对齐 double 和 long long 的结构，对齐要求可能因设备代码和主机代码而异。这样的结构也可能以不同的方式填充。例如，以下结构在主机代码中没有填充，但在设备代码中在字段 f 之后填充了 12 个字节，因为字段 f4 的对齐要求是 16。

typedef struct {
    float  f;
    float4 f4;
} myStruct;

3.2.17.4 运行时和驱动 API 之间的互操作​

应用程序可以使用运行时 API（Runtime API） 和驱动 API（Driver API）编程，设备和版本管理部分的所有功能都可以互换使用。
runtime API 会隐式创建 context，但如果通过 driver API 创建 context 并设置其为当前 context，则后续 runtime API 调用时将选择此context，而不会创建新的 context。

调用 runtime API 会做隐式初始化，由 runtime 隐式创建的 context 称为 primary context，driver API 可以管理 primary context 的所有功能。如果 runtime 已初始化，则 muCtxGetCurrent() 可用于检索在初始化期间创建的 context，此 context 可由后续 driver API 调用。

用户可以使用 runtime 或 driver 中任一分配和释放设备内存的 API。MUdevicptr 可以转换为常规指针，反之亦然：

    MUdeviceptr devPtr;
    float* d_data;

    // Allocation using driver API
    muMemAlloc(&devPtr, size);
    d_data = (float*)devPtr;
    
    // Allocation using runtime API
    mudaMalloc(&d_data, size);
    devPtr = (MUdeviceptr)d_data;

这也意味着使用 driver API 编写的应用程序可以调用 runtime API 编写的库（例如 muBLAS、muFFT、muRand 等）。

3.3. MUSA 编译器简介​

MUSA 应用程序源码由 Host 侧传统的 c++ 语言，Device 侧 C++ 语言扩展组成，这使得一个 MUSA 应用程序源码的编译相对传统的 C++ 源码而言变得更加复杂。mcc 是 MUSA 编程语言配套的编译器工具链，用于将 MUSA 应用程序源码编译为可以被 GPU 识别的可执行文件，为用户隐藏复杂的 MUSA 应用程序源码编译流程。此外，mcc 将针对不同 MTGPU 硬件架构对用户 MUSA 代码进行优化，保证一定的性能。

3.3.1. 编译流程​

mcc 将分别对 MUSA 源码进行“两趟”编译，分别处理 host 侧代码与 device 侧代码，整体流程如下图所示：

img3.3.1

编译流程图

3.4. MUSA 数学加速库简介​

摩尔线程数学库当前提供了基础线性代数算子库 muBLAS（MUSA Basic Linear Algebra Subprograms）, 快速傅里叶变换库 muFFT（MUSA Fast Fourier Transform）和基础性能元库 muPP (MUSA Performance Primitives)。

3.4.1. muBLAS​

muBLAS 提供 BLAS-1/2/3 三种级别的算子。

  • BLAS-1 算子包括各种标量、向量和向量－向量计算，例如缩放、Givens 旋转和点积等；

  • BLAS-2 算子包括各种矩阵－向量计算，例如矩阵向量乘、Hermitian 秩 1 修正、三角方程组求解等；

  • BLAS-3 算子提供了各种矩阵－矩阵计算，包括各种类型的矩阵乘和 Hermitian 秩 K 修正等。

muBLAS 算子提供了各种基本线性代数计算的 GPU 加速, 开发者可以利用 BLAS 算子构建基于摩尔线程 GPU 产品的复杂数值计算方法。

当前版本 muBLAS 算子当前支持单精度实数和单精度复数两种数据类型，未来版本会添加更多数据类型的支持。

以下示例代码以一段单精度矩阵向量乘算子调用为例，使用 muBLAS 的基本流程是：

  1. 创建并配置 muBLASHandle 句柄；
  2. 分配必要的 GPU 存储资源；
  3. 调用相应的 muBLAS 算子；
  4. 运算结束后，释放资源。
{
  mublasHandle_t handle;
  mublasCreate(&handle);
  mublasSetStream(handle, streamIdx);
  float *A, *X, *Y;
  float alpha = 1.0f;
  float beta = 1.0f
  int m = 8192;
  int n = 8192;
  int lda = 8192;
    ...
  mublasSgemv(handle, MUBLAS_OP_N, m, n, &alpha, A, lda, X, 1, &beta, Y, 1);
  musaStreamSynchronize(streamIdx);
    ...
  mublasDestroy(handle);
  musaFree(A);
  musaFree(X);
  musaFree(Y);
}

3.4.2. muFFT​

muFFT 提供了 1/2/3 维快速傅里叶变换算子，可以支持单精度实数和单精度复数两种数据类型和不同数据大小的离散傅里叶变换，muFFT 针对最常见的 2 幂长度进行了特别的优化。

当前版本的 muFFT 支持单精度实数－单精度复数、单精度复数－单精度实数和单精度复数－单精度复数变换，并且支持多个相同大小的问题的批量处理，以提高计算效率。

以下是一段一维傅里叶变换的代码示例，NX 为所变换数组的长度，BATCH 为批处理的数量，代码中可以看出使用 muFFT 进行快速傅里叶变换的主要流程为：

  1. 创建并配置 muFFTHandle 句柄，并为问题申请必要的 GPU 存储资源；
  2. 通过 mufftPlan 配置所需要处理的问题；
  3. 调用问题对用的 mufftExecC2C（例子中为单精度复数类型）执行计算；
  4. 在程序运行完成后释放资源。
#define NX 256
#define BATCH 10
...
{
  mufftHandle plan;
  mufftComplex *data;
    ...
  musaMalloc((void**)&data, sizeof(mufftComplex)*NX*BATCH);
  mufftPlan1D(&plan, NX, MUFFT_C2C, BATCH);
    ...
  mufftExecC2C(plan, data, data, MUFFT_FORWARD);
  musaDeviceSynchronize();
    ...
  mufftDestroy(plan);
  musaFree(data);
}

3.4.3. muPP​

muPP 是摩尔线程公司推出的基础性能元算子库，提供了各种图像处理和数据处理算子，主要包括：

  • 代数算子，如 SubC, DivC 等颜色通道代数操作及阿尔法复合等；
  • 颜色转换算子，如 RGB 转灰度图，RGB 转 YUV 算子等；
  • 数据转换和初始化算子，包括拷贝、通道转换、截取等；
  • 统计算子，包括大小值和范数等；
  • 滤波算子，包括中值滤波和高斯滤波等。

以实现 RGB 转 YUV 的应用为例，调用 muPP。muPP 库的调用也非常简单，在分配好必要的资源后，直接调用库中的相应算子即可：

...
{
  float *mu_src_4k, *mu_dst_4k;
  MUppiSize mu_roi_4k{col, row};
  int mupp_src_step = mu_roi_4k.width * cn * sizeof(unsigned char);
  int mupp_dst_step = mu_roi_4k.width * cn * sizeof(unsigned char);
    ...
  muppiRGBToYUV_8u_C3R(reinterpret_cast<const MUpp8u*>(mu_src_4k), mupp_src_step,    reinterpret_cast<MUpp8u*>(mu_dst_4k), mupp_dst_step, mu_roi_4k);
    ...
  musaDeviceSynchronize();
  musaFree(mu_src_4k);
  musaFree(mu_dst_4k);
}

版权所有 2024 摩尔线程，本文档受国际版权法保护。

摩尔线程和摩尔线程徽标是摩尔线程智能科技 (北京)有限责任公司的注册商标。

免责声明

本文档提供有关摩尔线程产品的信息。本文档并未授权任何知识产权的许可，并未以明示或暗示，或以禁止反言或其他方式授予任何知识产权许可。