cuda 的 malloc和memcpy很慢，什么原因？

daxiaoyuyu · daxiaoyuyu

最近小弟真的是非常郁闷，发现在cuda中用malloc，memcpy，这样的函数效率非常低。各路高手走过路过不要错过，给小弟一点帮助吧，哪怕一点建议都行，不胜感激。

我的程序中需要在一张表中提出100个参数（这不是固定的），需要将椭圆分成n个面积为sideLength* sideLength的微小平面，并且把这些平面的中心坐标存放起来供给下一个函数使用。

由于sideLength和ellipseA， 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.047s，CPU快91倍。

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中的线程分配过多之后，反而更慢了？是不是分配的线程过多后，会产生更多的显存碎片，从而会更慢？

daxiaoyuyu · daxiaoyuyu

GPU代码：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
#include <stdlib.h>
#include <stdio.h>
#define PI 3.1415926535897
#define MIN(x, y) ((x)>(y)?(y):(x))
#define MINLEN 256
#define EXPANED_TOTLEN 1
#define CVPUSHBACK 1; //内存回退
#define CVSUCESS 0; //内存分配成功
struct Point3D
{
float x, y, z;
int attr;
};
struct Para
{
// x^2/g_ellipseA + y^2/g_ellipseB <= 1
float sideLength;
float ellipseA;
float ellipseB;
};
struct CVector
{
void *pData;//指向所需的数据
size_t len,totLen,typeSize;
};
//-------c语言vector版--------
__device__ CVector* CVectorCreate(size_t const size)
{
CVector *pCV = (CVector *)malloc(sizeof(CVector));
if (!pCV)
{
return NULL;
}
pCV->pData = malloc( MINLEN * size );
if (!pCV->pData)
{
free(pCV);
return NULL;
}
pCV->len = 0;
pCV->totLen = MINLEN;
pCV->typeSize = size;
return pCV;
}
__device__ int CVectorPushBack(CVector * const pCV, void const * const pMem)
{
if (pCV->len >= pCV->totLen)
{
void *pDataSave = pCV->pData;
int totLenSave = pCV->totLen;
pCV->totLen <<= EXPANED_TOTLEN;
pCV->pData = malloc(pCV->totLen * pCV->typeSize);
if ( !pCV->pData )//如果分配新的内存不成功，滚回。
{
pCV->pData = pDataSave;
pCV->totLen = totLenSave;
return CVPUSHBACK;
}
//把原来的数据复制进去
memcpy(pCV->pData, pDataSave, totLenSave * pCV->typeSize);
free(pDataSave);
}
//把新数据复制进去。
memcpy( (char *)pCV->pData + pCV->len * pCV->typeSize, pMem, pCV->typeSize );
pCV->len ++;
return CVSUCESS;
}
__device__ size_t CVectorLength(CVector const * const pCV)
{
return pCV->len;
}
__device__ CVector * CVectorCopy(CVector const * const pCV)
{
CVector * pNewCV = (CVector *)malloc(sizeof(CVector));
if (!pNewCV)
{
return NULL;
}
pNewCV->pData = malloc( pCV->typeSize * pCV->totLen);
if ( !pNewCV->pData )
{
free(pNewCV);
return NULL;
}
memcpy(pNewCV->pData, pCV->pData, pCV->typeSize * pCV->totLen);
pNewCV->len = pCV->len;
pNewCV->totLen = pCV->totLen;
pNewCV->typeSize = pCV->typeSize;
return pNewCV;
}
//取得指定位置的指针
__device__ void* GetCVecElement(CVector * const pCV, size_t index)
{
return (char*)pCV->pData + pCV->typeSize * index;
}
__device__ void CVectorDestroy(CVector * const pCV)
{
free(pCV->pData);
free(pCV);
return;
}
//---获取椭圆范围内微平面的中心坐标点。
__device__ void GetCodLeafAndCam(CVector * const pCodLeafFacet, const Para& para)
{
//先提取出 xMax * yMax 矩形的微微平面中心坐标点。
//再判断这些坐标点是否落在椭圆内，从而提取出椭圆范围内的中心坐标点。
float xMax = 3*18;
float yMax = 3*53;
//按sideLength分割 xMax。
CVector *pXnCenter = CVectorCreate(sizeof(float));
for(float i = -xMax; i <= xMax; i += para.sideLength)
{
CVectorPushBack(pXnCenter, &i);
}
//按sideLength分割 yMax
CVector *pYnCenter = CVectorCreate(sizeof(float));
for(float i = -yMax; i <= yMax; i += para.sideLength)
{
CVectorPushBack(pYnCenter, &i);
}
//提取出坐标点。
int numXnCenter = CVectorLength(pXnCenter);
int numYnCenter = CVectorLength(pYnCenter);
for ( int i = 0; i != numXnCenter; i++)
{
for ( int j = 0; j != numYnCenter; j++)
{
float tmpXn = *((float *)GetCVecElement(pXnCenter,i));
float tmpYn = *((float *)GetCVecElement(pYnCenter,j));
if (tmpXn*tmpXn/para.ellipseA + tmpYn*tmpYn/para.ellipseB <= 1)
{
Point3D tmpPoint;
tmpPoint.x = tmpXn;
tmpPoint.y = tmpYn;
tmpPoint.z = 0;
CVectorPushBack(pCodLeafFacet, &tmpPoint);
}
}
}
CVectorDestroy(pXnCenter);
CVectorDestroy(pYnCenter);
}
__global__ static void ReceiveFxArray(Para *pcuPara)
{
CVector *pCodLeafCam = CVectorCreate( sizeof(Point3D) );
CVector *pCodLeafFacet = CVectorCreate( sizeof(Point3D) );
GetCodLeafAndCam(pCodLeafFacet, *pcuPara);
CVectorDestroy(pCodLeafFacet);
}
int main()
{
size_t memSize;
cudaDeviceSetLimit(cudaLimitMallocHeapSize, 512*1024*1024);
cudaDeviceGetLimit(&memSize, cudaLimitMallocHeapSize);
Para tmpPara;
tmpPara.sideLength = 0.25;
tmpPara.ellipseA = 324;
tmpPara.ellipseB = 2704;
Para *pcuPara;
cudaMalloc( (void**)&pcuPara, sizeof(Para) );
cudaMemcpy( pcuPara, &tmpPara, sizeof(Para), cudaMemcpyHostToDevice );
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
ReceiveFxArray<<<6,100,0>>> ( pcuPara);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("elapsedTime = %.5f\n",elapsedTime);
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaFree( pcuPara );
system("pause");
return 0;
}

复制代码

daxiaoyuyu · daxiaoyuyu

CPU代码：

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <time.h>
#define PI 3.1415926535897
#define MIN(x, y) ((x)>(y)?(y):(x))
#define MINLEN 256
#define EXPANED_TOTLEN 1
#define CVPUSHBACK 1; //内存回退
#define CVSUCESS 0; //内存分配成功
struct Point3D
{
double x, y, z;
};
struct Para
{
// x^2/g_ellipseA + y^2/g_ellipseB <= 1
float sideLength;
float ellipseA;
float ellipseB;
};
struct CVector
{
void *pData;//指向所需的数据
size_t len,totLen,typeSize;//pData里面有多少个数据，最多能容纳多少个数据，数据类型的字节数。
};
CVector* CVectorCreate(size_t const size)
{
CVector *pCV = (CVector *)malloc(sizeof(CVector));//sizeof是用来确定类型的字节数的。
if (!pCV) //如果内存分配不成功，则 pCV 为NULL，即为0;
{
return NULL;
}
pCV->pData = malloc( MINLEN * size );
if (!pCV->pData)
{
free(pCV);
return NULL;
}
pCV->len = 0;
pCV->totLen = MINLEN;
pCV->typeSize = size;
return pCV;
}
int CVectorPushBack(CVector * const pCV, void const * const pMem)
{
if (pCV->len >= pCV->totLen)
{
void *pDataSave = pCV->pData;
int totLenSave = pCV->totLen;
pCV->totLen <<= EXPANED_TOTLEN; //往前移位，意味着扩大一倍。
pCV->pData = malloc(pCV->totLen * pCV->typeSize);
if ( !pCV->pData )//如果分配新的内存不成功，滚回。
{
pCV->pData = pDataSave;
pCV->totLen = totLenSave;
return CVPUSHBACK;
}
//把原来的数据复制进去
memcpy(pCV->pData, pDataSave, totLenSave * pCV->typeSize);
free(pDataSave);
}
//把新数据复制进去。
memcpy( (char *)pCV->pData + pCV->len * pCV->typeSize, pMem, pCV->typeSize );
pCV->len ++;
return CVSUCESS;
}
size_t CVectorLength(CVector const * const pCV)
{
return pCV->len;
}
CVector * CVectorCopy(CVector const * const pCV)
{
CVector * pNewCV = (CVector *)malloc(sizeof(CVector));
if (!pNewCV)
{
return NULL;
}
pNewCV->pData = malloc( pCV->typeSize * pCV->totLen);
if ( !pNewCV->pData )
{
free(pNewCV);
return NULL;
}
memcpy(pNewCV->pData, pCV->pData, pCV->typeSize * pCV->totLen);
pNewCV->len = pCV->len;
pNewCV->totLen = pCV->totLen;
pNewCV->typeSize = pCV->typeSize;
return pNewCV;
}
//取得指定位置的数据
void* GetCVecElement(CVector * const pCV, size_t index)
{
return (char*)pCV->pData + pCV->typeSize * index;
}
void CVectorDestroy(CVector * const pCV)
{
free(pCV->pData);
free(pCV);
return;
}
void GetCodLeafAndCam(CVector * const pCodLeafFacet, const Para& pPara)
{
float xMax = 3*18;
float yMax = 3*53;
//得到y轴按pPara.sideLength分割的每个微平面的中心点坐标。
CVector *pYnCenter = CVectorCreate(sizeof(float));
for(float i = -yMax; i <= yMax; i += pPara.sideLength)
{
CVectorPushBack(pYnCenter, &i);
}
CVector *pXnCenter = CVectorCreate(sizeof(float));
for(float i = -xMax; i <= xMax; i += pPara.sideLength)
{
CVectorPushBack(pXnCenter, &i);
}
int numXnCenter = CVectorLength(pXnCenter);
int numYnCenter = CVectorLength(pYnCenter);
for ( int i = 0; i != numXnCenter; i++)
{
for ( int j = 0; j != numYnCenter; j++)
{
float tmpXn = *((float *)GetCVecElement(pXnCenter,i));
float tmpYn = *((float *)GetCVecElement(pYnCenter,j));
if (tmpXn*tmpXn/pPara.ellipseA + tmpYn*tmpYn/pPara.ellipseB <= 1)
{
Point3D tmpPoint;
tmpPoint.x = tmpXn;
tmpPoint.y = tmpYn;
tmpPoint.z = 0;
CVectorPushBack(pCodLeafFacet, &tmpPoint);
}
}
}
CVectorDestroy(pXnCenter);
CVectorDestroy(pYnCenter);
}
void ReceiveFxArray(Para *pPara)
{
CVector *pCodLeafCam = CVectorCreate( sizeof(Point3D) );
CVector *pCodLeafFacet = CVectorCreate( sizeof(Point3D) );
GetCodLeafAndCam(pCodLeafFacet, *pPara);
CVectorDestroy(pCodLeafFacet);
}
int main()
{
Para tmpPara;
tmpPara.sideLength = 0.25;
tmpPara.ellipseA = 324;
tmpPara.ellipseB = 2704;
clock_t startTime = clock();
for(int i = 0; i < 600; i++)
{
ReceiveFxArray(&tmpPara);
}
clock_t elapsedTime = clock() - startTime;
printf("elapsedTime = %.5f\n",(float)elapsedTime/CLOCKS_PER_SEC);
system("pause");
return 0;
}

复制代码

		自动登录	找回密码
密码			立即注册


20 CUDA币	回复本帖可获得 20 CUDA币奖励! 每人限 5 次