Создается небольшой поект. Одной из задач проекта является морфинг поверхности, состоящей из полигонов. В связи с большим количеством точек поверхности и медленным выполнением кода было решено искать пути ускорения процесса. Переход на OpenMP дал прирост производительности порядка 2-х раз на 4-х ядерном процессоре Intel Q6800. В настоящее время эксперементирую с CUDA. Для сравнения скорости привожу кусок кода проверки попадания точек полигонов в усеченый конус.  

 

Данные о процессоре:

 

Number of CPU(s) One Physical Processor / 2 Cores / 2 Logical Processors / 64 bits

CPU Full Name Intel(R) Pentium(R) D CPU 3.20GHz

CPU Name Intel Pentium D 940

CPU Code Name Presler

Platform Name Socket 775 LGA

Revision B1

Technology 65 nm

Instructions MMX, SSE, SSE2, SSE3, ET64, VMX

Original Clock 3200 MHz

Original System Clock 200 MHz

Original Multiplier 16.0

CPU Clock 3201 MHz

System Clock 200.0 MHz

FSB 800.1 MHz

Core 0 Speed 3200.6 MHz

Core 0 Multiplier 16.0

Core 1 Speed 3200.6 MHz

Core 1 Multiplier 16.0

L1 Data Cache 2 x 16 KBytes

L1 Trace Cache 2 x 12 Kuops

L2 Cache 2 x 2048 Kbytes

 

Данные о видео:

 

Number of GPUs 2

Display Adapter NVIDIA GeForce 9400 GT

Video Processor GeForce 9400 GT

Video memory size 1024 MBytes

Adapter DAC Type Integrated RAMDAC

BIOS String Version 62.94.72.00.C1

BIOS Date 08/28/09

Display Drivers nv4_disp.dll

Driver Version 6.14.12.5721

Driver Date 2010-06-07 23:57:00

Inf File Name oem104.inf

Inf Section Section005

Display Adapter BB Capture Driver

Driver Version 1.00

Driver Date 2008-05-19 04:30:53

Inf File Name oem42.inf

Inf Section bbcap

 

В таблице приведены результаты.

Первый столбик: 0 без CUDA/1 с CUDA.

Второй столбик: количество полигонов.

Третий столбик: общее количество точек.

Четвертый столбик: затраченое время в тиках процессора.

Как видно из таблицы, при изменении количества точек в полигоне, среднее время, затраченное на одну точку остается постоянным (примерно 3100-3200 на точку), а с использованием CUDA при увеличении количества точек, имеется прирост производительности до 10 раз(около 400 на точку).

 

0   16   18512   59129664

0   16   18512   58459152

0   16   18512   58712096

0   16   18512   60967168

0   16   18512   57244720

 

1   16   18512   502950240 (первый длительный цикл с CUDA связан с инициализацией GPU)

1   16   18512   52466864

1   16   18512   53369744

1   16   18512   50947872

1   16   18512   49795088

 

0   16   69872   224359936

0   16   69872   216826112

0   16   69872   218405712

0   16   69872   228633536

0   16   69872   225921376

 

1   16   69872   510279616

1   16   69872   62507984

1   16   69872   59920208

1   16   69872   58379904

1   16   69872   72676896

 

0   16   270992   865252016

0   16   270992   871903296

0   16   270992   875956464

0   16   270992   866460912

0   16   270992   871480384

 

1   16   270992   546036912

1   16   270992   102836928

1   16   270992   95773008

1   16   270992   101953184

1   16   270992   96566416

 

 

 

 

Код Microsoft Visual C++ 2008:

 

__int64 t1,t2,t3;

t2=0;

t3=0;

for (int i=0; ipMesh.size(); i++){

Model->pMesh[i]->GetVertexBuffer(&mVB);

mVB->Lock( 0, 0, ( void** )&m, D3DLOCK_READONLY );

DWORD i2=Model->pMesh[i]->GetNumVertices();

for (DWORD i1=0; i1

{

p[i1].x=m[i1].p.x;

p[i1].y=m[i1].p.y;

p[i1].z=m[i1].p.z;

t3++;

}

 

t1=__rdtsc();

if (Model->CUDA==TRUE)

{

// cuda

if (Cuda_PoligonInstr(i2, p, _matr1, _a, _d)==1)

_Count.push_back(i);

}

else

{

for (int i1=0; i1

D3DXVECTOR3 _vecSource;

D3DXVec3TransformCoord(&_vecSource,&m[i1].p,&Matr1);

if (_vecSource.z  > 0 ){

D3DXVECTOR2 _l=D3DXVECTOR2(_vecSource.x, _vecSource.y);

float l=D3DXVec2Length(&_l);

float l1=_a*_vecSource.z+_d;

if (l

_Count.push_back(i);

i1=i2;

}

}

}

}

t2+=__rdtsc()-t1;

 

mVB->Unlock();

mVB->Release();

m=NULL;

}

delete(p);

 

//////////////////////

std::ofstream out;

 

out.open("prim.txt",std::ios.app);

out

 

out.close();

//////////////////////

 

 

Код CUDA:

__global__ void PoligonInstrKernel(float3* point, float* matr, float a, float r, int* len)

{

unsigned int x1 = threadIdx.x+blockIdx.x*blockDim.x;

float3 p;

p.x = matr[0]*point[x1].x+matr[1]*point[x1].y+matr[2]*point[x1].z+matr[12];

p.y = matr[4]*point[x1].x+matr[5]*point[x1].y+matr[6]*point[x1].z+matr[13];

p.z = matr[8]*point[x1].x+matr[9]*point[x1].y+matr[10]*point[x1].z+matr[14];

point[x1]=p;

if (point[x1].z>0)

if (x1

{

float l=sqrt(point[x1].x*point[x1].x+point[x1].y*point[x1].y);

float l1=a*point[x1].z+r;

if (l

}

}

 

 

_declspec(dllexport) int Cuda_PoligonInstr(int len, float3* point, float* matr, float a, float r)

{

float3* point1 = new float3[len];

float3* devpoint;

cutilSafeCall(cudaMalloc((void**)&devpoint, sizeof(float3) * len));

cutilSafeCall( cudaMemset(devpoint, 0, sizeof(float3) * len) );

cutilSafeCall(cudaMemcpy(devpoint, point, sizeof(float3) * len, cudaMemcpyHostToDevice));

 

float* devmatr;

cutilSafeCall(cudaMalloc((void**)&devmatr, sizeof(float) * 16));

cutilSafeCall(cudaMemcpy(devmatr, matr, sizeof(float) * 16, cudaMemcpyHostToDevice));

int* l= new int[2];

l[0]=0;

l[1]=len;

int* devlen;

cutilSafeCall(cudaMalloc((void**)&devlen, sizeof(int)*2));

cutilSafeCall(cudaMemcpy(devlen, l, sizeof(int)*2, cudaMemcpyHostToDevice));

 

dim3 gridSize = dim3(1, 1, 1);

dim3 blockSize  = dim3(len, 1, 1);

 

PoligonInstrKernel(devpoint, devmatr, a, r, devlen);

 

cudaEvent_t syncEvent;

cudaEventCreate(&syncEvent);

cudaEventRecord(syncEvent, 0);

cudaEventSynchronize(syncEvent);

cudaMemcpy(point1, devpoint, sizeof(float3) * len, cudaMemcpyDeviceToHost);

 

cutilSafeCall(cudaMemcpy(l, devlen, sizeof(float)*2, cudaMemcpyDeviceToHost));

cudaFree(devpoint);

cudaFree(devmatr);

delete point1;

 

int _a;

if (l[0]>0) _a=1; else _a=0;

delete l;

return _a;

}