CUDA GPU编程中使用结构体传递函数参数

最新推荐文章于 2023-10-23 11:25:35 发布
最新推荐文章于 2023-10-23 11:25:35 发布
阅读量2.3k 收藏 12
点赞数 3
分类专栏： CUDA编程学习 文章标签： cuda 函数

CUDA GPU编程中使用结构体传递函数参数

 

         CUDA GPU编程中，虽然统一寻址(Unified Memory)技术能够大大简化编程难度和代码复杂度，但是速度略有牺牲，同时对运行环境提出更多的要求。而在不使用这项技术时，编程时需要同时创建CPU（host）和GPU（device）端的变量指针，然后为其分别分配内存。操作完成后，再分别释放内存。CUDA工程的范例程序中，单独编写了__global__函数，并将其和设备内存的创建，使用和释放包装在了一个函数里面，最终将这个包装后的函数提供给用户。用户只需管理自己的CPU内存，以及简单地调用该封装函数即可实现GPU加速。

但实际处理过程中，由于数据量大或者希望使用CUDA流来实现CPU和GPU并行，从而最大限度提高运行效率，由于内存的创建和释放效率极低，上诉方式并不可取。而是在程序开始创建变量指针并分配内存，随后对该内存进行多次使用，在程序运行结束时再释放资源。这样之前的一个封装函数显然不够用，现在需要用户自己创建各项主机和设备端内存指针，然后分别提供相应封装函数以实现内存的创建、核心的GPU并行计算代码、内存释放若干操作。应用程序往往需要使用若干数据，而分别要在主机和设备端建立一个指针，因此造成函数参数数量过多的问题，我的一些程序中经常超过20个参数，这对程序的调用和管理造成极大的不便。

         为了解决这个问题，考虑把相关必要的参数封装在一个结构体里面，通过该结构体的指针进行传输传递，而封装函数通过该结构体获取必要的数据，从而函数的参数大大降低。而进一步，把结构体内部指针变量对应的内存创建、释放封装起来，用户只需要创建相应的结构体再调用这些封装函数即可，对代码的易用性和易维护性均有很大的提升。

         下面首先提供一个范例，这是根据cuda工程创建的范例进行的修改。

[cpp] view plain copy
print ?
  1. #include “cuda_runtime.h”  
  2. #include “device_launch_parameters.h”  
  3.   
  4. #include <stdio.h>  
  5. #include <string.h>  
  6.   
  7. struct ProcPara  
  8. {  
  9.     int h_a; // host CPU  
  10.     int *h_b; // host CPU  
  11.     int *h_c; // host CPU  
  12.   
  13.     int *d_a; // device GPU  
  14.     int *d_b; // device GPU  
  15.     int *d_c; // device GPU  
  16.   
  17. };  
  18.   
  19. global void addKernel(ProcPara para)  
  20. {  
  21.     int i = threadIdx.x;  
  22.     para->d_c[i] = para->d_a[i] + para->d_b[i];  
  23. }  
  24.   
  25. void InitProcPara(ProcPara* para, int arraySize)  
  26. {  
  27. cudaMallocHost((void)&para->h_a, arraySize  sizeof(int));  
  28. cudaMallocHost((void)&para->h_b, arraySize  sizeof(int));  
  29. cudaMallocHost((void)&para->h_c, arraySize  sizeof(int));  
  30.   
  31.   
  32. cudaMalloc((void)&para->d_a, arraySize  sizeof(int));  
  33. cudaMalloc((void)&para->d_b, arraySize  sizeof(int));  
  34. cudaMalloc((void)&para->d_c, arraySize  sizeof(int));  
  35.   
  36. }  
  37.   
  38. void DeinitProcPara(ProcPara* para)  
  39. {  
  40.     // free host memory  
  41.     cudaFreeHost(para->h_a);  
  42.     cudaFreeHost(para->h_b);  
  43.     cudaFreeHost(para->h_c);  
  44.     // free device memory  
  45.     cudaFree(para->d_a);  
  46.     cudaFree(para->d_b);  
  47.     cudaFree(para->d_c);  
  48. }  
  49.   
  50. void addWithCuda(ProcPara* para, unsigned int arraySize)  
  51. {  
  52.     cudaSetDevice(0);  
  53.     cudaMemcpy(para->d_a, para->h_a, arraySize  sizeof(int), cudaMemcpyHostToDevice);  
  54.     cudaMemcpy(para->d_b, para->h_b, arraySize  sizeof(int), cudaMemcpyHostToDevice);  
  55.   
  56.     // Launch a kernel on the GPU with one thread for each element.  
  57.     addKernel << <1, arraySize >> >(para);  
  58.   
  59.     cudaMemcpy(para->h_c, para->d_c, arraySize  sizeof(int), cudaMemcpyDeviceToHost);  
  60.     cudaDeviceSynchronize();  
  61. }  
  62.   
  63.   
  64.   
  65.   
  66.   
  67. int main()  
  68. {  
  69.     const int arraySize = 5;  
  70.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
  71.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
  72.     int c[arraySize] = { 0 };  
  73.   
  74.     ProcPara para = new ProcPara;  
  75.   
  76.     InitProcPara(para, arraySize);  
  77.   
  78.   
  79.     memcpy(para->h_a, a, arraySize  sizeof(int));  
  80.     memcpy(para->h_b, b, arraySize  sizeof(int));  
  81.   
  82.     addWithCuda(para, arraySize);  
  83.   
  84.     memcpy(c, para->h_c, arraySize * sizeof(int));  
  85.   
  86.     printf(”{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n”,  
  87.         c[0], c[1], c[2], c[3], c[4]);  
  88.   
  89.     DeinitProcPara(para);  
  90.     delete para;  
  91.   
  92.     return 0;  
  93. }  
#include “cuda_runtime.h”

但是，这个程序并不能得到正确的结果。

这是因为，在传递给__global__函数的结构体指针的内存实体在主机内存中。而在非Unified Memory的条件下，GPU是不能访问到CPU端的内存数据的。由此，要继续使用结构体传递参数，需要进行如下的操作：

分别创建主机端和设备端结构体指针，并分配相应的内存。为主机端结构体内部成员变量分别分配主机端和设备端内存。然后将主机端结构体数据复制给设备端结构体。在核心计算代码中，主机端程序使用主机端结构体指针来获取所有数据的指针（包括数据拷贝），设备端程序使用设备端结构体指针获取设备端数据的指针。修改后的代码如下，这里需要注意的是，分配内存用到了指向指针的指针。因为指针本身是个变量，其保存的是个地址，分配内存要对该变量的内容进行更新，必须传递该变量的指针才有效，否则只是改变了参数变量的内容。这也是为什么cudaMalloc使用的是指针的指针的原因。

 

[cpp] view plain copy
print ?
  1. #include “cuda_runtime.h”  
  2. #include “device_launch_parameters.h”  
  3.   
  4. #include <stdio.h>  
  5. #include <string.h>  
  6.   
  7. struct ProcPara  
  8. {  
  9.     int h_a; // host CPU  
  10.     int *h_b; // host CPU  
  11.     int *h_c; // host CPU  
  12.   
  13.     int *d_a; // device GPU  
  14.     int *d_b; // device GPU  
  15.     int *d_c; // device GPU  
  16.   
  17. };  
  18.   
  19. global void addKernel(ProcPara d_para)  
  20. {  
  21.     int i = threadIdx.x;  
  22.     d_para->d_c[i] = d_para->d_a[i] + d_para->d_b[i];  
  23.   
  24. }  
  25.   
  26. void InitProcPara(ProcPara**ha_para, ProcPara**da_para, int arraySize)  
  27. {  
  28.     //allocate stucture memory  
  29.     cudaMallocHost((void)ha_para, sizeof(ProcPara));  
  30.     cudaMalloc((void)da_para, sizeof(ProcPara));  
  31.   
  32.   
  33.     ProcPara*h_para = ha_para;  
  34.   
  35.     cudaMallocHost((void)&h_para->h_a, arraySize  sizeof(int));  
  36.     cudaMallocHost((void)&h_para->h_b, arraySize  sizeof(int));  
  37.     cudaMallocHost((void)&h_para->h_c, arraySize  sizeof(int));  
  38.   
  39.     cudaMalloc((void)&h_para->d_a, arraySize  sizeof(int));  
  40.     cudaMalloc((void)&h_para->d_b, arraySize  sizeof(int));  
  41.     cudaMalloc((void)&h_para->d_c, arraySize  sizeof(int));  
  42.   
  43.     // exchange data  
  44.     cudaMemcpy(*da_para, *ha_para, sizeof(ProcPara), cudaMemcpyHostToDevice);  
  45.   
  46. }  
  47.   
  48. void DeinitProcPara(ProcPara*h_para, ProcPara*d_para)  
  49. {  
  50.     // free host memory  
  51.     cudaFreeHost(h_para->h_a);  
  52.     cudaFreeHost(h_para->h_b);  
  53.     cudaFreeHost(h_para->h_c);  
  54.     // free device memory  
  55.     cudaFree(h_para->d_a);  
  56.     cudaFree(h_para->d_b);  
  57.     cudaFree(h_para->d_c);  
  58.   
  59.     //release stucture memory  
  60.     cudaFreeHost(h_para);  
  61.     cudaFree(d_para);  
  62.   
  63. }  
  64.   
  65. void addWithCuda(ProcPara h_para, ProcPara* d_para, unsigned int arraySize)  
  66. {  
  67.     cudaSetDevice(0);  
  68.   
  69.     cudaMemcpy(h_para->d_a, h_para->h_a, arraySize  sizeof(int), cudaMemcpyHostToDevice);  
  70.     cudaMemcpy(h_para->d_b, h_para->h_b, arraySize  sizeof(int), cudaMemcpyHostToDevice);  
  71.   
  72.     cudaDeviceSynchronize();  
  73.   
  74.     // Launch a kernel on the GPU with one thread for each element.  
  75.     addKernel << <1, 5 >> >(d_para);  
  76.   
  77.     cudaMemcpy(h_para->h_c, h_para->d_c, arraySize  sizeof(int), cudaMemcpyDeviceToHost);  
  78.     cudaDeviceSynchronize();  
  79.   
  80. }  
  81.   
  82.   
  83.   
  84. int main()  
  85. {  
  86.     const int arraySize = 5;  
  87.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
  88.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
  89.     int c[arraySize] = { 0 };  
  90.   
  91.     ProcPara *h_para;  
  92.     ProcPara *d_para;  
  93.   
  94.   
  95.     InitProcPara(&h_para, &d_para, arraySize);  
  96.   
  97.   
  98.     memcpy(h_para->h_a, a, arraySize  sizeof(int));  
  99.     memcpy(h_para->h_b, b, arraySize  sizeof(int));  
  100.   
  101.     addWithCuda(h_para, d_para, arraySize);  
  102.   
  103.     memcpy(c, h_para->h_c, arraySize  sizeof(int));  
  104.   
  105.     printf(”{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n”,c[0], c[1], c[2], c[3], c[4]);  
  106.   
  107.     DeinitProcPara(h_para, d_para);  
  108.   
  109.     return 0;  
  110. }  

命中缺汤
关注
  • 3
    点赞
  • 踩
  • 12
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
专栏目录
CUDA传递结构体数组
进击的小白
09-22 5378
在把CPU的数组结构体传递到GPU内存中需要好多问题,找不到太多的资料,最后只能靠自己慢慢摸索 CPU上的代码我们一般称为 Host代码 GPU上面的一般叫Device 代码,之所以这样区分是由于GPU的内存和CPU的内存是相互独立的,它们的数据没法直接相互调用,只能通过cuda的一些方法进行操作,最常用的是 cudaMalloc和cudaMemcpy cudaMalloc是用来在GPU上申请
使用pybind11封装C++结构体作为参数的函数实现步骤
08-24
使用pybind11封装C++结构体作为参数的函数实现步骤 pybind11是轻量级的只包含头文件的库，它主要是用来在已有的C++代码的基础上做扩展。pybind11的优点是对C++ 11支持很好，API比较简单。pybind11摒弃了对旧版...
在cuda中,使用自定义结构体,
闲暇余客
08-25 3789
在使用GPU加速中,有时需要传入结构体来进行统一计算,完后再同意导出,所以,在参考了一些资料,做了一下测试, 并记录,以供以后继续看
CUDA GPU编程中使用结构体/类传递函数参数
LuchangLi 的专栏
09-21 1万+
CUDA GPU编程中使用结构体传递函数参数 CUDAGPU编程中，虽然统一寻址(Unified Memory)技术能够大大简化编程难度和代码复杂度，但是速度略有牺牲，同时对运行环境提出更多的要求。而在不使用这项技术时，编程时需要同时创建CPU（host）和GPU（device）端的变量指针，然后为其分别分配内存。操作完成后，再分别释放内存。CUDA工程的范例程序中，单
向CUDA中传入多维结构体
AndrewBW666的博客
12-30 622
向CUDA中传入多维结构体 逻辑相对复杂，一些应用中也不一定用的上传入多维结构体，所以不做文字解释。一个样例程序放在这里为大家提供思路，供大家参考。如果有什么疑问欢迎大家评论，我会及时回复。 /** * Copyright 1993-2012 NVIDIA Corporation. All rights reserved. * * Please refer to the NVIDIA e...
cuda 从CPU到GPU的结构体数组传输
UserName11的博客
05-10 3055
结构体的数组传输，目前的理解是在GPU开辟一块显存，和CPU中的数据的排列顺序一一对应，并在GPU定义一个结构体，能够灵活的访问数据。现在就以pycuda的DemoStruct为例来进行说明。目前指针只用到了一维的，在GPU中定义结构体： 在这个结构体中有三个变量datalen，__padding，ptr。其中ptr是一个指针，datalen指的当前ptr中的数据长度，__padding为占位符无...
CUDA数据结构体在GPU和CPU之间的传输
Alfafar的博客
03-21 1825
参照《cuda c编程权威指南》上写出的，在CPU和GPU互相传输结构体。 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #include<stdlib.h> #define maxnumber 100 typedef struct node { int x[maxnumber]; int y[maxnumber]; int test[maxnumbe
C++结构体作为函数参数传参的实例代码
01-21
*结构体作为函数参数传参 */ //值传递 void printStufdent1(struct Student st3) { cout << 子函数 << endl; st3.age = 100; cout << 名字： << st3.name << 年龄： << ...
结构体类型数据作为函数参数(三种方法)
01-01
（1）用结构体变量名作为参数。 代码如下:#include<iostream>#include<string>using namespace std;struct Student{ string name; int score; };int main(){ Student one; void Print(Student one); one.name=...
C#调用C++动态库，执行回调函数并传递结构体参数
12-25
c#调用C++动态库、执行回调函数，并回传结构体参数数据。vs2017环境编写C#和C++动态库，这个为完整工程例子，可供相关人员学习参考。
CUDA结构体或类的传递内存实验
Dezeming的博客
06-29 1809
我想把指针传送到CUDA内存区域中，但是遇到了一定的问题，最终找到了解决方法，这里进行一下记录： 一、普通传递没有任何问题： cudaMalloc((void**)&a_dev, 50 * sizeof(int)); a_host = (int*)malloc(sizeof(int) * 50); for(inti = 0;i < 50;i++) a_host[i] = i; cudaMemcpy(a_dev, a_host, 50 * sizeof(int), cudaMem..
结构体在cuda核函数里面不能用吗_细品Redis高性能数据结构之SDS
weixin_39945445的博客
12-08 150
背景redis之所以快，除了他是基于内存存储的，还有优秀的IO框架外更离不了其底层高性能数据结构的设计。现在我们来细细品一下redis的高新能数据结构是如何设计的。Redis数据结构redis有五种数据结构，分别是String，list（列表），Set（集合），Hash（哈希），Zset（有序结合）。这五种数据结构代表的是他的Value值的数据结构。今天我们来看redis的String类型。动态字...
结构体在cuda核函数里面不能用吗_深度学习知识点总结二（深度网络结构与池化）...
weixin_39668479的博客
11-21 295
关于深度学习网络结构的一些问题1、卷积原理，工程实现？参考：https://blog.csdn.net/wweiainn/article/details/802317922、CNN为什么对图像work的特别好？强大的非线性拟合能力，结构化数据操作。3、卷积在图像中有什么直观作用？4、卷积核是否越大越好？卷积核不是越大越好，大多数情况下，通过堆叠较小的卷积核比直接采用单个更大的卷积核会更加有效。一般...
cuda 为包含指针的结构数组指针分配统一内存(UM)
最新发布
G751888629的博客
10-23 188
c/c++ cuda unified memory
给cuda核函数传递二维数组的一种方法
linger(心怀梦想，活在当下，知乎ID：linger liu)
04-25 9567
#include /** * 需求：需要把若干个一维数组传给核函数 * 实现方法：在gpu生成一个一维的指针数组，每个元素指向一个普通一维数组。 * 把该指针数组的地址传递给核函数。 * 其实该指针数组充当二维数组的角色。 */ __global__ void testKernel(float ** pointerArray) { printf("(%d,%d):%f\n",th
cuda从入门到精通（三）之核函数和参数配置
热门推荐
xiangxianghehe的博客
06-13 1万+
这里继续讲上一节cuda从入门到精通（二）之函数关键字的一个遗留问题：__global__修饰函数中的<<<1,1>>>到底是什么。
cuda 学习之内存层次结构
xukang95的博客
03-23 2834
文章目录Cuda 内存层次结构registers：寄存器local memory：局部存储器shared memory：共享存储器共享内存-Bank Conflictconstant memory：常数存储器texture memory：纹理存储器global memory：全局存储器全局内存-对齐与合并host memory :主机端内存总结 Cuda 内存层次结构 要想编写高效的程序，那么一...
cuda数据传输之cudaMemcpy()和cudaMemcpy2D()详解
qq_45241855的博客
05-11 5973
核函数中：a = (double*)((char*)dev_a + i*pitch)，这样对a的遍历方式是获取数组每一行的数据，同时计算a的结果直接作用在dev_a 上，而不需要再将a赋值给dev_a！总结：无论一维还是二维数组，都变为一维数组，进行传递，二维数组要注意数组的宽度大小。2. cudaMemcpy2D()传递一维数组、二维数组，已经核函数索引遍历的方法。再核函数中对一维数组索引，直接是0下标索引。
arduino 定义一个结构体 并且函数中传递改结构体参数
05-30
在 Arduino 中，定义一个结构体可以使用 C 语言的结构体定义方式。例如，定义一个包含两个整型变量和一个字符型数组的结构体可以使用以下代码： ``` struct MyStruct { int a; int b; char c[10]; }; ``` 这里我们定义了一个名为 `MyStruct` 的结构体，包含两个整型变量 `a` 和 `b`，以及一个长度为 10 的字符型数组 `c`。 接下来，我们可以在函数中传递这个结构体参数。例如，定义一个函数 `myFunction`，该函数接受一个 `MyStruct` 类型的参数： ``` void myFunction(MyStruct myStruct) { // Do something with myStruct } ``` 在函数内部，可以像使用普通变量一样使用 `myStruct`，例如： ``` void myFunction(MyStruct myStruct) { Serial.print("a = "); Serial.println(myStruct.a); Serial.print("b = "); Serial.println(myStruct.b); Serial.print("c = "); Serial.println(myStruct.c); } ``` 这里我们使用 `Serial` 对象输出结构体中的成员变量。 最后，在调用函数时，需要传递一个 `MyStruct` 类型的参数。例如： ``` MyStruct myStruct = {1, 2, "hello"}; myFunction(myStruct); ``` 在上面的例子中，我们创建了一个名为 `myStruct` 的结构体变量，并初始化其成员变量。然后，将其作为参数传递给 `myFunction` 函数。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值