Rocm-HIP kernel language

HIP的内核启动语法hipLaunchKernelGGL是一个宏，可以作为启动内核的替代方式，它接受启动配置的参数（网格维度、分组维度、流、动态共享大小）以及任意数量的内核参数。这个宏可以替代CUDA中的三连字符（<<< >>>）启动语法。

HIP-Clang作为ROCm平台的一部分，是用于编译HIP程序的新编译器，它使用与GCC兼容的API，允许由不同GCC兼容编译器生成的代码相互链接。

一、Function-Type Qualifiers

在HIP（Heterogeneous-compute Interface for Portability）中，Function-Type Qualifiers是一组用于定义函数在主机（host）或设备（device）上执行的关键字。

1. __device__

• 标记为 __device__ 的函数只能在设备（即GPU）上执行。
• 这些函数只能从设备代码中调用。
• __device__ 关键字可以与 __host__ 关键字结合使用，这样标记的函数将同时编译为在主机和设备上执行。但当这样结合使用时，函数不能使用HIP的设备端坐标函数，例如 hipThreadIdx_x。如果需要在函数内部使用这些坐标信息，可以将它们作为参数传递给函数。

2. __global__

• 标记为 __global__ 的函数也在设备上执行，但与 __device__ 不同的是，它们是从主机代码中“启动”或“调用”的。

• HIP中的 __global__ 函数必须有 void 返回类型。

• 这些函数通常指的是内核（kernel）函数，它们的执行需要主机端提供执行配置，包括网格（grid）和块（block）的维度信息。执行配置还可以包括其他信息，比如分配的共享内存大小和内核执行的流（stream）。

• __global__ 函数的启动可以通过HIP提供的 hipLaunchKernelGGL 函数进行，这个函数接受内核名称、网格维度、块维度、动态共享内存大小和流等参数，然后是内核参数。

HIP目前不支持动态并行性，这意味着 __global__ 函数不能从设备代码中调用。

3. __host__

• 执行位置：标记为 __host__ 的函数在主机上执行。
• 调用位置：这些函数只能从主机代码中调用。

HIP 编译器解析 __noinline__ 和 __forceinline__ 关键字，并将它们转换为适当的 Clang 编译器属性。这些关键字分别用于建议编译器不要内联函数和强制内联函数。然而，值得注意的是，在使用 HCC（HIP 编译器的后端之一）时，所有设备函数都会被内联，因此这些关键字实际上可能被忽略。

调用核函数

在HIP中，内核可以通过以下两种方式启动：

• 使用传统的CUDA语法 <<< >>> 来启动内核。
• 使用 hipLaunchKernel 函数，这是HIP提供的C++标准调用约定，用于传递执行配置。

hipLaunchKernel 宏总是以五个基本参数开始，然后是内核参数。

（1）内核名称（kernelName）：使用HIP_KERNEL_NAME宏来支持包含逗号的模板内核。hipify工具会自动插入这个宏。

（2）网格维度（gridDim）：使用dim3类型的变量来指定3D网格维度，这决定了要启动的块的数量。

（3）块维度（blockDim）：同样使用dim3类型的变量来指定每个块中的线程数量。

（4）动态共享内存（dynamicShared）：指定启动内核时要分配的额外共享内存的大小。

（5）流（hipStream_t）：指定内核应该在哪个流中执行。值为0对应于NULL流，NULL流是GPU上的默认执行流。

（6）内核参数：在上述五个参数之后，跟随的是传递给内核的参数。

note：dim3 构造函数可以接受零到三个参数，未指定的维度将默认初始化为1。这意味着你可以指定一维、二维或三维的网格和块维度。

举个栗子：

二、Variable-Type Qualifiers

在HIP（Heterogeneous-compute Interface for Portability）中，变量类型限定符（Variable-Type Qualifiers）用于定义变量的作用域和生命周期，以及它们是如何在主机（host）和设备（device）之间使用的。

1. __constant__

使用 __constant__ 关键字声明的变量存储在GPU的常量内存中。这些变量在主机代码中初始化，并在内核启动之前加载到GPU。从GPU内核中访问这些变量时，它们是只读的。这意味着在内核执行期间，不能修改这些变量的值。

HIP提供了以下函数来访问常量内存中的变量：

• hipGetSymbolAddress(): 获取常量内存中符号的地址。
• hipGetSymbolSize(): 获取常量内存中符号的大小。
• hipMemcpyToSymbol(): 从主机内存复制数据到常量内存。
• hipMemcpyToSymbolAsync(): 异步地从主机内存复制数据到常量内存。
• hipMemcpyFromSymbol(): 从常量内存复制数据到主机内存。
• hipMemcpyFromSymbolAsync(): 异步地从常量内存复制数据到主机内存。

2. __shared__

在HIP中，__shared__ 关键字用于声明在内核（kernel）中由块（block）内所有线程共享的内存。这种内存在内核执行期间分配，并在内核执行结束后释放。

1. 使用 extern __shared__ 允许主机端动态分配共享内存，并且共享内存的大小作为内核启动参数来指定。
2. 在调用 hipLaunchKernel 或使用 <<< >>> 语法启动内核时，需要指定共享内存的大小。例如，在使用 hipLaunchKernel 时，可以通过一个参数来指定所需的共享内存字节数。
3. 在使用共享内存时，需要确保所有线程在访问共享内存之前已经同步，通常使用 __syncthreads() 函数来实现。

__global__ void myKernel(int *output, int *input, int N) {
    extern __shared__ int sdata[]; // 声明共享内存
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 将数据从全局内存复制到共享内存
    sdata[tid] = input[i];
    __syncthreads(); // 确保所有线程完成复制

    // 执行一些操作...
    // 例如：sdata[tid] 与 sdata[tid + 1] 相加

    // 将结果从共享内存复制到全局内存
    output[i] = sdata[tid];
}

int main() {
    // 假设已经分配并初始化了 input, output 等变量
    int N = 256;
    size_t sharedMemSize = N * sizeof(int); // 计算共享内存大小
    hipLaunchKernelGGL(
        myKernel, dim3(N / 256), dim3(256), sharedMemSize, 0, output, input, N);
    // ...
}

3. __managed__

1. 使用 __managed__ 声明的内存由HIP自动管理，它会根据需要自动在主机和设备之间迁移数据。
2. __managed__ 内存可以被主机和设备代码访问。在设备代码中访问 __managed__ 内存时，无需进行额外的同步操作。

3. _managed__ 内存可以使用 hipMallocManaged() 函数进行分配，使用 hipFree() 函数进行释放。

4. 内存属性：__managed__ 内存具有一些属性，例如是否可以被设备全局访问，是否支持内存映射等。这些属性可以在分配内存时通过标志来指定。

__managed__ int* managedArray;

// 在主机代码中分配内存
hipMallocManaged(&managedArray, N * sizeof(int));

// 初始化内存
for (int i = 0; i < N; i++) {
    managedArray[i] = i;
}

// 定义设备内核，可以直接访问 managedArray
__global__ void kernelExample(int* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    // 直接访问 managedArray，无需额外同步
    data[idx] *= 2;
}

// 启动内核
kernelExample<<<gridSize, blockSize>>>(managedArray);

// 同步，确保设备完成执行
hipDeviceSynchronize();

// 使用内存数组
for (int i = 0; i < N; i++) {
    // managedArray 已经被内核修改
    printf("%d ", managedArray[i]);
}

// 释放内存
hipFree(managedArray);

4. __restrict__

在HIP中，__restrict__ 关键字是一个提示（hint），告诉编译器该指针在当前函数或内核中不会与其他指针别名（alias），即不会有多个指针指向同一块内存。这个提示可以帮助编译器进行更有效的优化，因为它允许编译器假设对这些指针的访问不会相互冲突。

三、Built-In Variables

1. Coordinate Built-Ins

2. warpSize

在HIP中，warpSize 是一个内建变量，用于确定目标设备上的 warp（在NVIDIA GPU中通常称为warp，在AMD GPU中称为wavefront）大小，即每个warp包含的线程数。这个变量的类型是 int，并且其值依赖于GPU架构：

• 对于NVIDIA设备，warpSize 通常返回 32，因为NVIDIA GPU的warp大小是32个线程。
• 对于AMD设备，warpSize 返回 64，因为AMD GPU的wavefront大小是64个线程。

warpSize 应在设备代码中使用，而不是在主机代码中，因为它依赖于执行内核的GPU的特性。

四、Vector Types

在HIP中，向量类型（Vector Types）是一组在 hip_runtime.h 头文件中定义的数据结构，它们用于表示和操作向量数据。

1. short vector types

在HIP中，短向量类型（Short vector types）是基本整数和浮点类型派生的复合数据结构，它们在hip_vector_types.h头文件中定义。

以下是短向量类型的关键特性：

1. 派生自基本类型：短向量类型基于基本的整数（如int）和浮点（如float）类型。

2. 结构体定义：这些类型作为结构体在hip_vector_types.h中定义，每个结构体包含四个字段：x, y, z, 和 w，分别对应向量的第1、2、3、4个分量。

3. 分量访问：向量的分量可以通过点操作符访问，例如，vector.x访问向量的第一个分量。

4. 构造函数：所有短向量类型都支持一个构造函数，其形式为make_<type_name>()。这个构造函数用于创建具有给定分量值的新向量实例。

HIP支持以下短向量类型：

举个栗子：

#include <hip/hip_vector_types.h>

// 使用构造函数创建一个float4类型的向量
float4 myVector = make_float4(1.0f, 2.0f, 3.0f, 4.0f);

// 访问向量的分量
float x = myVector.x;
float y = myVector.y;
float z = myVector.z;
float w = myVector.w;

__global__ void vectorKernel(float4* output, float4* input) {
    int idx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
    output[idx] = input[idx] * myVector; // 向量乘法示例
}

int main() {
    // 假设分配内存和初始化数据的代码已经完成
    float4 *dev_output, *dev_input;
    // ... 初始化和分配内存 ...

    // 启动内核
    dim3 blockSize(256);
    dim3 gridSize((N + 255) / 256);
    hipLaunchKernelGGL(vectorKernel, gridSize, blockSize, 0, 0, dev_output, dev_input);

    // 销毁资源的代码
    // ... 释放内存 ...
    return 0;
}

2. dim3

dim3 是 HIP 中的一个三维度的整型向量类型，它通常用于指定 GPU 内核启动时的网格（grid）和块（block）的维度。dim3 向量可以非常方便地表达一维、二维或三维的并行结构。

1. 三维度：dim3 包含三个整数字段：x，y，和 z，分别代表三个维度的大小。
2. 默认初始化：如果在使用 dim3 构造函数时没有指定所有三个维度的大小，未指定的维度将默认初始化为 1。这意味着如果你只指定了 x 维度，y 和 z 维度将自动设置为 1。
3. 灵活性：dim3 允许你为 GPU 内核定义从一维到三维的并行执行空间。
4. 使用场景：
   1. 在设置内核的网格维度时，dim3 用来指定每个维度上的块数量。
   2. 在设置内核的块维度时，dim3 用来指定每个块在每个维度上的线程数量。

五、Memory-Fence Instructions

在HIP（Heterogeneous-compute Interface for Portability）中，内存屏障函数__threadfence()和__threadfence_block()用于确保在并行线程中的内存操作顺序性。这些函数类似于CUDA中的__threadfence()和__threadfence_block()，它们分别确保所有先前的内存访问（包括所有类型的内存）和当前线程块内所有先前的内存访问（只包括全局和共享内存）完成。

__threadfence()：这个函数确保当前线程中的所有先前的内存访问在任何后续内存访问之前完成。它影响所有类型的内存，包括全局、共享、局部和常量内存。

__threadfence_block()：这个函数仅确保当前线程块内的所有先前的内存访问在任何后续内存访问之前完成。它主要影响全局和共享内存，而不包括局部和常量内存。

threadfence_system() 工作区：HIP提供了对threadfence_system()的替代实现，这在HIP-Clang编译器路径下可用。为了启用这个工作区，需要在构建HIP时设置环境变量HIP_COHERENT_HOST_ALLOC。

1. 使用条件：要使用__threadfence_system()的内核应该只操作细粒度系统内存，这种内存应该使用hipHostMalloc()进行分配。

2. memcpy的移除：对于使用__threadfence_system()的细粒度系统内存区域，应该移除所有相关的memcpy操作。

3. 内存一致性：当使用hipHostMalloc()分配的内存时，通常这种内存是一致的，意味着它不会被GPU缓存（当 GPU 访问这种内存时，它总是直接从内存中读取数据，而不是从自己的缓存中读取，这确保了对内存的写入对 GPU 来说是立即可见的，无需等待缓存刷新或同步操作），从而保证了数据的一致性。如果需要，可以覆盖这一行为，允许GPU缓存这些内存，但这需要开发者根据性能需求和数据一致性要求来决定。

六、Math Functions

1. Single Precision Mathematical Functions

2. Double Precision Mathematical Functions

3. Integer Intrinsics：supported on devices only.

4. Floating-point Intrinsics：supported on devices only.

5. Texture Functions

6. Surface Functions: Surface functions are not supported.

7. Timer Functions

1. clock(): 这个函数返回一个clock_t类型的值，它是一个在每个时钟周期递增的计数器的值。通过计算两次调用clock()返回值的差值，可以确定内核执行所消耗的时钟周期数。

2. clock64(): 类似于clock()，clock64()返回一个long long int类型的值，表示64位的计时器计数。这允许在可能的情况下测量更长的时间跨度，而不会因为计数器溢出而丢失信息。

#include <hip/hip_runtime.h>

__global__ void myKernel() {
    // 内核代码
}

int main() {
    // 启动内核之前记录计时器
    clock_t start = clock();

    // 启动内核
    hipLaunchKernelGGL(myKernel, dim3(256), dim3(256), 0, 0);

    // 内核执行完成后记录计时器
    clock_t end = clock();

    // 计算内核执行所需的时钟周期数
    double kernelTime = (double)(end - start);

    printf("Kernel took %f clock cycles to execute.\n", kernelTime);

    return 0;
}

8. Atomic Functions

原子函数（Atomic functions）是用于在并行环境中执行原子操作的函数，这些操作是不可分割的，即在操作完成之前，不会被其他线程或设备中断。在HIP（Heterogeneous-compute Interface for Portability）中，原子函数通常用于全局内存或共享内存中，以确保数据的一致性和线程安全。

1. 原子操作：原子函数执行读-修改-写（read-modify-write）操作，保证在并行环境中对内存位置的访问是原子的。

2. 内存位置的独占性：在原子操作执行期间，没有其他设备或线程能够观察或修改该内存位置。

3. 指令序列化：如果来自不同设备或线程的多个指令目标是同一内存位置，这些指令将被序列化，但序列化顺序是未定义的。

4. 系统范围的原子操作：HIP引入了带有_system后缀的新API，以支持系统范围的原子操作。这意味着原子操作不仅局限于GPU设备，还可以扩展到系统范围，包括从GPU设备到系统中的其他CPU和GPU设备。

9. Warp Cross-Lane Functions

Warp Cross-Lane Functions 是一类在同一个Warp内所有线程上执行的函数，这些函数不需要额外的同步机制或共享内存，因为硬件保证了所有Warp内的线程会锁定步调（lockstep）执行。（即要么所有线程执行指令，要么都不执行。）

Warp Cross-Lane Functions的要点：

1. 跨通道操作：这些函数在Warp内的所有线程上操作，可以执行跨线程的数据传输或同步操作。

2. 无需同步：由于硬件保证了Warp内线程的执行同步，因此使用这些函数时不需要额外的同步指令。

3. 不使用共享内存：Warp Cross-Lane Functions执行时不依赖于共享内存。

4. 不同GPU架构的Warp大小：NVIDIA和AMD GPU的Warp大小可能不同。NVIDIA GPU通常使用32个线程的Warp，而AMD GPU可能使用64个线程的Wavefront。

5. 可移植代码：为了编写可在不同GPU架构上运行的可移植代码，应使用warpSize内建变量查询Warp大小，而不是硬编码为32或64。

6. "Wave-aware"代码：如果代码假设Warp大小为32，在支持64个线程的Wavefront的AMD GPU上运行时，可能只利用了一半的资源。

7. 内建变量：warpSize内建变量只能在设备函数中使用，并且它的值取决于GPU架构。

8. 主机端查询：主机端函数应使用hipGetDeviceProperties来获取GPU设备的默认Warp大小。

   hipDeviceProp_t props;
   hipGetDeviceProperties(&props, deviceID);
   int warpSize = props.warpSize;
   // 基于warpSize实现可移植算法

9. 汇编内核：一些为特定Warp大小编写的汇编内核可能与默认Warp大小不同。

举个栗子：

#include <hip/hip_runtime.h>

__device__ __forceinline__ int getLaneIndex() {
    return hipThreadIdx_x & (warpSize - 1);
}

__global__ void warpCrossLaneKernel() {
    int laneIndex = getLaneIndex();
    // 使用Warp Cross-Lane Functions执行操作
    // 例如：使用shfl, popc等函数
}

int main() {
    // 假设已经设置了设备和内核参数
    hipLaunchKernelGGL(warpCrossLaneKernel, grid, block, 0, 0);
    return 0;
}

在这个示例中，getLaneIndex函数用于获取线程在其Warp内的索引。然后可以在内核中使用这个索引来使用Warp Cross-Lane Functions。

开发者需要注意，当从CUDA路径迁移代码到HIP时，需要仔细检查并确保代码没有假设Warp大小为32，以确保在不同的GPU架构上都能正确运行。

10. Warp Vote and Ballot Functions

在GPU编程中，一个Warp（在NVIDIA GPU中）或Wavefront（在AMD GPU中）是一组同时执行相同指令的线程，它们在逻辑上被视为一个单元。这些线程在执行时是"锁定步调"的，即要么所有线程执行指令，要么都不执行。这种特性使得Warp或Wavefront内的线程可以进行协作操作，而不必担心执行顺序或线程间的同步问题。

Warp Vote和Ballot Functions 就是在这种锁定步调的执行模型下，用于在同一个Warp或Wavefront内进行线程间协作的一组函数。

1. Warp Vote Functions：这些函数允许Warp内的每个线程对某个条件进行投票，并根据投票结果进行操作。
   1. __any(int predicate)：如果Warp中至少有一个线程的谓词为真（非零），则返回1，表示至少有一个线程满足条件。
   2. __all(int predicate)：如果Warp中所有线程的谓词都为真（非零），则返回1，表示所有线程都满足条件。
2. Warp Ballot Function：这个函数收集Warp内每个线程的投票结果，并生成一个位掩码，其中每个位对应一个线程的投票结果。
   1. __ballot(int predicate)：返回一个位掩码，其中第n位表示第n个线程的投票结果，1表示线程的谓词为真，0表示为假。
3. 这些函数高效地在Warp内广播每个线程的谓词值到所有线程。
4. 平台支持检测：
   1. 使用hasWarpVote设备属性或HIP_ARCH_HAS_WARP_VOTE编译器定义来检测目标平台是否支持__all和__any指令。
   2. 使用hasWarpBallot设备属性或HIP_ARCH_HAS_WARP_BALLOT编译器定义来检测目标平台是否支持__ballot指令。
5. Warp大小：与CUDA相比，HIP的__ballot函数支持64位的返回值，这允许它支持更大的Warp大小。
6. 代码移植：从CUDA移植到HIP的代码应该考虑到HIP版本支持的更大的Warp大小，并相应地调整。

11. Warp Shuffle Functions

Warp Shuffle Functions 是一类在 GPU 编程中用于在同一个 Warp（或 Wavefront）内线程之间交换数据的函数。这些函数允许线程读取其他线程的寄存器内容，从而可以执行跨线程的数据操作，如归约（reduction）、扫描（scan）或排序（sorting）。

1. 在 HIP 中，半精度浮点数（half-float，即 16 位浮点数）的 shuffle 操作是不被支持的。
2. shuffle 操作的默认宽度是 warpSize，即一个 Warp 内线程的数量。这意味着 shuffle 函数可以在这个数量范围内的线程之间交换数据。
3. 开发者在使用 shuffle 函数时，不应当假设 warpSize 一定是 32 或 64。不同的 GPU 架构可能有不同的 warp 大小，因此应该使用 HIP 运行时 API 来查询 warpSize 的实际值。
4. 与 Warp Shuffle Functions 类似，Warp Cross-Lane Functions 也允许在 Warp 内进行跨线程的操作，但它们通常用于逻辑操作如投票（vote）和选举（elect），而不是数据交换。

int __shfl(int var, int srcLane, int width=warpSize); 
float __shfl(float var, int srcLane, int width=warpSize); 
int __shfl_up(int var, unsigned int delta, int width=warpSize); 
float __shfl_up(float var, unsigned int delta, int width=warpSize); 
int __shfl_down(int var, unsigned int delta, int width=warpSize); 
float __shfl_down(float var, unsigned int delta, int width=warpSize) ; 
int __shfl_xor(int var, int laneMask, int width=warpSize);
float __shfl_xor(float var, int laneMask, int width=warpSize); 

12. Cooperative Groups Functions

Cooperative Groups 是 CUDA 和 HIP 中的一个特性，它允许开发者以新的维度组织和协调 GPU 上的线程，从而实现更细粒度的并行计算控制。

• Cooperative Groups 允许开发者将来自不同 block 的线程组合成一个逻辑上的“超级 block”，这个更大的组可以执行同步操作，就像它们是同一个 block 中的线程一样。
• 在 Cooperative Groups 中，线程可以跨越它们原本的 block 边界进行通信和同步。这意味着，如果一个线程需要等待其他线程完成某个操作，它可以跨越 block 边界等待来自同一个 Cooperative Group 内其他 block 的线程。

Cooperative Group 类型：HIP 支持几种 Cooperative Group 类型，包括：

• hipGroup
• hipDeviceGroup
• hipGridGroup

使用场景：

• Cooperative Groups 特别适用于那些需要跨多个 block 进行同步或协作的操作，如大规模归约（reduction）、排序（sorting）或搜索（searching）算法。
• 假设你有一个由多个 block 组成的 grid，每个 block 执行一部分计算任务。使用 Cooperative Groups，你可以创建一个设备组（device group），它跨越了所有这些 block。然后，你可以在这个设备组内执行一个归约操作，而不需要在 host 端进行多次内存访问和同步。

代码实现：

• 在 CUDA 和 HIP 中，Cooperative Groups API 提供了如 cooperativeGroupArray、cooperativeSubGroup 和 cooperativeGroup 等函数和类型，允许你查询和使用这些逻辑上的线程组。

性能优势：

• 使用 Cooperative Groups 可以减少 host 到 device 的同步次数，降低内存访问延迟，并可能提高内存访问的效率，因为可以在不离开 GPU 的情况下完成更多的工作。

 

13. Warp Matrix Functions

Warp Matrix Functions 是 CUDA 中的一个特性，它允许一个 warp 内的线程协同工作，以处理分布在各个线程上的小型矩阵。这些矩阵的元素在不同的线程（lanes）之间以未指定的方式分布。这种机制可以用于实现高效的并行线性代数运算，如矩阵乘法或其他涉及矩阵的计算。Warp Matrix Functions 是从 CUDA 9 开始引入的特性。

14. Independent Thread Scheduling

Independent Thread Scheduling（独立线程调度）是一种硬件特性，它在某些支持 CUDA 的 GPU 架构中被引入。这项特性允许线程独立于彼此进展，即使它们属于同一个 warp。这意味着，即使在 warp 内，线程也可以有不同的执行路径，而不是像以前那样必须执行相同的指令。

HIP（Heterogeneous-compute Interface for Portability）目前不支持这种类型的线程调度。

七、Device-Side Dynamic Global Memory Allocation

Device-side dynamic global memory allocation is under development.

八、__launch_bounds__

这是一个与__global__函数一起使用的函数属性，用于提供关于资源使用（主要是寄存器）的提示。它帮助编译器根据预期的最大线程块大小和每个执行单元（EU）的最小warp数量来优化代码。

__launch_bounds__参数：

• MAX_THREADS_PER_BLOCK：程序员保证内核将以不超过此限制的线程数启动。这允许编译器使用比无限制时更多的资源。
• MIN_WARPS_PER_EU：此参数告诉编译器最小化资源使用，以便指定数量的warp可以同时处于活动状态。它是可选的，默认值为1，如果没有指定。

使用HIP API启动内核时，它会检查指定的launch_bounds以确保内核的维度大小没有超过指定的限制。如果超出了，启动将失败，并记录错误消息，包括帮助调试的详细信息。

1. Compiler Impact

编译器使用这些参数的方式如下：

• 编译器仅使用这些提示来管理寄存器使用，并不会自动减少共享内存或其他资源。
• 如果编译器无法生成满足指定启动界限要求的内核，编译将失败。
• 从MAX_THREADS_PER_BLOCK参数中，编译器推导出可以在启动时使用的warp/block的最大数量。MAX_THREADS_PER_BLOCK的值小于默认值时，允许编译器使用更大的寄存器池：每个warp使用寄存器，这个提示将启动限制为小于最大值的warp/block大小。
• 从MIN_WARPS_PER_EU参数中，编译器推导出内核可以使用的最大寄存器数量（以满足所需的#同时活动块数）。如果MIN_WARPS_PER_EU为1，则内核可以使用多处理器支持的所有寄存器。

编译器确保内核使用的寄存器少于这两个参数允许的最大值，具体做法可能包括：

• 寄存器溢出（Spilling）：如果内核所需的寄存器数量超过了GPU的物理寄存器限制，编译器会将一些数据从寄存器移动到共享内存或全局内存中。这个过程称为寄存器溢出。溢出会增加内存访问次数，可能会降低程序的性能。

• 使用更多指令：为了减少寄存器的使用，编译器可能会生成更多的指令来执行相同的任务。例如，它可能会使用多个指令来逐步计算一个原本可以存储在单个寄存器中的值。

• 优化寄存器分配：编译器会尝试优化寄存器的使用，例如通过重用寄存器来存储不同的变量，或者通过调整代码来减少同时活跃的变量数量。

总的来说，这句话的意思是编译器需要在有限的寄存器资源和内核的性能之间找到平衡点。通过各种策略，编译器努力确保内核能够有效地使用寄存器，同时避免因资源限制而导致的性能下降。

2. Porting from CUDA __launch_bounds

CUDA 中定义的 __launch_bounds__ 属性旨在控制多处理器上的占用率（occupancy），其格式如下：

__launch_bounds(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MULTIPROCESSOR) 

这里有两个参数：

• MAX_THREADS_PER_BLOCK：每个线程块中的最大线程数。
• MIN_BLOCKS_PER_MULTIPROCESSOR：每个多处理器上期望的最小线程块数。

HIP（Heterogeneous-compute Interface for Portability）中的 __launch_bounds__ 属性格式略有不同，使用 warps 和执行单元（execution units）而不是块和多处理器：

__hip_launch_bounds__(MIN_WARPS_PER_EXECUTION_UNIT, MIN_WARPS_PER_MULTIPROCESSOR)

3. Maxregcount

在CUDA编程中，--maxrregcount 是一个编译器选项，它允许开发者指定在使用__launch_bounds__时内核可以使用的最大寄存器数量。这个选项在NVIDIA的编译器nvcc中可用，但HIP-Clang不支持这个选项。

HIP-Clang是AMD为HIP（Heterogeneous-compute Interface for Portability）提供的编译器，它旨在提供跨平台的编程接口，使得开发者可以编写能够在NVIDIA和AMD的GPU上运行的代码。由于HIP-Clang不支持--maxregcount选项，它推荐使用hip_launch_bounds指令来控制内核的资源使用。

九、Register Keyword

在C++中，register关键字曾经用于建议编译器将局部变量存储在寄存器中，以便快速访问。然而，在现代编译器中，这个关键字已经不推荐使用，并且在C++17标准中被正式弃用。在CUDA编译器nvcc和HIP-Clang中，register关键字会被静默忽略，不会对编译过程产生任何影响。

如果你的代码中使用了register关键字，并且你希望在编译时获得警告信息，可以使用以下编译器选项：

• 对于nvcc编译器，可以使用-Wdeprecated-declarations选项来启用对弃用声明的警告，这包括register关键字的使用。

• 对于HIP-Clang编译器，可以使用-Wdeprecated-register选项来专门针对register关键字的使用发出警告。

十、Pragma Unroll

#pragma unroll 是一种编译器指令，用于向编译器提供循环展开的提示。#pragma unroll 后面跟的数字表示编译器应该展开循环的迭代次数。这个数字必须是编译时已知的常量。

#pragma unroll 16 /* hint to compiler to unroll next loop by 16 */ 
for (int i=0; i<16; i++) ... 
#pragma unroll 1  /* tell compiler to never unroll the loop */ 
for (int i=0; i<16; i++) ... 
#pragma unroll /* hint to compiler to completely unroll next loop. */ 
for (int i=0; i<16; i++) ... 

十一、In-Line Assembly

内联汇编（Inline Assembly）是一种在高级语言代码中嵌入汇编语言指令的技术，允许开发者直接控制硬件层面的操作。在GPU编程中，特别是使用AMD的GCN（Graphics Core Next）架构时，内联汇编可以用来执行特定的硬件指令，优化性能。

asm volatile ("v_mac_f32_e32 %0, %2, %3" : "=v" (out[i]) : "0"(out[i]), "v" (a), "v" (in[i]));

• asm volatile：asm关键字用于声明内联汇编块，volatile关键字告诉编译器这个汇编代码可能会有副作用，因此编译器不应尝试优化或重新排序这些指令。

• "v_mac_f32_e32 %0, %2, %3"：这是GCN ISA中的一条汇编指令，v_mac_f32_e32代表一个浮点乘累加（Multiply-Add）操作，%0、%2和%3是操作数占位符，编译器将根据后面的约束代码将它们替换为实际的寄存器。

• : "=v" (out[i]) : "0"(out[i]), "v" (a), "v" (in[i])：这是内联汇编的输入输出约束列表。约束列表定义了汇编指令中使用的寄存器和变量：

  • "=v" (out[i])：指定out[i]作为输出操作数，=表示写入，v是约束代码，表示32位的虚拟通用寄存器（VGPR）。
  • "0"(out[i])：表示使用与输出相同的寄存器作为输入，0表示这是列表中的第一个约束。
  • "v" (a)和"v" (in[i])：分别指定变量a和in[i]作为输入操作数，它们也被映射到32位的VGPR。

• % followed by a position in the list of operands：在汇编指令字符串中，%后面跟着的数字表示约束列表中对应位置的操作数。

• "v" is the constraint code：v是约束代码，用于AMD GPU编程，表示目标是32位的VGPR。

• Output Constraints：输出约束，如上所示，使用=前缀，表示汇编代码将写入此操作数，然后将其作为asm表达式的返回值。

• Input constraints：输入约束没有前缀，只有约束代码。

使用内联汇编时，需要非常小心，因为它绕过了高级语言的许多安全特性，并且对硬件有直接的影响。此外，内联汇编代码通常与特定的硬件架构紧密相关，这可能会降低代码的可移植性。开发者在使用内联汇编时应确保熟悉目标硬件的指令集架构（ISA），并参考相应的硬件手册，例如AMD GCN3 ISA架构手册。

十二、C++ Support

The following C++ features are not supported:

• Run-time-type information (RTTI)

• Virtual functions

• Try/catch

十三、Kernel Compilation

内核编译是将C++/HIP（Heterogeneous-compute Interface for Portability）编写的内核编译成二进制代码对象的过程。使用AMD的HIP-Clang编译器hipcc，现在支持将内核编译成.co文件格式，即代码对象（Code Object）。.co是二进制代码对象的文件扩展名，它包含了编译后的内核代码，可以被GPU执行。

`hipcc --genco --offload-arch=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]` 
[TARGET GPU] = GPU architecture 
[INPUT FILE] = Name of the file containing kernels 
[OUTPUT FILE] = Name of the generated code object file

• hipcc：HIP-Clang编译器的命令行工具。
• --genco：指示编译器生成二进制代码对象。
• --offload-arch=[TARGET GPU]：指定目标GPU的架构。例如，--offload-arch=gfx908针对特定型号的AMD GPU。
• [INPUT FILE]：包含要编译的内核的源文件名。
• -o [OUTPUT FILE]：指定输出的二进制代码对象文件名。

优势：

• 编译成二进制代码对象可以提高应用程序的加载和启动速度，因为编译器不需要在运行时重新编译内核代码。
• 它还有助于隐藏应用程序的源代码，因为分发的是编译后的二进制形式。

使用场景：

• 当你需要将应用程序部署到多个平台或设备上，并且希望避免在每个设备上重复编译内核时，使用二进制代码对象非常有用。

十四、gfx-arch-specific-kernel

在Clang编译器中，定义了一些宏（macros），允许开发者在内核代码中根据特定的GFX（Graphics Core Next，AMD GPU架构）架构执行不同的代码分支。这些宏以__gfx*__开头，可以用来检测当前编译的GPU架构，并据此包含或排除特定代码。

以下是一些常见的__gfx*__宏示例：

• __AMDGPU__：如果编译目标是AMD GPU，这个宏会被定义。
• __gfx600__、__gfx601__、__gfx700__等：这些宏分别对应不同的GFX架构版本。例如，__gfx600__用于GFX6架构。

使用这些宏，开发者可以编写条件编译代码，以适应不同的GPU架构。例如：

#if defined(__gfx600__)
    // GFX600-specific code
#elif defined(__gfx700__)
    // GFX700-specific code
#else
    // Generic or fallback code
#endif

在HIP编程模型中，可以使用这些宏来包含特定于GFX架构的内核代码。例如，如果某个功能只存在于GFX7架构中，可以使用以下方式：

__global__ void myKernel() {
    #if defined(__gfx700__)
    // Code that uses GFX700-specific features
    #endif
}

使用这些宏的好处包括：

• 架构特定的优化：开发者可以根据特定架构的特性来优化代码，提高性能。
• 代码的可移植性：通过条件编译，同一段代码可以适应不同的GPU架构。
• 易于维护：集中管理架构特定的代码，简化了代码的维护和更新。