मैंने ओपनसीएल सीखना शुरू कर दिया है और मैं वर्तमान में परीक्षण करने की कोशिश करता हूं कि मैं एक साधारण कंकाल एनीमेशन एल्गोरिदम के लिए प्रदर्शन में कितना सुधार कर सकता हूं। ऐसा करने के लिए मैंने एक प्रोग्राम लिखा है जो यादृच्छिक रूप से जेनरेट किए गए कोने और रूपांतरण मैट्रिक्स से दो बार कंकाल एनीमेशन करता है, एक बार एसएसई-अनुकूलित रैखिक बीजगणित लाइब्रेरी के साथ सादे सी ++ में, और एक बार GPU पर अपने स्वयं के ओपनसीएल कर्नेल का उपयोग करके (मैं एक परीक्षण कर रहा हूं एनवीडिया जीटीएक्स 460)।ओपनसीएल प्रदर्शन अनुकूलन
मैंने एक साधारण कर्नेल के साथ शुरू किया जहां प्रत्येक कार्य-वस्तु वास्तव में एक वर्टेक्स को बदलती है, जिसमें सभी मान वैश्विक स्मृति से पढ़े जाते हैं। क्योंकि मैं इस कर्नेल के प्रदर्शन से संतुष्ट नहीं था, इसलिए मैंने थोड़ा अनुकूलित करने की कोशिश की। मेरे वर्तमान गिरी इस तरह दिखता है:
inline float4 MultiplyMatrixVector(float16 m, float4 v)
{
return (float4) (
dot(m.s048C, v),
dot(m.s159D, v),
dot(m.s26AE, v),
dot(m.s37BF, v)
);
}
kernel void skelanim(global const float16* boneMats, global const float4* vertices, global const float4* weights, global const uint4* indices, global float4* resVertices)
{
int gid = get_global_id(0);
int lid = get_local_id(0);
local float16 lBoneMats[NUM_BONES];
async_work_group_copy(lBoneMats, boneMats, NUM_BONES, 0);
barrier(CLK_LOCAL_MEM_FENCE);
for (int i = 0 ; i < NUM_VERTICES_PER_WORK_ITEM ; i++) {
int vidx = gid*NUM_VERTICES_PER_WORK_ITEM + i;
float4 vertex = vertices[vidx];
float4 w = weights[vidx];
uint4 idx = indices[vidx];
resVertices[vidx] = (MultiplyMatrixVector(lBoneMats[idx.x], vertex * w.x)
+ MultiplyMatrixVector(lBoneMats[idx.y], vertex * w.y)
+ MultiplyMatrixVector(lBoneMats[idx.z], vertex * w.z)
+ MultiplyMatrixVector(lBoneMats[idx.w], vertex * w.w));
}
}
अब मैं प्रति काम आइटम कोने की एक निरंतर संख्या की प्रक्रिया, और मैं एक काम के आइटम, जो मैं नेतृत्व करेंगे मानना था के लिए केवल एक बार स्थानीय स्मृति में सब हड्डी मैट्रिक्स prefetch बेहतर प्रदर्शन करने के लिए क्योंकि कई चरम के लिए मैट्रिस को बाद में स्थानीय स्थानीय स्मृति से पढ़ा जा सकता है। दुर्भाग्यवश, यह कर्नेल मेरे पहले प्रयास से भी बदतर है, और सीपीयू-केवल कार्यान्वयन से भी बदतर है।
प्रदर्शन के साथ इतना बुरा प्रदर्शन क्यों होना चाहिए?
यह, यहाँ मदद करता है कि कैसे मैं गिरी निष्पादित है:
#define NUM_BONES 50
#define NUM_VERTICES 30000
#define NUM_VERTICES_PER_WORK_ITEM 100
#define NUM_ANIM_REPEAT 1000
uint64_t PerformOpenCLSkeletalAnimation(Matrix4* boneMats, Vector4* vertices, float* weights, uint32_t* indices, Vector4* resVertices)
{
File kernelFile("/home/alemariusnexus/test/skelanim.cl");
char opts[256];
sprintf(opts, "-D NUM_VERTICES=%u -D NUM_REPEAT=%u -D NUM_BONES=%u -D NUM_VERTICES_PER_WORK_ITEM=%u", NUM_VERTICES, NUM_ANIM_REPEAT, NUM_BONES, NUM_VERTICES_PER_WORK_ITEM);
cl_program prog = BuildOpenCLProgram(kernelFile, opts);
cl_kernel kernel = clCreateKernel(prog, "skelanim", NULL);
cl_mem boneMatBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_BONES*sizeof(Matrix4), boneMats, NULL);
cl_mem vertexBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*sizeof(Vector4), vertices, NULL);
cl_mem weightBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*4*sizeof(float), weights, NULL);
cl_mem indexBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*4*sizeof(uint32_t), indices, NULL);
cl_mem resVertexBuf = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, NUM_VERTICES*sizeof(Vector4), NULL, NULL);
uint64_t s, e;
s = GetTickcount();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &boneMatBuf);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &vertexBuf);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &weightBuf);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &indexBuf);
clSetKernelArg(kernel, 4, sizeof(cl_mem), &resVertexBuf);
size_t globalWorkSize[] = { NUM_VERTICES/NUM_VERTICES_PER_WORK_ITEM };
size_t localWorkSize[] = { NUM_BONES };
for (size_t i = 0 ; i < NUM_ANIM_REPEAT ; i++) {
clEnqueueNDRangeKernel(cq, kernel, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
}
clEnqueueReadBuffer(cq, resVertexBuf, CL_TRUE, 0, NUM_VERTICES*sizeof(Vector4), resVertices, 0, NULL, NULL);
e = GetTickcount();
return e-s;
}
मुझे लगता है कि वहाँ अधिक चीजें हैं जो अनुकूलित किया जा सकता है, हो सकता है अन्य वैश्विक एक साथ पढ़ता के कुछ batching रहे हैं, लेकिन पहली बार मैं वास्तव में की तरह होगा यह जानने के लिए कि यह पहला अनुकूलन क्यों काम नहीं करता था।
मैं प्रदर्शन के बारे में पता नहीं है, लेकिन तुम क्या कर रहे अपरिभाषित परिणाम हो रहा है । आप एक अवरोध के बाद एक async_copy ऑपरेशन का उपयोग करें। अवरोध एसिंक कॉपी को समाप्त करने की प्रतीक्षा नहीं करेगा - जैसे ही सभी कार्य आइटम उस बिंदु तक पहुंच जाएंगे। Spec के अनुसार, आपको async_copy के बाद अपने कर्नेल में wait_group_events फ़ंक्शन का उपयोग करना होगा, या परिणाम अपरिभाषित हैं। यह समझ में आता है, क्योंकि async_copy हो रहा है जबकि शेष कर्नेल निष्पादित हो रहा है, इसलिए प्रतीक्षा_group_events कर्नेल को यह सुनिश्चित करने के लिए मजबूर करेगा कि स्मृति प्रतिलिपि पूरी हो गई है। –