一些零散的和编译相关的语法/flash-attn涉及语法扫盲

rebekk
于 2024-10-09 19:20:29 发布
阅读量550
点赞数 11
分类专栏： C++ cuda编程文章标签： C++ CUDA 编译
本文链接：https://blog.csdn.net/qq_43527718/article/details/142696426
版权
C++ 同时被 2 个专栏收录
10 篇文章 0 订阅
订阅专栏
cuda编程
2 篇文章 0 订阅
订阅专栏
                    
                        
                    
                    #pragma once：一个编译指令，用于防止头文件被多次包含。当编译器遇到#pragma once时，它会确保该头文件在一个编译单元（一个.cpp文件及其包含的所有文件）中只会被包含一次。即使该文件被间接包含多次，编译器也会忽略多余的包含操作，从而加快编译速度并减少编译错误 
  #pragma once需要编译器追踪文件路径以确保文件只被包含一次，这增加了实现的复杂度
C和C++以及它们的标准早期并未包括#pragma once这种指令，所以为了保持与旧版代码的兼容性，和确保跨不同平台和编译器的代码能正常运行，编译器通常不会默认强制使用#pragma once
 
__VA_ARGS__：在宏定义中，是用来表示可变参数的预处理器标记。 
  调用BOOL_SWITCH时传递了一个lambda表达式：BOOL_SWITCH(flag, BoolConst, [&] { return some_function<BoolConst>(); }); 
BOOL_SWITCH定义如下：其中__VA_ARGS__就代表宏的可变参数，__VA_ARGS__()表示将这些可变参数当作一个可调用对象（通常是一个函数或lambda表达式）来调用#define BOOL_SWITCH(COND, CONST_NAME, ...)      \
  [&] {                                         \
    if (COND) {                                 \
      constexpr static bool CONST_NAME = true;  \
      return __VA_ARGS__();                     \
    } else {                                    \
      constexpr static bool CONST_NAME = false; \
      return __VA_ARGS__();                     \
    }                                           \
  }()
 
则调用BOOL_SWITCH后展开后的代码就是：这里__VA_ARGS__()就代表了[&] { return some_function<BoolConst>(); }这个lambda函数的调用，()的作用就是让__VA_ARGS__所代表的lambda函数立即执行[&] {
  if (flag) {
    constexpr static bool BoolConst = true;
    return [&] { return some_function<BoolConst>(); }();
  } else {
    constexpr static bool BoolConst = false;
    return [&] { return some_function<BoolConst>(); }();
  }
}()
 
 
函数模板：允许用户编写可以处理多种数据类型的通用函数，而无需为每种类型都分别编写函数实现。函数模板的基本定义形式如下：template<typename T>
void function_name(T param){
	// 函数体
}
 
  typename T表示T是一个类型参数，可以用它来代表函数中使用的任何数据类型
当调用一个函数模板时，编译器根据用户提供的参数类型自动推断出模板参数，并生成一个具体的函数实例。即，模版是一种编译时的机制，它不执行具体的代码，直到用户实际调用它时才会生成具体的代码
 
在flash-attention的flash_fwd_kernel.h文件中的combine_attn_seqk_parallel函数解释中，G老师提到了“代码中使用了大量的模板和元编程技术，使得函数能够灵活应对不同的配置参数” 
  模板：是一种让程序可以编写通用代码的机制。用户可以在编写模板时不指定具体类型，而是在使用模板时指定，这样函数代码就可以处理多种类型
元编程：允许程序在编译时而不是运行时生成或修改代码，这通常通过使用模板来实现，使得代码能够在编译时进行类型推导和其他计算。元编程是一种“编写程序来生成程序”的方法，元编程使得程序能够在编译时或运行时检查、生成或修改其他代码 
    如static_assert是C++中一种元编程技术，用于在编译时检查条件。下面这行代码检查如果kMaxSplits大于128，编译器就会在编译阶段报错static_assert(kMaxSplits <= 128, "kMaxSplits must be <= 128");
 
如下面是一个简单的使用bash脚本编写的元程序示例，同时也是一个生成式编程的例子#!/bin/bash
# metaprogram
echo '#!/bin/bash' >program
for ((I=1; I<=992; I++)) do
    echo "echo $I" >>program
done
chmod +x program
 
 
 
宏 
  是什么：是在C和C++中由预处理器处理的指令。通过#define关键字定义，允许为代码中的表达式、常量或代码片段取别名，甚至定义具有参数的宏。宏会在预处理阶段展开，在编译前会被预处理器替换为其定义的内容
举例：下面的SQUARE(x)是一个宏函数，接收参数x并将其替换为(x)*(x)#define SQUARE(x) ((x) * (x))
 
条件编译：可使用预处理器的条件控制语句（如#ifdef、ifndef），将宏用于条件编译#ifdef DEBUG
printf("Debug mode is enabled\n");
#endif 
 
缺点： 
    无类型检查：宏在预处理阶段展开，而非编译时，所以编译器无法对宏参数进行类型检查。如上面的SQUARE("test")传入一个字符串，但预处理器仍会将“test”替换到宏定义中，导致无法预测的行为
调试困难：由于宏是在预处理阶段替换的，调试时代码看起来会和原始代码不同，可能让问题更难追踪
 
宏和const的区别 
    const是编译时常量，且具有类型信息。编译器可在编译阶段进行类型检查以确保传入的值类型正确
const变量可在调试时查看其值，而宏常量会在预处理阶段被替换，调试工具很难捕捉它们
const遵循C++的作用域规则，只在声明的作用域内可见，而宏是全局的。若多个文件中定义相同名称的宏，会导致冲突和错误 
      常见作用域：【1】局部作用域（在函数、代码块等局部范围内声明的变量，只能在该范围内访问）、【2】类作用域（类的成员变量或函数，只能在该类的对象和成员函数中使用）、【3】命名空间作用域（在命名空间内声明的变量或函数，只能在同一命名空间或通过作用域解析符访问）
全局：全局作用域中的标识符可在整个程序中访问。宏定义（#define）是一种预处理器指令，作用类似于全局替换。宏一旦定义，就会在整个源文件中全局替换
原则上可以在b.cpp中include a.cpp来访问a.cpp中定义的宏，但这并不是推荐的做法（被include的.cpp文件可能会被编译多次，导致符号重复定义问题从而导致编译错误）。对于这种跨源文件使用的宏，一般会将宏定义放入一个共享的头文件（.h）中，并在各个源文件中include这个头文件
 
 
下面是flash-attention中用宏进行模板化的CUDA内核函数声明的操作例子，感觉很新颖很灵活，用户可通过这种方式在不同的配置或类型下灵活定义内核函数，而不需要手动写大量重复的代码// Use a macro to clean up kernel definitions
#define DEFINE_FLASH_FORWARD_KERNEL(kernelName, ...) \  // 这行代码定义了一个宏DEFINE_FLASH_FORWARD_KERNEL，它接受两个参数kernelName（内核函数名字）和__VA_ARGS__（这是个可变参数宏，允许传入不定数量的参数）
template<typename Kernel_traits, __VA_ARGS__> \  // 这行代码使用模板的语法，Kernel_traits常用于在CUDA中描述内核函数的特性（如线程布局、块布局等）
__global__ void kernelName(KERNEL_PARAM_MODIFIER const Flash_fwd_params params)  // 这行代码声明了一个CUDA内核函数
 
    使用：
 // 使用宏
DEFINE_FLASH_FORWARD_KERNEL(MyKernel, typename T) 

// 上面的宏展开后会变成：
template<typename Kernel_traits, __VA_ARGS__>
__global__ void MyKernel(KERNEL_PARAM_MODIFIER const Flash_fwd_params params) 
// 这就定义了一个名为MyKernel的CUDA内核函数，它接受模板参数Kernel_traits和T，并可以在调用时传入Flash_fwd_params结构体
 
 
编译和调试 
  编译是把源代码编程二进制obj的过程（链接后成为可执行文件），会检查有无简单的语法问题（要不然编译器不认识）
调试的话，先要提前生成二进制代码，所以需要先进行编译和链接，然后到断点后，调试器会帮你加int3中断，就停住了。
调试是在程序运行后，根据运行状况来检查错误，是对已经存在的二进制文件进行调试，目的在于查找性能瓶颈和跟踪软件bug；编译器是在程序没有运行的时候帮你检查错误，目的在于把代码编译（再汇编，这里不太严谨，具体可看这篇）成二进制文件，即可执行的程序
 
vscode每次调试都要重新编译项目，这是因为在项目运行时更改了系统时间，导致vs编译日志认为文件需要重新编译。解决办法见这篇
 
pybind11库： 
  是什么：是一个轻量级的头文件库，用于在C++和python之间创建绑定。它允许用户将C++函数、类等轻松暴露给python，使得python可以调用C++代码，而不用编写繁琐的python C API代码。
使用说明：在CPP文件中使用pybind11将C++函数绑定为python函数；然后在py文件中import刚才绑定的库，并使用
示例： 
    C++文件：
 #include <pybind11/pybind11.h>

int add(int i, int j){
	return i+j;
}

PYBIND11_MODULE(example, m){
	m.def("add", &add, "A function which adds two numbers");
}
 
    python中使用：
 import example
result = example.add(2, 3)
print(result)
 
 
内联函数inline： 
  是什么：在编译阶段，编译器会尝试将内联函数的调用语句替换成函数体本身，从而避免实际的函数调用（参数传递、函数跳转、函数返回），如下例编译器会将square(5)替换成5*5，这样就避免了函数调用的过程inline int square(int x) {
    return x * x;
}

int main() {
    int result = square(5);
}
 
使用场景：内联函数通常用于一些短小、性能敏感、执行频繁的函数，以避免重复进行函数调用的开销。模板函数通常会被定义为内联函数，它们的实例化会在编译期展开
代码膨胀：内联函数会导致代码膨胀，因为函数体被复制到被一个调用点，这会增加可执行文件的体积
 
extern、__shared__关键字 
  在C/C++中，extern关键字用来声明一个全局变量或函数，表示这个变量或函数是在其他文件或作用域中定义的
在CUDA中，当extern与__shared__关键字一起使用时，表示动态分配的共享内存，所以shared memory的大小是在运行时确定的而不是编译时（shared memory不是固定的？）。在调用CUDA内核时，可通过第三个参数<<<gridDim, blockDim, sharedMemSize>>>(...);来指定共享内存的大小。如下的代码，在调用内核时制定了共享内存大小，然后smem_[]就会在运行时变成大小为1024字节的数组__gloval__ void myKernel(...){
	extern __shared__ char smem_[];  // 声明动态shared memory
	// 使用shared memory
}

size_t sharedMemSize = 1024;
myKernel<<<gridDim, blockDim, sharedMemSize>>>(...);
 
__shared__关键字在CUDA中表示共享内存，供同一block（=SM，含多个warp）内的所有线程共享访问
 
constexpr关键字：
 是C++11引入的关键字，表示常量表达式（constant expression），作用是让变量或函数的值在编译时就能计算完成并确定其值，而不是在运行时。在性能敏感的代码中，constexpr是一个非常有效的工具 
  constexpr和const的区别：前者要求值在编译时计算得出；后者表示变量一旦初始化后就不能改变，但它不一定在编译时计算
 
using语法
 using是C++中的一种类型别名声明语法，通常用于为复杂的类型创建别名，使代码更加简洁和可读。用法就是：using 新类型名 = 原类型名;。如using Element = elem_type;的作用就是为elem_type创建一个别名，这个新的别名叫做Element。使用了这个using后，代码中的Element就相当于elem_type。

                