找回密码
 立即注册
回帖奖励 20 CUDA币 回复本帖可获得 20 CUDA币奖励! 每人限 5 次
daxiaoyuyu 该用户已被删除
发表于 2013-11-1 19:09:17
17382

最近小弟真的是非常郁闷,发现在cuda
中用mallocmemcpy,这样的函数效率非常低。各路高手走过路过不要错过,给小弟一点帮助吧,哪怕一点建议都行,不胜感激。
我的程序中需要在一张表中提出100个参数(这不是固定的),需要将椭圆分成n个面积为sideLength* sideLength的微小平面,并且把这些平面的中心坐标存放起来供给下一个函数使用。
由于sideLengthellipseA ellipseB都是可以任意指定,所以一个椭圆存放多少个微小平面的中心坐标是未知的。所以我写了一个动态数组,类似C++vector,用来存放微平面中心点坐标。
struct  Para
{
       // x^2/g_ellipseA + y^2/g_ellipseB <= 1
       float sideLength; //表示边长
       float ellipseA; //表示椭圆短轴的平方。
       float ellipseB; //表示椭圆张轴的平方。
};
我做了一组测试
cuda中:(block * thread 数)
1 ReceiveFxArray()调用1*1: 4.3s
400 ReceiveFxArray()调用2*200: 10.9s  4*100: 10.7s
600 ReceiveFxArray()调用:3*200: 45.7.6s  6*100: 47.03s
1000 ReceiveFxArray()调用:5*200: 201.4s 10 * 100: 197.3s 1*1000: 203.5s
CPU中:
1次ReceiveFxArray()调用:0.047sCPU91倍。
400次ReceiveFxArray()调用:20s,GPU快1.8倍。GPU终于比CPU块了
600次ReceiveFxArray()调用:31s,CPU快1.5倍
1000 ReceiveFxArray()调用:51s,CPU快4倍。
问题:
1.malloc,memcpy等函数在GPU中为什么会这么慢?
2.为什么GPU中的线程分配过多之后,反而更慢了?是不是分配的线程过多后,会产生更多的显存碎片,从而会更慢?

使用道具 举报 回复
发表于 2013-11-1 19:10:45
GPU代码:
  1. 1
  2. 2
  3. 3
  4. 4
  5. 5
  6. 6
  7. 7
  8. 8
  9. 9
  10. 10
  11. 11
  12. 12
  13. 13
  14. 14
  15. 15
  16. 16
  17. 17
  18. 18
  19. 19
  20. 20
  21. 21
  22. 22
  23. 23
  24. 24
  25. 25
  26. 26
  27. 27
  28. 28
  29. 29
  30. 30
  31. 31
  32. 32
  33. 33
  34. 34
  35. 35
  36. 36
  37. 37
  38. 38
  39. 39
  40. 40
  41. 41
  42. 42
  43. 43
  44. 44
  45. 45
  46. 46
  47. 47
  48. 48
  49. 49
  50. 50
  51. 51
  52. 52
  53. 53
  54. 54
  55. 55
  56. 56
  57. 57
  58. 58
  59. 59
  60. 60
  61. 61
  62. 62
  63. 63
  64. 64
  65. 65
  66. 66
  67. 67
  68. 68
  69. 69
  70. 70
  71. 71
  72. 72
  73. 73
  74. 74
  75. 75
  76. 76
  77. 77
  78. 78
  79. 79
  80. 80
  81. 81
  82. 82
  83. 83
  84. 84
  85. 85
  86. 86
  87. 87
  88. 88
  89. 89
  90. 90
  91. 91
  92. 92
  93. 93
  94. 94
  95. 95
  96. 96
  97. 97
  98. 98
  99. 99
  100. 100
  101. 101
  102. 102
  103. 103
  104. 104
  105. 105
  106. 106
  107. 107
  108. 108
  109. 109
  110. 110
  111. 111
  112. 112
  113. 113
  114. 114
  115. 115
  116. 116
  117. 117
  118. 118
  119. 119
  120. 120
  121. 121
  122. 122
  123. 123
  124. 124
  125. 125
  126. 126
  127. 127
  128. 128
  129. 129
  130. 130
  131. 131
  132. 132
  133. 133
  134. 134
  135. 135
  136. 136
  137. 137
  138. 138
  139. 139
  140. 140
  141. 141
  142. 142
  143. 143
  144. 144
  145. 145
  146. 146
  147. 147
  148. 148
  149. 149
  150. 150
  151. 151
  152. 152
  153. 153
  154. 154
  155. 155
  156. 156
  157. 157
  158. 158
  159. 159
  160. 160
  161. 161
  162. 162
  163. 163
  164. 164
  165. 165
  166. 166
  167. 167
  168. 168
  169. 169
  170. 170
  171. 171
  172. 172
  173. 173
  174. 174
  175. 175
  176. 176
  177. 177
  178. 178
  179. 179
  180. 180
  181. 181
  182. 182
  183. 183
  184. 184
  185. 185
  186. 186
  187. 187
  188. 188
  189. 189
  190. 190
  191. 191
  192. 192
  193. 193
  194. 194
  195. 195
  196. 196
  197. 197
  198. 198
  199. 199
  200. 200
  201. 201
  202. 202
  203. 203
  204. 204
  205. 205
  206. 206
  207. 207
  208. 208
  209. 209
  210. 210
  211. #include <stdlib.h>
  212. #include <stdio.h>

  213. #define PI 3.1415926535897
  214. #define MIN(x, y) ((x)>(y)?(y):(x))

  215. #define MINLEN 256
  216. #define EXPANED_TOTLEN 1
  217. #define CVPUSHBACK 1; //内存回退
  218. #define CVSUCESS 0; //内存分配成功

  219. struct Point3D
  220. {
  221.     float  x, y, z;
  222.     int attr;
  223. };

  224. struct  Para
  225. {
  226.     // x^2/g_ellipseA + y^2/g_ellipseB <= 1
  227.     float sideLength;
  228.     float ellipseA;
  229.     float ellipseB;
  230. };

  231. struct CVector
  232. {
  233.     void *pData;//指向所需的数据
  234.     size_t len,totLen,typeSize;
  235. };

  236. //-------c语言vector版--------
  237. __device__ CVector*  CVectorCreate(size_t const size)
  238. {
  239.     CVector *pCV = (CVector *)malloc(sizeof(CVector));

  240.     if (!pCV)
  241.     {
  242.         return NULL;
  243.     }

  244.     pCV->pData = malloc( MINLEN * size );

  245.     if (!pCV->pData)
  246.     {
  247.         free(pCV);
  248.         return NULL;
  249.     }

  250.     pCV->len = 0;
  251.     pCV->totLen = MINLEN;
  252.     pCV->typeSize = size;

  253.     return pCV;
  254. }

  255. __device__ int CVectorPushBack(CVector * const pCV, void const * const pMem)
  256. {
  257.     if (pCV->len >= pCV->totLen)
  258.     {
  259.         void *pDataSave = pCV->pData;
  260.         int totLenSave = pCV->totLen;
  261.         pCV->totLen <<= EXPANED_TOTLEN;
  262.         pCV->pData = malloc(pCV->totLen * pCV->typeSize);

  263.         if ( !pCV->pData )//如果分配新的内存不成功,滚回。
  264.         {
  265.             pCV->pData = pDataSave;
  266.             pCV->totLen = totLenSave;
  267.             return CVPUSHBACK;
  268.         }
  269.         //把原来的数据复制进去
  270.         memcpy(pCV->pData, pDataSave, totLenSave * pCV->typeSize);
  271.         free(pDataSave);
  272.     }
  273.     //把新数据复制进去。
  274.     memcpy( (char *)pCV->pData + pCV->len * pCV->typeSize, pMem, pCV->typeSize );
  275.     pCV->len ++;
  276.     return CVSUCESS;
  277. }

  278. __device__ size_t CVectorLength(CVector const * const pCV)
  279. {
  280.     return pCV->len;
  281. }

  282. __device__ CVector * CVectorCopy(CVector const * const pCV)
  283. {
  284.     CVector * pNewCV = (CVector *)malloc(sizeof(CVector));
  285.     if (!pNewCV)
  286.     {
  287.         return NULL;
  288.     }
  289.     pNewCV->pData = malloc( pCV->typeSize * pCV->totLen);

  290.     if ( !pNewCV->pData )
  291.     {
  292.         free(pNewCV);
  293.         return NULL;
  294.     }

  295.     memcpy(pNewCV->pData, pCV->pData, pCV->typeSize * pCV->totLen);
  296.     pNewCV->len = pCV->len;
  297.     pNewCV->totLen = pCV->totLen;
  298.     pNewCV->typeSize = pCV->typeSize;
  299.     return pNewCV;
  300. }

  301. //取得指定位置的指针
  302. __device__ void* GetCVecElement(CVector * const pCV, size_t index)
  303. {
  304.     return (char*)pCV->pData + pCV->typeSize * index;
  305. }

  306. __device__ void CVectorDestroy(CVector * const pCV)
  307. {
  308.     free(pCV->pData);
  309.     free(pCV);
  310.     return;
  311. }

  312. //---获取椭圆范围内微平面的中心坐标点。
  313. __device__ void GetCodLeafAndCam(CVector * const pCodLeafFacet, const Para& para)
  314. {
  315.     //先提取出 xMax * yMax 矩形的微微平面中心坐标点。
  316.     //再判断这些坐标点是否落在椭圆内,从而提取出椭圆范围内的中心坐标点。
  317.     float xMax = 3*18;
  318.     float yMax = 3*53;

  319.     //按sideLength分割 xMax。
  320.     CVector *pXnCenter = CVectorCreate(sizeof(float));
  321.     for(float i = -xMax; i <= xMax; i += para.sideLength)
  322.     {
  323.         CVectorPushBack(pXnCenter, &i);
  324.     }

  325.     //按sideLength分割 yMax
  326.     CVector *pYnCenter = CVectorCreate(sizeof(float));
  327.     for(float i = -yMax; i <= yMax; i += para.sideLength)
  328.     {
  329.         CVectorPushBack(pYnCenter, &i);
  330.     }

  331.     //提取出坐标点。
  332.     int numXnCenter = CVectorLength(pXnCenter);
  333.     int numYnCenter = CVectorLength(pYnCenter);
  334.     for ( int i = 0; i != numXnCenter; i++)
  335.     {
  336.         for ( int j = 0; j != numYnCenter; j++)
  337.         {
  338.             float tmpXn = *((float *)GetCVecElement(pXnCenter,i));
  339.             float tmpYn = *((float *)GetCVecElement(pYnCenter,j));
  340.             if (tmpXn*tmpXn/para.ellipseA + tmpYn*tmpYn/para.ellipseB <= 1)
  341.             {   
  342.                 Point3D tmpPoint;
  343.                 tmpPoint.x = tmpXn;
  344.                 tmpPoint.y = tmpYn;
  345.                 tmpPoint.z = 0;
  346.                 CVectorPushBack(pCodLeafFacet, &tmpPoint);
  347.             }
  348.         }
  349.     }
  350.     CVectorDestroy(pXnCenter);
  351.     CVectorDestroy(pYnCenter);
  352. }

  353. __global__ static void ReceiveFxArray(Para *pcuPara)
  354. {
  355.     CVector *pCodLeafCam = CVectorCreate( sizeof(Point3D) );
  356.     CVector *pCodLeafFacet = CVectorCreate( sizeof(Point3D) );
  357.     GetCodLeafAndCam(pCodLeafFacet, *pcuPara);
  358.     CVectorDestroy(pCodLeafFacet);
  359. }

  360. int main()
  361. {

  362.     size_t memSize;
  363.     cudaDeviceSetLimit(cudaLimitMallocHeapSize, 512*1024*1024);
  364.     cudaDeviceGetLimit(&memSize, cudaLimitMallocHeapSize);

  365.     Para tmpPara;
  366.     tmpPara.sideLength = 0.25;
  367.     tmpPara.ellipseA = 324;
  368.     tmpPara.ellipseB = 2704;

  369.     Para *pcuPara;
  370.     cudaMalloc( (void**)&pcuPara,  sizeof(Para) );
  371.     cudaMemcpy( pcuPara, &tmpPara, sizeof(Para), cudaMemcpyHostToDevice );
  372.     cudaEvent_t start, stop;
  373.     cudaEventCreate(&start);
  374.     cudaEventCreate(&stop);
  375.     cudaEventRecord(start, 0);

  376.     ReceiveFxArray<<<6,100,0>>> ( pcuPara);

  377.     cudaEventRecord(stop, 0);
  378.     cudaEventSynchronize(stop);
  379.     float elapsedTime;
  380.     cudaEventElapsedTime(&elapsedTime, start, stop);
  381.     printf("elapsedTime = %.5f\n",elapsedTime);
  382.      
  383.     cudaEventDestroy(start);
  384.     cudaEventDestroy(stop);
  385.      
  386.     cudaFree( pcuPara );

  387.     system("pause");
  388.     return 0;
  389. }
复制代码
使用道具 举报 回复 支持 反对
发表于 2013-11-1 19:11:42
CPU代码:
  1. #include <stdlib.h>
  2. #include <stdio.h>
  3. #include <string.h>
  4. #include <time.h>

  5. #define PI 3.1415926535897
  6. #define MIN(x, y) ((x)>(y)?(y):(x))

  7. #define MINLEN 256
  8. #define EXPANED_TOTLEN 1
  9. #define CVPUSHBACK 1; //内存回退
  10. #define CVSUCESS 0; //内存分配成功

  11. struct Point3D
  12. {
  13.     double x, y, z;
  14. };

  15. struct  Para
  16. {
  17.     // x^2/g_ellipseA + y^2/g_ellipseB <= 1
  18.     float sideLength;
  19.     float ellipseA;
  20.     float ellipseB;
  21. };

  22. struct CVector
  23. {
  24.     void *pData;//指向所需的数据
  25.     size_t len,totLen,typeSize;//pData里面有多少个数据,最多能容纳多少个数据,数据类型的字节数。
  26. };

  27. CVector*  CVectorCreate(size_t const size)
  28. {
  29.     CVector *pCV = (CVector *)malloc(sizeof(CVector));//sizeof是用来确定类型的字节数的。

  30.     if (!pCV) //如果内存分配不成功,则 pCV 为NULL, 即为0;
  31.     {
  32.         return NULL;
  33.     }

  34.     pCV->pData = malloc( MINLEN * size );

  35.     if (!pCV->pData)
  36.     {
  37.         free(pCV);
  38.         return NULL;
  39.     }

  40.     pCV->len = 0;
  41.     pCV->totLen = MINLEN;
  42.     pCV->typeSize = size;

  43.     return pCV;
  44. }

  45. int CVectorPushBack(CVector * const pCV, void const * const pMem)
  46. {
  47.     if (pCV->len >= pCV->totLen)
  48.     {
  49.         void *pDataSave = pCV->pData;
  50.         int totLenSave = pCV->totLen;
  51.         pCV->totLen <<= EXPANED_TOTLEN; //往前移位,意味着扩大一倍。
  52.         pCV->pData = malloc(pCV->totLen * pCV->typeSize);

  53.         if ( !pCV->pData )//如果分配新的内存不成功,滚回。
  54.         {
  55.             pCV->pData = pDataSave;
  56.             pCV->totLen = totLenSave;
  57.             return CVPUSHBACK;
  58.         }
  59.         //把原来的数据复制进去
  60.         memcpy(pCV->pData, pDataSave, totLenSave * pCV->typeSize);
  61.         free(pDataSave);
  62.     }
  63.     //把新数据复制进去。
  64.     memcpy( (char *)pCV->pData + pCV->len * pCV->typeSize, pMem, pCV->typeSize );
  65.     pCV->len ++;
  66.     return CVSUCESS;
  67. }

  68. size_t CVectorLength(CVector const * const pCV)
  69. {
  70.     return pCV->len;
  71. }

  72. CVector * CVectorCopy(CVector const * const pCV)
  73. {
  74.     CVector * pNewCV = (CVector *)malloc(sizeof(CVector));
  75.     if (!pNewCV)
  76.     {
  77.         return NULL;
  78.     }
  79.     pNewCV->pData = malloc( pCV->typeSize * pCV->totLen);

  80.     if ( !pNewCV->pData )
  81.     {
  82.         free(pNewCV);
  83.         return NULL;
  84.     }

  85.     memcpy(pNewCV->pData, pCV->pData, pCV->typeSize * pCV->totLen);
  86.     pNewCV->len = pCV->len;
  87.     pNewCV->totLen = pCV->totLen;
  88.     pNewCV->typeSize = pCV->typeSize;
  89.     return pNewCV;
  90. }

  91. //取得指定位置的数据
  92. void* GetCVecElement(CVector * const pCV, size_t index)
  93. {
  94.     return (char*)pCV->pData + pCV->typeSize * index;
  95. }

  96. void CVectorDestroy(CVector * const pCV)
  97. {
  98.     free(pCV->pData);
  99.     free(pCV);
  100.     return;
  101. }

  102. void GetCodLeafAndCam(CVector * const pCodLeafFacet, const Para& pPara)
  103. {
  104.     float xMax = 3*18;
  105.     float yMax = 3*53;
  106.     //得到y轴按pPara.sideLength分割的每个微平面的中心点坐标。
  107.     CVector *pYnCenter = CVectorCreate(sizeof(float));
  108.     for(float i = -yMax; i <= yMax; i += pPara.sideLength)
  109.     {
  110.         CVectorPushBack(pYnCenter, &i);
  111.     }

  112.     CVector *pXnCenter = CVectorCreate(sizeof(float));
  113.     for(float i = -xMax; i <= xMax; i += pPara.sideLength)
  114.     {
  115.         CVectorPushBack(pXnCenter, &i);
  116.     }

  117.     int numXnCenter = CVectorLength(pXnCenter);
  118.     int numYnCenter = CVectorLength(pYnCenter);
  119.     for ( int i = 0; i != numXnCenter; i++)
  120.     {
  121.         for ( int j = 0; j != numYnCenter; j++)
  122.         {
  123.             float tmpXn = *((float *)GetCVecElement(pXnCenter,i));
  124.             float tmpYn = *((float *)GetCVecElement(pYnCenter,j));
  125.             if (tmpXn*tmpXn/pPara.ellipseA + tmpYn*tmpYn/pPara.ellipseB <= 1)
  126.             {   
  127.                 Point3D tmpPoint;
  128.                 tmpPoint.x = tmpXn;
  129.                 tmpPoint.y = tmpYn;
  130.                 tmpPoint.z = 0;
  131.                 CVectorPushBack(pCodLeafFacet, &tmpPoint);
  132.             }
  133.         }
  134.     }
  135.     CVectorDestroy(pXnCenter);
  136.     CVectorDestroy(pYnCenter);
  137. }

  138. void ReceiveFxArray(Para *pPara)
  139. {
  140.     CVector *pCodLeafCam = CVectorCreate( sizeof(Point3D) );
  141.     CVector *pCodLeafFacet = CVectorCreate( sizeof(Point3D) );
  142.     GetCodLeafAndCam(pCodLeafFacet, *pPara);
  143.     CVectorDestroy(pCodLeafFacet);
  144. }

  145. int main()
  146. {
  147.     Para tmpPara;
  148.     tmpPara.sideLength = 0.25;
  149.     tmpPara.ellipseA = 324;
  150.     tmpPara.ellipseB = 2704;

  151.     clock_t startTime = clock();
  152.     for(int i = 0; i < 600; i++)
  153.     {
  154.         ReceiveFxArray(&tmpPara);
  155.     }
  156.     clock_t elapsedTime = clock() - startTime;
  157.     printf("elapsedTime = %.5f\n",(float)elapsedTime/CLOCKS_PER_SEC);

  158.     system("pause");
  159.     return 0;
  160. }
复制代码
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册

zzczczxczxczx