`
schi
  • 浏览: 200236 次
  • 性别: Icon_minigender_1
  • 来自: 上海
社区版块
存档分类
最新评论

Maya plugin with OpenCL C++ example

阅读更多
GPU, Python and Maya使用的是python和pyopencl,这次我们使用C++来实现。

在GPU上面运行的kerne代码
// kerne code
const char kerneSrc[] = "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n"
"__kernel void ytwist(__global const double4 *pos,\n"
"__global double4 *newPos,\n"
"double magnitude,\n"
"float envelope)\n"
"{\n"
"    int gid = get_global_id(0);\n"
"    newPos[gid] = pos[gid];\n"
"    float ff = magnitude * pos[gid].y * envelope;\n"
"    if (ff != 0.f)\n"
"    {\n"
"        float cct = cos(ff);\n"
"        float cst = sin(ff);\n"
"        newPos[gid].x = pos[gid].x * cct - pos[gid].z * cst;\n"
"        newPos[gid].z = pos[gid].x * cst + pos[gid].z * cct;\n"
"    }\n"
"}";

对比python的版本你会发现有点不一样,python的版本使用的float和float4,而这次我们使用double和double4,为了能使用double类型,我们需要OpenCL启动对double的支持
#pragma OPENCL EXTENSION cl_khr_fp64: enable


因为我们使用的是OpenCL的c++蒙板,为了能捕捉错误需要
#define __CL_ENABLE_EXCEPTIONS


查询并选择一个平台
// find a opencl platform
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.size() == 0)
{
    MGlobal::displayError("Platform size 0");
    return MS::kFailure;
}

什么是平台(platform)?
在桌面电脑上是Nvidia,AMD,Intel。

创建context,因为我们需要的是GPU,所以使用CL_DEVICE_TYPE_GPU
// create a context
cl_context_properties properties[] = { CL_CONTEXT_PLATFORM,
                                       (cl_context_properties)(platforms[0])(), 
                                       0 };
cl::Context context(CL_DEVICE_TYPE_GPU, properties);


获取所有显卡设备
// get devices
std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();


从kernel代码创建GPU程序并生成
// create a program
cl::Program::Sources source(1, std::make_pair(kerneSrc, strlen(kerneSrc)));
cl::Program program(context, source);
program.build(devices);


获取所有的顶点位移,这里我使用的是cl_double4类型,它是OpenCL里定义的,但需要注意的是它在GNUC编译器下才跟maya里的MPoint的一样可以使用x,y,z,w属性,而在MSVC却只能使用s[i]的形式
// position data
int numPoints = iter.count();
cl_double4 *allPos = new cl_double4[numPoints];
cl_double4 *outPos = new cl_double4[numPoints];

// iterate through each point in the geometry
// get all position data
int i = 0;
for (; !iter.isDone(); iter.next()) {
    MPoint pt = iter.position();
    #if defined(__GNUC__)
    allPos[i].x = pt.x;
    allPos[i].y = pt.y;
    allPos[i].z = pt.z;
    allPos[i].w = pt.w;
    #else
    // for MSVC
    allPos[i].s[0] = pt.x;
    allPos[i].s[1] = pt.y;
    allPos[i].s[2] = pt.z;
    allPos[i].s[3] = pt.w;
    #endif

    i++;
}


创建GPU内存对象,并把数据传到GPU内存上
// create buffers
cl::Buffer posBuffer = cl::Buffer(context,
    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    numPoints * sizeof(cl_double4), allPos);
cl::Buffer nposBuffer = cl::Buffer(context,
    CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
    numPoints * sizeof(cl_double4), outPos);


创建我们的kernel,并设置kernel需要的参数,kernel可以理解为GPU程序里面的一个函数
// create a kernel
cl::Kernel kernel(program, "ytwist", &err);
// set kernel argmuments
kernel.setArg(0, posBuffer);
kernel.setArg(1, nposBuffer);
kernel.setArg(2, magnitude);
kernel.setArg(3, env);


创建commandQueue,并执行我们的kernel,kernel只能通过commandQueue来执行,线程数也是通过commandQueue来指定的
// create a command queue
cl::Event event;
cl::CommandQueue queue(context, devices[0], 0, &err);
// execut the kernel
queue.enqueueNDRangeKernel(kernel,
    cl::NullRange,
    cl::NDRange(numPoints),
    cl::NullRange,
    NULL,
    &event);
event.wait();


所有线程都运行完了,我们需要把数据从GPU内存读回来,然后再设置所有的顶点位置。
// read the data back
queue.enqueueReadBuffer(nposBuffer, CL_TRUE, 0, numPoints * sizeof(cl_double4), outPos);

// set the new positions
i = 0;
iter.reset();
for (; !iter.isDone(); iter.next()) {
    MPoint pt = iter.position();
    #if defined(__GNUC__)
    pt.x = outPos[i].x;
    pt.z = outPos[i].z;
    #else
    pt.x = outPos[i].s[0];
    pt.z = outPos[i].s[2];
    #endif
    iter.setPosition(pt);
    i++;
}


完整的代码
https://github.com/mackst/myRandomStuff/tree/master/plugins/yTwistNode_ocl
分享到:
评论

相关推荐

Global site tag (gtag.js) - Google Analytics