1.第一个CUDA程序
1 #include <iostream>
2
3 __global__ void kernel(void) { //__global__告知编译器函数kernel用设备代码编辑器
4 }
5
6 int main() { //默认主机编译
7 kernel << <1, 1 >> > ();
8 printf("HelloWorld");
9 return 0;
10 }CUDA提供与C在语言级别上集成，在主机代码中调用设备代码
尖括号内参数用来确定运行时如何启动设备代码
2.关键词
1 #include <iostream>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4 //#include <book.h>
5
6 __global__ void add(int a, int b, int *c) {
7 *c = a + b;
8 }
9
10 int main() {
11 int c;
12 int *dev_c;
13 cudaMalloc((void**)&dev_c, sizeof(int));
14
15 add <<<1, 1 >>> (2, 7, dev_c);
16
17 cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);
18 printf("2 + 7=%d\\n", c);
19
20 cudaFree(dev_c);
21 return 0;
22 }
- 像调用C函数一样将参数传递给核函数
- 设备执行操作时需要分配内存
使用cudaMalloc()分配内存
作用：使CUDA在运行时在设备上分配内存
cudaMalloc((void**)&dev_c, sizeof(int)
- 第一个参数是指针，用来
保存新分配内存地址变量 - 第二个参数是分配内存的大小
- 返回类型为void*
不能在主机代码中对cudaMalloc()返回的指针进行解引用(Dereference)。
主机代码可以将这个指针作为参数传递，对其进行算术运算，转换为另一种不同类型，但是不可以使用这个指针来进行读取或者写入内存
设备指针使用限制：
- 可以将
cudaMalloc()分配的指针传递给设备上执行的函数 - 可以将
cudaMalloc()分配的指针传递给在主机上执行的函数 - 可以在设备代码中使用
cudaMalloc()分配的指针进行内存读/写 - 不能在主机代码中使用
cudaMalloc()分配的指针进行内存读/写
不能用标准C的free()释放cudaMalloc()分配的内存，需要调用cudaFree()
主机上不能对设备上的内存做任何修改
访问设备内存两种方法
- 在设备代码中使用设备指针
- 主机指针只能访问主机代码中的内存
- 设备指针只能访问设备代码中的内存
- 主机调用
cudaMemcpy()
cudaMemcpy()
类似标准C中的memcpy()，多了一个指定设备内存指针(源指针/目标指针)的参数
void *memcpy(void *dest, const void *src, size_t n); 由src指向地址为起始地址的连续n个字节的数据复制到以destin指向地址为起始地址的空间内。 #include<string.h> 函数返回一个指向dest的指针。
cudaMemcpyDeviceToHostcudaMemcpyHostToDevice-
cudaMemcpyDeviceToDevice告诉运行时两个指针均位于设备上 - 若源指针和目标指针均位于主机上，可以直接调用memcpy()函数
3.查询设备信息
调用cudaGetDeviceCount,返回结构参数如图：
代码
1 #include <iostream>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4 int main() {
5 cudaDeviceProp prop;
6
7 int count;
8 cudaGetDeviceCount(&count);
9 for (int i = 0;i < count;i++) {
10 cudaGetDeviceProperties(&prop, i);
11
12 printf(" ---General Information for Device %d---\\n", i);
13 printf("Name: %s\\n", prop.name);
14 printf("Compute capability: %d.%d\\n", prop.major,prop.minor);
15 printf("Clock rate: %d\\n", prop.clockRate);
16 printf("Device copy overlap: ");
17 if (prop.deviceOverlap)
18 printf("Enabled\\n");
19 else
20 printf("Disabled\\n");
21 printf("Kernel execition timeout : ");
22 if (prop.kernelExecTimeoutEnabled)
23 printf("enabled\\n");
24 else
25 printf("Disabled\\n");
26 printf("\\n");
27
28 printf("---Memory Information for device %d---\\n", i);
29 printf("Total global Mem:%ld\\n", prop.totalGlobalMem);
30 printf("Total constant Mem:%ld\\n", prop.totalConstMem);
31 printf("Max mem pitch:%ld\\n", prop.memPitch);
32 printf("Texture Alignment:%ld\\n", prop.textureAlignment);
33 printf("\\n");
34
35 printf("---MP Information for device %d---\\n", i);
36 printf("Multiprocessor count :%d\\n", prop.multiProcessorCount);
37 printf("Shared mem per mp:%ld\\n", prop.sharedMemPerBlock);
38 printf("Registers per mp: %d\\n", prop.regsPerBlock);
39 printf("Threads in warp: %d\\n", prop.warpSize);
40 printf("Max threads per block: %d\\n", prop.maxThreadsPerBlock);
41 printf("Max thread dimensions:(%d, %d, %d)\\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
42 printf("Max grid dimensions:(%d, %d, %d)\\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
43 printf("\\n");
44
45 }
46 }
自动寻找设备
- 将目标属性填充到cudaDeviceProp结构
cudaDeviceProp prop;
memset(&prop,0,sizeof(cudaDeviceProp));
prop.major=1;
prop.minor=3;- 将其传递给cudaChooseDevice()
- cudaChooseDevice()返回满足条件的设备ID
- 将ID传递给cudaSetDevice()，之后所有操作在此设备上进行
完整程序
1 #include <iostream>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4 int main() {
5 cudaDeviceProp prop;
6 int dev;
7
8 cudaGetDevice(&dev);
9 printf("ID of current CUDA device: %d\\n", dev);
10
11 memset(&prop, 0, sizeof(cudaDeviceProp));
12 prop.major = 1;
13 prop.minor = 3;
14 cudaChooseDevice(&dev, &prop);
15 printf("ID of CUDA device closest to reviaion 1.3: %d\\n", dev);
16 cudaSetDevice(dev);
17 }
设备使用
速度快->多核处理器的GPU
核函数与CPU有密集交互->在集成的GPU上运行代码，因为其可与CPU共享内存
NVIDIA的SLI（Scalable Link Interface，可伸缩链路接口）技术使得多个独立的GPU可以并排排列。
无论是哪种情况，应用程序都可以从多个GPU中选择最适合的GPU。
如果应用程序依赖于GPU的某些特定属性，或者需要在系统中最快的GPU上运行，此API有帮助，因为CUDA运行时本身并不能保证为应用程序选择最优或者最合适的GPU。
小结
CUDA C/C++只是对标准C/C++进行了语言级扩展，利用修改符指定代码在主机或设备上运行。
__global__指明函数在GPU上运行
使用GPU上内存，通过与C相关API对应的CUDA的API
4.CUDA C并行编程
GPU计算应用前景取决于能否从问题中发掘出大规模并行性
书籍P29，对CPU上并行进行了否定
1 #include <iostream>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4
5 #define N 10000
6
7 __global__ void add(int *a, int *b, int *c) {
8 int tid = blockIdx.x; //计算位于此索引处的数据
9 if (tid < N)
10 c[tid] = a[tid] + b[tid];
11 }
12
13 int main() {
14 int a[N], b[N], c[N];
15 int *dev_a, *dev_b, *dev_c;
16
17 //GPU上分配内存
18 cudaMalloc((void**)&dev_a, N * sizeof(int));
19 cudaMalloc((void**)&dev_b, N * sizeof(int));
20 cudaMalloc((void**)&dev_c, N * sizeof(int));
21
22 //对数组a,b赋值
23 for (int i = 0;i < N;i++) {
24 a[i] = -i;
25 b[i] = i*i;
26 }
27
28 //HostToDevice
29 cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
30 cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
31
32 add << <N, 1 >> > (dev_a, dev_b, dev_c);
33
34 //将结果从GPU复制到CPU
35 cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
36
37 //输出结果
38 for (int i = 0;i < N;i++) {
39 printf("%d + %d = %d\\n", a[i], b[i], c[i]);
40 }
41 //释放内存
42 cudaFree(dev_a);
43 cudaFree(dev_b);
44 cudaFree(dev_c);
45
46 return 0;
47 }
上例仅给出函数main()，其在GPU上的实现与在CPU上的实现是不同的，但此时无差别
kernel<<<N，1>>>(dev_a, dev_b, dev_c);第一个参数表示设备在执行核函数时使用的并行线程块数量，运行N个核函数副本,前行线程块集合也称为一个线程格grid
- 在核函数中，通过变量
blockIdx.x确定当前运行区块 -
blockIdx.x为当前执行设备代码的线程块的索引
e.g.N=4，此时4个线程的的blockIdx.x值分别为0，1，2，3
每个线程块实际执行的代码如下：
4.1实例
Julia集：通过迭代等式对复平面中的等求值。
- 迭代等式计算结果发散，朝无穷大的方向增长，此点不属于Julia集合
- 迭代等式收敛，位于某个边界满园之内，此点属于Julia集合
迭代等式：
$$Z_{n+1}^2=Z_{n}^2+C$$
4.1.1基于CPU的Julia集
1 #include <stdio.h>
2
3 #include <cuda_runtime.h>
4 #include <device_launch_parameters.h>
5
6 #include "D:\\common\\book.h"
7 #include "D:\\common\\cpu_bitmap.h"
8
9 #define DIM 1000
10
11 //计算在复数上进行，定义结构保存复数
12 //定义复数的加法和乘法运算
13 struct cuComplex {
14 float r;//实部r
15 float i;//虚部i
16 cuComplex(float a,float b):r(a),i(b){}
17 float magnitude2() { return r / r + i + i; }
18 cuComplex operator*(const cuComplex &a) {
19 return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
20 }
21 cuComplex operator*(const cuComplex &a) {
22 return cuComplex(r + a.r, i + a.i);
23 }
24 };
25
26 int julia(int x, int y) {
27 //实现图形绽放的scale因数
28 const float scale = 1.5;
29
30 //将像素坐标转换为空间坐标
31 //像素移动DIM/2个单位，将复平面原点定位在图像中心
32 //图像范围在-1.0到1.0，图像坐标绽放了DIM/2倍
33 float jx = scale*(float)(DIM / 2 - x) / (DIM / 2);
34 float jy = scale*(float)(DIM / 2 - y) / (DIM / 2);
35
36 //迭代公式中的C为-0.5+0.156i
37 cuComplex c(-0.8, 0.156);
38 cuComplex a(jx, jy);
39
40 int i = 0;
41 for (i = 0;i < 200;i++) {
42 //a = a*a + c;
43 if (a.magnitude2() > 1000)//迭代结果阈值
44 return 0;
45 }
46 return 1;
47 }
48
49
50
51
52 //核函数对绘制的所有点进行迭代
53 void kernel(unsigned char *ptr) {
54 for (int y = 0;y < DIM;y++) {
55 for (int x = 0;x < DIM;x++) {
56 int offset = x + y*DIM;
57
58 //调用julia()判断点是否属于Julia集
59 //是返回1，点为红色
60 //否返回0，点为黑色，可改
61 int juliaValue = julia(x, y);
62 ptr[ offset * 4 + 0 ] = 255 * juliaValue;
63 ptr[ offset * 4 + 1 ] = 0;
64 ptr[ offset * 4 + 2 ] = 0;
65 ptr[ offset * 4 + 31 ] = 255;
66 }
67 }
68 }
69
70 int main() {
71 CPUBitmap bitmap(DIM, DIM); //通过工具库创建位图图像
72 unsigned char *ptr = bitmap.get_ptr;
73
74 //将指向位图数据的指针传递给核函数
75 kernel(ptr);
76
77 bitmap.display_and_exit();
78
79 return 0;
80 }
4.1.2基于GPU的Julia集
1 #include <stdio.h>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4
5 #include "D:\\common\\book.h"
6 #include "D:\\common\\cpu_bitmap.h"
7
8 #define DIM 1000
9
10 //计算在复数上进行，定义结构保存复数
11 //定义复数的加法和乘法运算
12 struct cuComplex {
13 float r;//实部r
14 float i;//虚部i
15 __device__ cuComplex(float a, float b) :r(a), i(b) {}
16 __device__ float magnitude2() {
17 return r * r + i * i;
18 }
19 __device__ cuComplex operator*(const cuComplex &a) {
20 return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
21 }
22 __device__ cuComplex operator+(const cuComplex &a) {
23 return cuComplex(r + a.r, i + a.i);
24 }
25 };
26
27
28 //判断点是否属于Julia集
29 __device__ int julia(int x, int y) {
30 const float scale = 1.5;
31 float jx = scale*(float)(DIM / 2 - x) / (DIM / 2);
32 float jy = scale*(float)(DIM / 2 - y) / (DIM / 2);
33
34 cuComplex c(-0.8, 0.156);
35 cuComplex a(jx, jy);
36
37 int i = 0;
38 for (i = 0;i < 200;i++) {
39 a = a*a + c;
40 if (a.magnitude2() > 1000)
41 return 0;
42 }
43 return 1;
44 }
45
46
47
48
49 //不需要for()来生成像素索引传递给julia()
50 //cuda运行时在变量blockIdx中包含这些索引
51 //在声明线程格时，线程格每一维的大小与图像每一维的大小是相等的，因此
52 //在（0,1）到（DIM,DIM）之间每个像素点都能分配一个线程块
53
54 __global__ void kernel(unsigned char *ptr) {
55 //将threadIdx/BlockIdx映射到像素位置
56 int x = blockIdx.x;
57 int y = blockIdx.y;
58 //内置变量gridDim，常数，保存线程格每一维大小
59 //行索引乘以线程格宽度+列索引得到ptr唯一索引，范围(DIM*DIM-1)
60 int offset = x + y*gridDim.x;
61
62 //计算此位置上的值
63 int juliaValue = julia(x, y);
64 ptr[offset * 4 + 0] = 255 * juliaValue;
65 ptr[offset * 4 + 1] = 0;
66 ptr[offset * 4 + 2] = 0;
67 ptr[offset * 4 + 3] = 255;
68 }
69
70 int main() {
71 //创建DIM*DIM大小的位图图像
72 CPUBitmap bitmap(DIM, DIM);
73 //保存设备上数据的副本
74 unsigned char *dev_bitmap;
75
76 cudaMalloc((void**)&dev_bitmap, bitmap.image_size());
77
78 dim3 grid(DIM, DIM);
79 kernel << <grid, 1 >> >(dev_bitmap);
80
81 //返回计算结果
82 cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
83
84 bitmap.display_and_exit();
85 cudaFree(dev_bitmap);
86 }
计算结果
计算线程块需要的数据索引
- 核函数的每个副本可以通过内置变量blockIdx来判断哪个线程块在执行它
- 通过内置变量gridDim获得线程格的大小
5线程协作
kernel<<<N,1>>>- 第一个参数是启动的线程块数量
- CUDA运行时每个线程块中创建的线程数量
- 启动的总线程数量 N个线程块*1个线程/线程块=N个并行线程
5.1矢量求和
5.1.1使用线程实现GPU上矢量求和
改动：
- add<<<N,1>>>(dev_a,dev_b,dev_c) -> add<<<1,N>>>(dev_a,dev_b,dev_c)
- 数据索引方法线程块索引变为线程索引
int tid = blockIdx.x; -> int tid = threadIdx.x;完整程序
1 #include <iostream>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4
5 #define N 10000
6
7 __global__ void add(int *a, int *b, int *c) {
8 int tid = threadIdx.x; //计算位于此索引处的数据
9 if (tid < N)
10 c[tid] = a[tid] + b[tid];
11 }
12
13 int main() {
14 int a[N], b[N], c[N];
15 int *dev_a, *dev_b, *dev_c;
16
17 //GPU上分配内存
18 cudaMalloc((void**)&dev_a, N * sizeof(int));
19 cudaMalloc((void**)&dev_b, N * sizeof(int));
20 cudaMalloc((void**)&dev_c, N * sizeof(int));
21
22 //对数组a,b赋值
23 for (int i = 0;i < N;i++) {
24 a[i] = -i;
25 b[i] = i*i;
26 }
27
28 //HostToDevice
29 cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
30 cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
31
32 add << <1, N >> > (dev_a, dev_b, dev_c);
33
34 //将结果从GPU复制到CPU
35 cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
36
37 //输出结果
38 for (int i = 0;i < N;i++) {
39 printf("%d + %d = %d\\n", a[i], b[i], c[i]);
40 }
41 //释放内存
42 cudaFree(dev_a);
43 cudaFree(dev_b);
44 cudaFree(dev_c);
45
46 return 0;
47 }
GPU上对更长矢量求和
- 线程块每一维的数量限制为65535
- 启动核函数时每个线程块中的线程数量不能超过设备属性结构中maxThreadsPerBlock域的值 大部分是每个线程块512个线程
更改
核函数中的索引计算方法 核函数的调用方式
计算索引方法类似于将二维索引空间转换为线性空间的标准算法
int tid =threadIdx.x + blockIdx.x * blockDim.xgridDim 线程格中每一维的线程块数量 二维
blockDim 线程块中每一维的线程数量 三维
int offset = x + y * DIM;DIM表示线程块大小即线程的数量
y为线程块索引，x为线程块中的线程索引
计算得到索引：tid = threadIdx.x + blockIdx.x * blockDim.x
核函数调用
kernel <<<(N+127/128,128)>>>(dev_a，dev_b，dev_c)启动128个线程
N+127/128一种向上取整的算法，计算大于或等于N的128的最小倍数
对于多启动的线程，在访问输入数组和输出数组之前，检查线程的偏移是否位于0到N之间
if(tid<N) c[tid] = a[tid] + b[tid];当索引越过数组边界时，核函数将自动停止计算，核函数不对越过数组边界的内存进行读取或写入
GPU上对任意长度的矢量求和
线程块每一维的数量限制为65535
当矢量长度超过限制时，核函数调用会失败
解决方法：将并行线程的数量看作是处理器的数量
认定每个线程在逻辑上都可以并行执行，并且硬件可以调用这些线程以便实际执行。通过将并行化过程与硬件的实际执行过程解耦开来。
步骤：
- 计算每个并行线程的初始化索引，以及递增的线程
- 对线程索引和线程块索引进行线性化，使每个并行线程从不同的索引开始
起始索引：
int tid = threadIdx.x + blockIdx.c * blockDim.x;- 对索引进行递增，递增步长为线程格中正在运行的线程数量。此数值等于每个线程块中的线程数量乘以线程格中线程块的数量，即
tid += blockDim.x * gridDim.x;- 线程块数量确定没明确说明P59
add<<<128,128>>>(dev_a,dev_b,dev_c);总的程序：
1 #include <stdio.h>
2 #include <cuda_runtime.h>
3 #include <device_launch_parameters.h>
4
5 #include "D:\\common\\book.h"
6 //#include "D:\\common\\cpu_bitmap.h"
7
8 #define N (33*1024)
9
10 __global__ void add(int *a, int *b, int *c) {
11 int tid = threadIdx.x + blockIdx.x*blockDim.x;
12 while (tid < N) {
13 c[tid] = a[tid] + b[tid];
14 tid += blockDim.x*gridDim.x;
15 }
16 }
17
18 int main() {
19 int a[N], b[N], c[N];
20 int *dev_a, *dev_b, *dev_c;
21
22 //GPU上分配内存
23 cudaMalloc((void**)&dev_a, N * sizeof(int));
24 cudaMalloc((void**)&dev_b, N * sizeof(int));
25 cudaMalloc((void**)&dev_c, N * sizeof(int));
26
27 //CPU上为数组a,b赋值
28 for (int i = 0;i < N;i++) {
29 a[i] = i;
30 b[i] = i*i;
31 }
32
33 //将数组a,b复制到GPU
34 cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
35 cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
36 add<<<128,128>>>(dev_a, dev_b, dev_c);
37
38 //将数组c复制回CPU
39 cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
40
41 //验证GPU完成了工作
42 bool success = true;
43 for (int i = 0;i < N;i++) {
44 if ((a[i] + b[i]) != c[i]) {
45 printf("Error: %d + %d != %d\\n", a[i], b[i], c[i]);
46 success = false;
47 }
48 }
49 if (success)
50 printf("done\\n");
51
52 //释放GPU上内存
53 cudaFree(dev_a);
54 cudaFree(dev_b);
55 cudaFree(dev_c);
56
57 return 0;
58 }
5.2.2在GPU上使用线程实现波纹效果
1 #include "D:/common/book.h"
2 #include "D:/common/cpu_anim.h"
3
4 #define DIM 1024
5
6 struct DataBlock {
7 unsigned char *dev_bitmap;
8 CPUAnimBitmap *bitmap;
9 };
10
11 void cleanup(DataBlock *d) {
12 cudaFree(d->dev_bitmap);
13 }
14
15 __global__ void kernel(unsigned char* ptr, int ticks) {
16 //将threadIdx、BlockIdx映射到像素位置
17 //线程得到其在线程块中的索引，及此线程块在线程格中的索引，并将两值转换为图形的唯一索引(x,y)
18 int x = threadIdx.x + blockIdx.x * blockDim.x;
19 int y = threadIdx.y + blockIdx.y * blockDim.y;
20
21 //对x,y进行线性化得到输出缓冲区中的一个偏移
22 int offset = x + y * blockDim.x * gridDim.x;
23 //int offset = y + x * blockDim.y * gridDim.y;//这两个offset等效
24
25 float fx = x - DIM / 2;
26 float fy = y - DIM / 2;
27 float d = sqrtf(fx * fx + fy * fy);
28 unsigned char grey = (unsigned char)(128.0f + 127.0f * cos(d / 10.0f - ticks / 7.0f) / (d / 10.0f + 1.0f));
29 ptr[offset * 4 + 0] = grey;//grey 2D时间函数
30 ptr[offset * 4 + 1] = grey;
31 ptr[offset * 4 + 2] = grey;
32 ptr[offset * 4 + 3] = 255;
33 }
34
35 void generate_frame(DataBlock *d, int ticks) {
36 dim3 blocks(DIM / 16, DIM / 16); //声明一个二维变量，线程格中包含的并行线程块数量
37 dim3 threads(16, 16); //声明一个二维变量，线程块中包含的线程数量
38
39 //核函数来计算像素值
40 //指针指向保存输出像素值的设备内存，是全局变量，其指向的内存是在main()中华西的。全局性针对主机，参数要传递让设备能够访问到
41 //当前时效ticks传递给generate_frame(),核函数根据当前动画时间生成正确的帧
42 kernel << <blocks, threads >> > (d->dev_bitmap, ticks);
43 HANDLE_ERROR(cudaMemcpy(d->bitmap->get_ptr(),
44 d->dev_bitmap,
45 d->bitmap->image_size(),
46 cudaMemcpyDeviceToHost));
47 }
48
49 int main() {
50 DataBlock data;
51 CPUAnimBitmap bitmap(DIM, DIM, &data); //大部分复杂性隐藏在辅助类CPUAnimBitmap中
52 data.bitmap = &bitmap;
53 HANDLE_ERROR(cudaMalloc((void**)&data.dev_bitmap, bitmap.image_size()));
54
55 //将指向generate_frame()函数的指针传递给anim_and_exit()，每当生成一帧新的动画，都将调用generate_frame()
56 bitmap.anim_and_exit((void(*)(void*, int))generate_frame, (void(*)(void*))cleanup);
57 }
本文章为转载内容，我们尊重原作者对文章享有的著作权。如有内容错误或侵权问题，欢迎原作者联系我们进行内容更正或删除文章。
