Здравствуйте.
Мне для диплома нужны расчёты на GPU, но вот сделать их действительно быстрыми не получается.
У меня задача моделирования частиц которые взаимодействуют каждая с каждой, т.е. сложность N^2 где N количество частиц.
Демка nVidia nBodyCS на моей видюхе считает 64k точек с fps ~4
Моя же программа считает 20k точек с fps ~1
nBodyCS написана на DirectCompute и взаимодействи между частицами у меня сложнее поэтому напрямую сравнивать конечно нельзя но всё же.
Как сделать быстрее?
вот код:
- инициализация буферов
- clmemFV = clCreateBuffer(GPUContext, CL_MEM_READ_ONLY, sizeof(cl_float4)*m_paramColculate.countSteps*m_countParticleAddedAtStap, NULL, &err);
- clmemFC = clCreateBuffer(GPUContext, CL_MEM_READ_ONLY, sizeof(cl_float4)*m_paramColculate.countSteps*m_countParticleAddedAtStap, NULL, &err);
- clmemFA = clCreateBuffer(GPUContext, CL_MEM_READ_ONLY, sizeof(cl_float4)*m_paramColculate.countSteps*m_countParticleAddedAtStap, NULL, &err);
- clmemSV = clCreateBuffer(GPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float4)*m_paramColculate.countSteps*m_countParticleAddedAtStap, NULL, &err);
- clmemSC = clCreateBuffer(GPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float4)*m_paramColculate.countSteps*m_countParticleAddedAtStap, NULL, &err);
- clmemSA = clCreateBuffer(GPUContext, CL_MEM_WRITE_ONLY, sizeof (cl_float4)*m_paramColculate.countSteps*m_countParticleAddedAtStap, NULL, &err);
- расчёт:
- err += clEnqueueWriteBuffer ( GPUCommandQueue, clmemFV, CL_TRUE, 0, sizeof(cl_float4)*countPart, m_Velocity[m1s], 0, NULL, NULL);
- err += clEnqueueWriteBuffer ( GPUCommandQueue, clmemFC, CL_TRUE, 0, sizeof(cl_float4)*countPart, m_Coordinate[m1s], 0, NULL, NULL);
- err += clEnqueueWriteBuffer ( GPUCommandQueue, clmemFA, CL_TRUE, 0, sizeof(cl_float4)*countPart, m_Acceleration_m1[m1s], 0, NULL, NULL);
- err += clSetKernelArg(clTest, 0, sizeof(cl_mem), (void*)&clmemFV);
- err += clSetKernelArg(clTest, 1, sizeof(cl_mem), (void*)&clmemFC);
- err += clSetKernelArg(clTest, 2, sizeof(cl_mem), (void*)&clmemFA);
- err += clSetKernelArg(clTest, 3, sizeof(cl_mem), (void*)&clmemSV);
- err += clSetKernelArg(clTest, 4, sizeof(cl_mem), (void*)&clmemSC);
- err += clSetKernelArg(clTest, 5, sizeof(cl_mem), (void*)&clmemSA);
- err += clSetKernelArg(clTest, 6, sizeof(cl_int), (void*)&countPart);
- err += clSetKernelArg(clTest, 7, sizeof(cl_float), (void*)&t);
- err += clSetKernelArg(clTest, 8, sizeof(cl_float4), (void*)&VECTOR4( m_paramColculate.directionE.x*m_paramColculate.E,
- m_paramColculate.directionE.y*m_paramColculate.E,
- m_paramColculate.directionE.z*m_paramColculate.E,
- 0.0));
- err += clSetKernelArg(clTest, 9, sizeof(cl_float4), (void*)&VECTOR4( m_paramColculate.directionB.x*m_paramColculate.B,
- m_paramColculate.directionB.y*m_paramColculate.B,
- m_paramColculate.directionB.z*m_paramColculate.B,
- 0.0));
- err += clSetKernelArg(clTest, 10, sizeof(cl_float), (void*)&k);
- //***********************EXECUTING KERNEL**********************
- WorkSize[0] = m_countParticle;
- cl_event events[1];
- err += clEnqueueNDRangeKernel(GPUCommandQueue, clTest, 1, NULL, WorkSize, NULL, 0, NULL, &events[0]);
- err += clWaitForEvents(1, &events[0]);
- err += clReleaseEvent(events[0]);
- //***********************READING RESULTS**********************
- err += clEnqueueReadBuffer(GPUCommandQueue, clmemSV, CL_TRUE, 0, sizeof(cl_float4)*countPart, m_Velocity[cs], 0, NULL, NULL);
- err += clEnqueueReadBuffer(GPUCommandQueue, clmemSC, CL_TRUE, 0, sizeof(cl_float4)*countPart, m_Coordinate[cs], 0, NULL, NULL);
- err += clEnqueueReadBuffer(GPUCommandQueue, clmemSA, CL_TRUE, 0, sizeof(cl_float4)*countPart, m_Acceleration_m1[cs], 0, NULL, NULL);
- и ядро:
- inline float4 F( float4 currentV, float4 E, float4 B)
- {
- float4 a;
- float EV = (E.x*currentV.x + E.y*currentV.y + E.z*currentV.z)/8.9875517e+16f;
- float rel = -1.0f*sqrt(1.0f - (currentV.x*currentV.x + currentV.y*currentV.y + currentV.z*currentV.z)/8.9875517e+16f)*1.756e+11f;
- a.x = ( E.x + currentV.y*B.z - currentV.z*B.y - currentV.x*EV ) * rel;
- a.y = ( E.y + currentV.z*B.x - currentV.x*B.z - currentV.y*EV ) * rel;
- a.z = ( E.z + currentV.x*B.y - currentV.y*B.x - currentV.z*EV ) * rel;
- return a;
- }
- inline float4 mul(float4 A, float4 B)
- {
- float4 tmp;
- tmp.x = A.y*B.z - A.z*B.y;
- tmp.y = A.z*B.x - A.x*B.z;
- tmp.z = A.x*B.y - A.y*B.x;
- tmp.w = 0.0f;
- return tmp;
- }
- __kernel void computeRungeKutta2( __global float4* FV,
- __global float4* FC,
- __global float4* FA,
- __global float4* SV,
- __global float4* SC,
- __global float4* SA,
- const int size,
- const float t,
- const float4 E,
- const float4 B,
- const float k
- )
- {
- int index = get_global_id(0);
- if(index >= size) return;
- const float c = 2.99792458e+8f;
- const float PI = 3.141592653589f;
- float4 f0, f1, tmpV, Ecs = {0.0f, 0.0f, 0.0f, 0.0f}, Bcs = {0.0f, 0.0f, 0.0f, 0.0f};
- for(int j = 0; j < size; j++)
- {
- if(index != j)
- {
- float4 Rj;
- Rj = FC[index] - FC[j];
- float Rjl = sqrt(Rj.x*Rj.x + Rj.y*Rj.y + Rj.z*Rj.z);
- float tmp = 4.0f*PI*8.8541878f*pow((Rjl - (Rj.x*FV[j].x + Rj.y*FV[j].y + Rj.z*FV[j].z)/c),3.0f)/(-1.602e-7f*k);
- float tmp2 = 1.0f - (FV[j].x*FV[j].x + FV[j].y*FV[j].y + FV[j].z*FV[j].z)/(c*c);
- float4 tmp1 = mul(Rj, mul((Rj - FV[j] * Rjl/c), FA[j]));
- float4 Ecsj;
- Ecsj = ((Rj - FV[j] * Rjl/c) * tmp2 + tmp1/(c*c))/ tmp;
- Ecs += Ecsj;
- Bcs += mul(Rj,Ecsj)/(c*Rjl);
- }
- }
- Bcs.w = 0.0f;
- Ecs.w = 0.0f;
- f0 = F( FV[index], E + Ecs, B + Bcs);
- f1 = F( FV[index] + f0*t, E + Ecs, B + Bcs);
- SA[index] = f0;
- tmpV = FV[index]+ (f0 + f1)*t/2.0f;
- SC[index] = FC[index] + tmpV*t;
- SV[index] = tmpV;
- }
всем спасибо.