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