CUDA C++ Programming Guide

【CUDA编程】【29】【6.C++ Language Extensions】【Part4】

Address Space Conversion Functions，Alloca Function，Compiler Optimization Hint Functions，Warp Vote Functions，Warp Match Functions

Posted by x-jeff on January 23, 2025

【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。

1.Address Space Conversion Functions

本部分都是用于地址空间转换的函数。一些预备知识点：

  • CUDA编程中，通用地址是一种抽象的地址形式，可能指向device内存中的不同地址空间。
  • PTX全称是Parallel Thread Execution，是CUDA中的底层指令集。

1.1.__cvta_generic_to_global()

1

__device__ size_t __cvta_generic_to_global(const void *ptr);

该函数会执行PTX指令cvta.to.global，将输入的通用地址的指针ptr转换为全局内存的地址并返回。

1.2.__cvta_generic_to_shared()

1

__device__ size_t __cvta_generic_to_shared(const void *ptr);

该函数会执行PTX指令cvta.to.shared，将输入的通用地址的指针ptr转换为共享内存的地址并返回。

1.3.__cvta_generic_to_constant()

1

__device__ size_t __cvta_generic_to_constant(const void *ptr);

该函数会执行PTX指令cvta.to.const，将输入的通用地址的指针ptr转换为常量内存的地址并返回。

1.4.__cvta_generic_to_local()

1

__device__ size_t __cvta_generic_to_local(const void *ptr);

该函数会执行PTX指令cvta.to.local，将输入的通用地址的指针ptr转换为本地内存的地址并返回。

1.5.__cvta_global_to_generic()

1

__device__ void * __cvta_global_to_generic(size_t rawbits);

该函数会执行PTX指令cvta.global，将输入的全局内存地址rawbits转换为通用地址的指针并返回。

1.6.__cvta_shared_to_generic()

1

__device__ void * __cvta_shared_to_generic(size_t rawbits);

该函数会执行PTX指令cvta.shared，将输入的共享内存地址rawbits转换为通用地址的指针并返回。

1.7.__cvta_constant_to_generic()

1

__device__ void * __cvta_constant_to_generic(size_t rawbits);

该函数会执行PTX指令cvta.const，将输入的常量内存地址rawbits转换为通用地址的指针并返回。

1.8.__cvta_local_to_generic()

1

__device__ void * __cvta_local_to_generic(size_t rawbits);

该函数会执行PTX指令cvta.local，将输入的本地内存地址rawbits转换为通用地址的指针并返回。

2.Alloca Function

2.1.Synopsis

1

__host__ __device__ void * alloca(size_t size);

2.2.Description

对栈（stack）与堆（heap）的解释。

对栈帧（stack frame）的解释：栈帧是栈中的一个部分，每当函数被调用时，系统会在栈上为这个函数创建一个栈帧。栈帧包含了该函数的局部变量、函数调用的参数、返回地址等信息。每个栈帧对应一个函数调用，当函数返回时，相应的栈帧会被销毁。栈帧的管理是编译器和运行时环境的责任。

函数通过调用alloca()在其栈帧上分配size字节的内存。alloca()的返回值是一个指向已分配内存的指针，当从device代码调用时，分配的内存的起始地址是16字节对齐的。分配的内存在调用alloca()的函数返回时会自动释放。

注意：

在Windows平台上，使用alloca()前必须包含<malloc.h>头文件。使用alloca()可能会导致栈溢出，因此用户需要根据需求调整栈大小。

仅支持计算能力在5.2及以上的device。

2.3.Example

1
2
3
4
5

__device__ void foo(unsigned int num) {
    int4 *ptr = (int4 *)alloca(num * sizeof(int4));
    // use of ptr
    ...
}

3.Compiler Optimization Hint Functions

本部分的函数可以用于向编译器优化器提供附加信息。

3.1.__builtin_assume_aligned()

1

void * __builtin_assume_aligned (const void *exp, size_t align)

编译器会假设参数指针对齐到至少align字节，并返回参数指针。

举个例子：

1
2

void *res = __builtin_assume_aligned(ptr, 32); // compiler can assume 'res' is
                                               // at least 32-byte aligned

上述例子中，编译器可以假设res至少是32字节对齐的。注意，res和ptr实际上是相同的地址，只是编译器会基于开发者的提示假设res是至少32字节对齐的，因此它可能优化生成的代码（比如不再进行对齐检查），__builtin_assume_aligned()并不会修改内存对齐，也不会检查或调整指针的实际对齐方式，如果实际对齐方式和假设的对齐方式不一致，则可能会引发未定义的行为。

三参数版本：

1
2

void * __builtin_assume_aligned (const void *exp, size_t align,
                                 <integral type> offset)

编译器会假设(char *)exp - offset对齐到至少align字节，并返回参数指针。

举个例子：

1
2
3

void *res = __builtin_assume_aligned(ptr, 32, 8); // compiler can assume
                                                  // '(char *)res - 8' is
                                                  // at least 32-byte aligned.

上述例子中，编译器假设(char *)res - 8至少是32字节对齐的。

3.2.__builtin_assume()

1

void __builtin_assume(bool exp)

__builtin_assume()是一个编译器优化提示函数，用于向编译器显式声明一个条件始终为真。这种声明可以帮助编译器在优化代码时减少不必要的检查或生成更高效的代码。编译器会假定exp表达式总是为真，因此可以省略基于exp的任何检查或分支代码。如果exp在运行时计算结果为假，程序行为是未定义的。未定义行为的后果可能包括崩溃、数据错误或其他不可预测的问题。如果exp有副作用（side effects），例如修改变量值或调用函数，编译器可能会忽略这些副作用，导致程序行为异常，比如下面的例子：

1

__builtin_assume(x++ < 10); //错误用法，x++有副作用

一个正确的例子：

1
2
3
4

 __device__ int get(int *ptr, int idx) {
   __builtin_assume(idx <= 2);
   return ptr[idx];
}

3.3.__assume()

1

void __assume(bool exp)

__assume()的作用和__builtin_assume()是一样的。区别在于，__builtin_assume()是GCC和Clang编译器的内置函数，通常适用于跨平台；__assume()是MSVC（微软编译器）的内置函数，主要用于Windows平台。

一个例子：

1
2
3
4

 __device__ int get(int *ptr, int idx) {
   __assume(idx <= 2);
   return ptr[idx];
}

3.4.__builtin_expect()

1

long __builtin_expect (long exp, long c)

__builtin_expect()会提示编译器，表达式exp更有可能等于c，并返回exp的值。该函数通常用于向编译器提供分支预测信息，开发者通过__builtin_expect()告诉编译器，某个条件更可能发生，从而调整代码的布局或优化跳转指令，这可以帮助编译器生成更高效的代码。

一个例子：

1
2
3
4
5

// indicate to the compiler that likely "var == 0",
// so the body of the if-block is unlikely to be
// executed at run time.
if (__builtin_expect (var, 0))
  doit ();

3.5.__builtin_unreachable()

1

void __builtin_unreachable(void)

__builtin_unreachable()向编译器指示程序的控制流永远不会到达调用该函数的代码点。如果运行时控制流实际到达此代码点，则程序的行为是未定义的。

一个例子：

1
2
3
4
5
6

// indicates to the compiler that the default case label is never reached.
switch (in) {
case 1: return 4;
case 2: return 10;
default: __builtin_unreachable();
}

3.6.Restrictions

__assume()仅在使用cl.exe（Microsoft编译器）作为host编译器时支持。其他函数在所有平台上均受支持，但需遵守以下限制：

  • 如果host编译器支持该函数，则可以在翻译单元（translation unit）中的任何位置调用该函数。
  • 否则，仅当以下条件满足时，才能调用该函数：
    • 在__device__/__global__函数体内调用。
    • 或者，仅当宏__CUDA_ARCH__被定义时调用。

这里简单解释下翻译单元的生成过程：

  1. 编译器读取源文件。
    • 以.c或.cpp文件为起点。
  2. 预处理阶段。
    • 处理所有的#include指令，将头文件展开。
    • 解析宏定义和替换（例如#define）。
    • 处理条件编译指令（如#ifdef和#endif）。
    • 删除注释。
  3. 形成翻译单元。
    • 在预处理完成后，所有内容合并成一个统一的代码块，这个结果就是翻译单元。
  4. 翻译单元传递到编译器。
    • 翻译单元会被编译器处理，生成目标文件（如.o或.obj）。

4.Warp Vote Functions

1
2
3
4

int __all_sync(unsigned mask, int predicate);
int __any_sync(unsigned mask, int predicate);
unsigned __ballot_sync(unsigned mask, int predicate);
unsigned __activemask();

__any、__all和__ballot在CUDA 9.0中被废弃。

在计算能力7.x及以上的device中，__any、__all和__ballot不再可用，必须使用带有_sync后缀的变体。

__all_sync(unsigned mask, predicate)：

当且仅当所有活跃线程（由mask指定）的predicate都不为零时，函数返回非零值。

__any_sync(unsigned mask, predicate)：

只要有任意一个活跃线程的predicate不为零，函数就返回非零值。

__ballot_sync(unsigned mask, predicate)：

函数返回一个32比特的整数，每个比特位对应warp中的一个线程。如果第N个线程是活跃的且predicate不为零，则第N个比特位就是1，否则为0。

__activemask()：

函数返回一个32比特的整数，每个比特位对应warp中的一个线程。如果第N个线程在函数调用时是活跃的，则第N个比特位就是1，否则为0。

以上这些内置函数不保证存在内存屏障。

5.Warp Match Functions

__match_any_sync和__match_all_sync只支持计算能力在7.x及以上的device。

5.1.Synopsis

1
2

unsigned int __match_any_sync(unsigned mask, T value);
unsigned int __match_all_sync(unsigned mask, T value, int *pred);

类型T可以是int、unsigned int、long、unsigned long、long long、unsigned long long、float或double。

5.2.Description

__match_sync()的作用范围为一个warp内。

__match_any_sync的参数mask用于指定warp中的哪些线程参与比较，其返回值也是一个mask，用于指定对于同一个变量，哪些线程的变量值与当前线程的变量值value是相同的。

__match_all_sync的参数mask用于指定warp中的哪些线程参与比较，只有当所有参与比较的线程的变量值都与当前线程的变量值value相同时，才会返回一个mask，用于指定都有哪些线程有与当前线程相同的变量值，其实此时返回的mask和参数mask是一样的（此时参数pred被置为true）。如果不是所有参与比较的线程都与当前线程有着相同的变量值，则函数返回0（此时参数pred被置为false）。

以上这些内置函数不保证存在内存屏障。

CATALOG
    FEATURED TAGS
    README Machine Learning Series Mathematics Series Image Segmentation Python Series Fundamentals of Computers OpenCV Series C++ Series Deep Learning Series Conda Tensorflow Series Algorithms Series AI Papers Object Recognition Natural Language Processing Aha! Algorithms Camera Calibration How Program Works Optical Flow Estimation How Networks Work Computer Vision Source Code CUDA C++ Programming Guide C++ Concurrency IN ACTION