Создается небольшой поект. Одной из задач проекта является морфинг поверхности, состоящей из полигонов. В связи с большим количеством точек поверхности и медленным выполнением кода было решено искать пути ускорения процесса. Переход на 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; }