CUDA编程中 extern "c"用法解析

本文介绍了在CUDA编程中extern 'C'的关键作用,解释了它如何指示编译器按照C语言的方式处理函数,以解决C++与C语言混合编程时的链接问题。通过举例说明了未加extern 'C'声明时的链接错误,以及添加后的正确编译和链接方式,并探讨了extern 'C'的含义和extern类型的作用。
部署运行你感兴趣的模型镜像

简介

  extern “c”并不是什么新鲜的东西,但是最在编译CUDA程序时候出现了“error LNK2019,未定义的符号”的问题。这不用说了,链接器说找不到,那就是找不到咯。但是why?我的源码中明明已经有了定义。
后来发现是extern “c”的问题。

extern “C”

  C++保留了一部分过程式语言的特点,因而它可以定义不属于任何类的全局变量和函数。但是,C++毕竟是一种面向对象的程序设计语言,为了支持函数的重载,C++对全局函数的处理方式与C有明显的不同。
extern “C”的主要作用就是为了能够正确实现C++代码调用其他C语言代码。加上extern “C”后,会指示编译器这部分代码按C语言的进行编译,而不是C++的。这很重要!由于C++支持函数重载,因此编译器编译函数的过程中会将函数的参数类型也加到编译后的代码中,而不仅仅是函数名;而C语言并不支持函数重载,因此编译C语言代码的函数时不会带上函数的参数类型,一般之包括函数名。
  如果两份代码,一份以C标准编译,一份以C++标准编译,那么链接时候就会出现error LNK2019,未定义的符号”的问题。因为链接时,C版本和C++版本的实际签名不一样。同样的,你用C 开发了一个DLL 库,为了能够让C ++语言也能够调用你的DLL输出(Export)的函数,你需要用extern “C”来强制编译器不要修改你的函数名。

extern “C”真正含义

  请注意理解extern “C”的真正含义,extern “C” 包含两层意思:

  • 被它修饰的目标是“extern”类型的;
  • 其次,被它修饰的目标是“C-type”的。

因为C++通常被认为是C的超集,导致我们在理解时候有误,认为C的代码一定可以用C++编译器编译。现在看来,还是有很多误区。下面详细讲讲extern “C” 包含的两层意思。

extern类型

  extern是C/C++语言中表明函数和全局变量作用范围(可见性)的关键字,该关键字告诉编译器,其声明的函数和变量可以在本模块或其它模块中使用。
  通常,在模块的头文件中对本模块提供给其它模块引用的函数和全局变量以关键字extern声明。例如,如果模块B欲引用该模块A中定义的全局变量和函数时只需包含模块A的头文件即可。这样,模块B中调用模块A中的函数时,在编译阶段,模块B虽然找不到该函数,但是并不会报错;它会在链接阶段中从模块A编译生成的目标代码中找到此函数。
  与extern对应的关键字是static,被它修饰的全局变量和函数只能在本模块中使用。因此,一个函数或变量只可能被本模块使用时,其不可能被extern “C”修饰。

extern “C”是按照C语言方式编译和链接的

  首先看看C++中对类似C的函数是怎样编译的。作为一种面向对象的语言,C++支持函数重载,而过程式语言C则不支持。函数被C++编译后在符号库中的名字与C语言的不同。例如,假设某个函数的原型为:

void foo( int x, int y );

  该函数被C编译器编译后在符号库中的名字为_foo,而C++编译器则会产生像_foo_int_int之类的名字(不同的编译器可能生成的名字不同,但是都采用了相同的机制,生成的新名字称为“mangled name”)。
  _foo_int_int这样的名字包含了函数名、函数参数数量及类型信息,C++就是靠这种机制来实现函数重载的。 例如,在C++中,函数void foo( int x, int y )与void foo( int x, float y )编译生成的符号是不相同的,后者为_foo_int_float。
   同样地,C++中的变量除支持局部变量外,还支持类成员变量和全局变量。用户所编写程序的类成员变量可能与全局变量同名,我们以”.”来区分。而本质上,编译器在进行编译时,与函数的处理相似,也为类中的变量取了一个独一无二的名字,这个名字与用户程序中同名的全局变量名字不同。

举例说明

未加extern “C”声明时的连接方式

假设在C++中,模块A的头文件如下:

// 模块A头文件 moduleA.h
#ifndef MODULE_A_H
#define MODULE_A_H
int foo( int x, int y );
#endif

在模块B中引用该函数:

// 模块B实现文件 moduleB.cpp
#include "moduleA.h"
foo(2,3);

链接阶段,链接器会从模块A生成的目标文件moduleA.obj中寻找_foo_int_int这样的符号!

加extern “C”声明后的编译和链接方式

   加extern “C”声明后,模块A的头文件变为:

// 模块A头文件 moduleA.h
#ifndef MODULE_A_H
#define MODULE_A_H
extern "C" int foo( int x, int y );
#endif

在模块B的实现文件中仍然调用foo( 2,3 ),其结果是:

  1. A编译生成foo的目标代码时,没有对其名字进行特殊处理,采用了C语言的方式;

  2. 链接器在为模块B的目标代码寻找foo(2,3)调用时,寻找的是未经修改的符号名_foo。

如果在模块A中函数声明了foo为extern “C”类型,而模块B中包含的是extern int foo(int x, int y),则模块B找不到模块A中的函数;反之亦然。

extern “C”这个声明的真实目的是为了实现C++与C及其它语言的混合编程。

应用

  • C++代码调用C语言代码、在C++的头文件中使用

在C++中引用C语言中的函数和变量,在包含C语言头文件(假设为cExample.h)时,需进行下列处理:

extern "C"
{
#include "cExample.h"
}

而在C语言的头文件中,对其外部函数只能指定为extern类型,C语言中不支持extern “C”声明,在.c文件中包含了extern “C”时会出现编译语法错误。

/* c语言头文件:cExample.h */
#ifndef C_EXAMPLE_H
#define C_EXAMPLE_H
extern int add(int x,int y);     //注:写成extern "C" int add(int , int ); 也可以
#endif
/* c语言实现文件:cExample.c */
#include "cExample.h"
int add( int x, int y )
{
 return x + y;
}
// c++实现文件,调用add:cppFile.cpp
extern "C"
{
 #include "cExample.h"        //注:此处不妥,如果这样编译通不过,换成 extern "C" int add(int , int ); 可以通过
}
int main(int argc, char* argv[])
{
 add(2,3);
 return 0;
}
  • 如果C++调用一个C语言编写的.DLL时,当包括.DLL的头文件或声明接口函数时,应加extern “C”{}。

在C中引用C++语言中的函数和变量时,C++的头文件需添加extern “C”,但是在C语言中不能直接引用声明了extern “C”的该头文件,应该仅将C文件中将C++中定义的extern “C”函数声明为extern类型

//C++头文件 cppExample.h
#ifndef CPP_EXAMPLE_H
#define CPP_EXAMPLE_H
extern "C" int add( int x, int y );
#endif
//C++实现文件 cppExample.cpp
#include "cppExample.h"
int add( int x, int y )
{
 return x + y;
}
/* C实现文件 cFile.c
/* 这样会编译出错:#include "cExample.h" */
extern int add( int x, int y );
int main( int argc, char* argv[] )
{
 add( 2, 3 );
 return 0;
}
  • C++文件中引用.cu文件中定义的函数

把被调用的函数在.cu文件中用extern “C” 包含起来。并且在调用文件.cpp中用extern “C”声明该函数,然后调用。

您可能感兴趣的与本文相关的镜像

PyTorch 2.5

PyTorch 2.5

PyTorch
Cuda

PyTorch 是一个开源的 Python 机器学习库,基于 Torch 库,底层由 C++ 实现,应用于人工智能领域,如计算机视觉和自然语言处理

<think>好的,我现在需要解释CUDA中的规约(Reduce)操作。首先,我得回忆一下规约的基本概念。规约是一种将数据集通过某种操作(比如加法、最大值、最小值等)合并成单个结果的过程。在CUDA中,高效实现规约对性能至关重要,尤其是在处理大规模数据时。 用户可能想了解CUDA规约的实现原理和优化方法。我需要从基础开始,逐步深入。首先,应该介绍并行规约的基本思想:将数据分成块,每个线程块处理一部分数据,通过逐步合并部分结果来得到最终结果。 然后,应该提到线程块内的规约。使用共享内存来加速数据访问,因为共享内存的访问速度比全局内存快得多。这里可能会涉及线程同步的问题,需要用__syncthreads()来确保所有线程完成数据加载后再进行计算。 接下来,需要解释不同的优化策略。比如顺序访问与交错访问的区别,交错访问可以减少共享内存的bank冲突。Bank冲突会导致多个线程同时访问同一个bank,从而降低性能,因此优化访问模式很重要。 可能还需要提到循环展开,减少分支预测的开销。循环展开可以通过减少循环次数来提高指令级并行度,从而加速计算。 然后,要讨论多级规约。当数据量超过一个线程块能处理的范围时,需要在全局内存中进行多级规约。第一级每个线程块生成一个部分结果,第二级再对这些部分结果进行规约,直到得到最终结果。 另外,需要考虑线程块的数量和线程数的选择。线程块数量通常与SM(流多处理器)的数量相关,以充分利用GPU的并行计算能力。每个线程块中的线程数一般是256或512,这取决于GPU架构和共享内存的大小。 可能还需要提到原子操作,但要注意原子操作的效率较低,通常只在最后一步不得已时使用。例如,在全局内存中对部分结果进行原子加操作,但应尽量避免频繁使用。 最后,可以给出一个简单的代码示例,展示如何实现一个加法规约。示例中应包括内核函数的实现,共享内存的使用,以及逐步规约的过程。同时,指出代码中的优化点,比如避免bank冲突和循环展开。 需要注意的是,确保LaTeX公式的正确使用。比如,提到时间复杂度时,使用$$O(\log n)$$来表示对数复杂度。行内公式如每个线程处理$n$个元素,需要用$...$包裹。 可能用户还会关心性能优化的具体数据,比如使用优化后的规约比简单实现快多少倍。这时候可以提到一些实际测试结果,比如优化后的速度提升可达10倍以上,但具体数值取决于GPU架构和数据规模。 此外,需要提醒用户在实际编程中要注意共享内存的大小限制,以及线程块和网格的配置参数。不同的GPU架构可能有不同的最优配置,需要根据具体情况调整。 总结一下,解释CUDA规约的步骤应该是:基本概念、线程块内规约、共享内存使用、优化策略(访问模式、循环展开)、多级规约、线程配置和注意事项。最后通过代码示例和性能考量来巩固理解。</think>在CUDA中,**规约(Reduce)** 是一种将大量数据通过特定操作(如求和、求最大值等)合并为单个结果的并行算法。以下是逐步解析其核心原理和优化方法: --- ### **1. 基本思想** - **目标**:将数组元素通过二元操作(如加法)合并为一个值。 - **并行策略**:将数据划分为块,每个线程块处理局部规约,再对局部结果进行全局规约。 - **时间复杂度**:理想情况下为 $$O(\log n)$$,优于串行的 $$O(n)$$。 --- ### **2. 线程块内规约** #### **(1) 共享内存加速** - **步骤**: 1. 每个线程块将全局数据加载到共享内存。 2. 使用树状规约逐步合并数据(如二分法相加)。 3. 最终结果由线程块的第一个线程写入全局内存。 - **代码片段**: ```cpp __global__ void reduce_kernel(float *d_in, float *d_out) { extern __shared__ float s_data[]; int tid = threadIdx.x; int i = blockIdx.x * blockDim.x + threadIdx.x; s_data[tid] = d_in[i]; __syncthreads(); for (int s = blockDim.x/2; s > 0; s >>= 1) { if (tid < s) { s_data[tid] += s_data[tid + s]; } __syncthreads(); } if (tid == 0) d_out[blockIdx.x] = s_data[0]; } ``` #### **(2) 避免共享内存Bank冲突** - **问题**:连续线程访问同一共享内存bank会导致串行化。 - **优化**:使用交错访问(Strided Access)或调整步长。 --- ### **3. 多级规约** - **场景**:当数据规模超过单线程块容量时。 - **步骤**: 1. 第一级规约:每个线程块生成一个局部结果。 2. 第二级规约:对局部结果再次调用规约内核。 3. 重复直到得到最终结果。 --- ### **4. 优化策略** #### **(1) 循环展开(Unrolling)** - **作用**:减少循环开销,提高指令级并行度。 - **示例**:手动展开循环以减少条件判断。 #### **(2) 向量化访存** - **方法**:使用`float4`或`int4`类型一次性读取多个数据。 #### **(3) 原子操作(慎用)** - **适用场景**:最终合并少量局部结果时。 - **注意**:全局原子操作(如`atomicAdd`)性能较差,尽量避免高频使用。 --- ### **5. 性能影响因素** - **线程块大小**:通常选择256或512线程,需适配共享内存容量。 - **网格大小**:与GPU流多处理器(SM)数量匹配,确保满载。 - **数据对齐**:确保全局内存访问对齐,提高吞吐量。 --- ### **6. 代码示例(优化版)** ```cpp template <unsigned blockSize> __device__ void warpReduce(volatile float* s_data, int tid) { if (blockSize >= 64) s_data[tid] += s_data[tid + 32]; if (blockSize >= 32) s_data[tid] += s_data[tid + 16]; if (blockSize >= 16) s_data[tid] += s_data[tid + 8]; if (blockSize >= 8) s_data[tid] += s_data[tid + 4]; if (blockSize >= 4) s_data[tid] += s_data[tid + 2]; if (blockSize >= 2) s_data[tid] += s_data[tid + 1]; } __global__ void optimized_reduce(float *d_in, float *d_out) { extern __shared__ float s_data[]; int tid = threadIdx.x; int i = blockIdx.x * (blockDim.x * 2) + threadIdx.x; s_data[tid] = d_in[i] + d_in[i + blockDim.x]; __syncthreads(); if (blockSize >= 512) { if (tid < 256) s_data[tid] += s_data[tid + 256]; __syncthreads(); } if (blockSize >= 256) { if (tid < 128) s_data[tid] += s_data[tid + 128]; __syncthreads(); } // ... 类似处理64,32,16等 if (tid < 32) warpReduce<blockSize>(s_data, tid); if (tid == 0) d_out[blockIdx.x] = s_data[0]; } ``` --- ### **7. 性能对比** - **朴素规约**:约5-10 GB/s。 - **优化后规约**:可达50-100 GB/s(取决于GPU架构)。 --- ### **总结** CUDA规约的关键在于: 1. 利用共享内存减少全局访问。 2. 通过树状规约降低计算复杂度。 3. 优化内存访问模式(如避免bank冲突)。 4. 合理配置线程块和网格大小。 实际开发中需结合`nvprof`工具分析性能瓶颈,并针对特定硬件调整参数。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值