std::string OpenCLParser::getTypedKernelLine(std::string line) {
std::string kernel_name = getKernelName(line);
boost::regex split_bgn(kernel_name);
boost::regex split_end("\\;");
boost::sregex_token_iterator it;
it = boost::sregex_token_iterator(line.begin(),line.end(), split_bgn, -1);
it++;
it++;
line = (*it);
it = boost::sregex_token_iterator(line.begin(),line.end(), split_end, -1);
line = *it;
return line;
}
// Read the kernels that this plan uses from file, and store into the plan
clfftStatus FFTAction::writeKernel( const clfftPlanHandle plHandle, const clfftGenerators gen, const FFTKernelSignatureHeader* data, const cl_context& context, const cl_device_id &device )
{
FFTRepo& fftRepo = FFTRepo::getInstance( );
std::string kernelPath = getKernelName(gen, plHandle, true);
// Logic to write string contents out to file
tofstreamRAII< std::ofstream, std::string > kernelFile( kernelPath.c_str( ) );
if( !kernelFile.get( ) )
{
std::cerr << "Failed to open kernel file for writing: " << kernelPath.c_str( ) << std::endl;
return CLFFT_FILE_CREATE_FAILURE;
}
std::string kernel;
OPENCL_V( fftRepo.getProgramCode( gen, data, kernel, device, context ), _T( "fftRepo.getProgramCode failed." ) );
kernelFile.get( ) << kernel << std::endl;
return CLFFT_SUCCESS;
}
//------------------------------------------------------------------------------
int main(int argc, char ** argv) {
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA |GLUT_DOUBLE | GLUT_DEPTH);
glutInitWindowSize(1024, 1024);
glutCreateWindow("OpenSubdiv ptexViewer");
int lmenu = glutCreateMenu(levelMenu);
for(int i = 1; i < 8; ++i){
char level[16];
sprintf(level, "Level %d\n", i);
glutAddMenuEntry(level, i);
}
int smenu = glutCreateMenu(schemeMenu);
glutAddMenuEntry("Catmark", 0);
glutAddMenuEntry("Bilinear", 1);
// Register Osd compute kernels
OpenSubdiv::OsdCpuKernelDispatcher::Register();
#if OPENSUBDIV_HAS_GLSL
OpenSubdiv::OsdGlslKernelDispatcher::Register();
#endif
#if OPENSUBDIV_HAS_OPENCL
OpenSubdiv::OsdClKernelDispatcher::Register();
#endif
#if OPENSUBDIV_HAS_CUDA
OpenSubdiv::OsdCudaKernelDispatcher::Register();
// Note: This function randomly crashes with linux 5.0-dev driver.
// cudaGetDeviceProperties overrun stack..?
cudaGLSetGLDevice( cutGetMaxGflopsDeviceId() );
#endif
int kmenu = glutCreateMenu(kernelMenu);
int nKernels = OpenSubdiv::OsdKernelDispatcher::kMAX;
for(int i = 0; i < nKernels; ++i)
if(OpenSubdiv::OsdKernelDispatcher::HasKernelType(
OpenSubdiv::OsdKernelDispatcher::KernelType(i)))
glutAddMenuEntry(getKernelName(i), i);
glutCreateMenu(menu);
glutAddSubMenu("Level", lmenu);
glutAddSubMenu("Scheme", smenu);
glutAddSubMenu("Kernel", kmenu);
glutAttachMenu(GLUT_RIGHT_BUTTON);
glutDisplayFunc(display);
glutReshapeFunc(reshape);
glutMouseFunc(mouse);
glutKeyboardFunc(keyboard);
glutMotionFunc(motion);
glewInit();
initGL();
for (int i = 1; i < argc; ++i) {
if (!strcmp(argv[i], "-d"))
g_level = atoi(argv[++i]);
else if (!strcmp(argv[i], "-c"))
g_repeatCount = atoi(argv[++i]);
else if (g_ptexColorFile == NULL)
g_ptexColorFile = argv[i];
else if (g_ptexDisplacementFile == NULL)
g_ptexDisplacementFile = argv[i];
else if (g_ptexOcclusionFile == NULL)
g_ptexOcclusionFile = argv[i];
}
if (g_ptexColorFile == NULL) {
printf("Usage: %s <color.ptx> [<displacement.ptx>] [<occlusion.ptx>] \n", argv[0]);
return 1;
}
createOsdMesh(g_level, g_kernel);
fitFrame();
glutIdleFunc(idle);
glutMainLoop();
quit();
}
//------------------------------------------------------------------------------
void
display() {
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glViewport(0, 0, g_width, g_height);
double aspect = g_width/(double)g_height;
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
gluPerspective(45.0, aspect, g_size*0.001f, g_size+g_dolly);
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
glTranslatef(-g_pan[0], -g_pan[1], -g_dolly);
glRotatef(g_rotate[1], 1, 0, 0);
glRotatef(g_rotate[0], 0, 1, 0);
glTranslatef(-g_center[0], -g_center[1], -g_center[2]);
glUseProgram(g_program);
{
// shader uniform setting
GLint position = glGetUniformLocation(g_program, "lightSource[0].position");
GLint ambient = glGetUniformLocation(g_program, "lightSource[0].ambient");
GLint diffuse = glGetUniformLocation(g_program, "lightSource[0].diffuse");
GLint specular = glGetUniformLocation(g_program, "lightSource[0].specular");
glProgramUniform4f(g_program, position, 0, 0.2f, 1, 0);
glProgramUniform4f(g_program, ambient, 0.4f, 0.4f, 0.4f, 1.0f);
glProgramUniform4f(g_program, diffuse, 0.3f, 0.3f, 0.3f, 1.0f);
glProgramUniform4f(g_program, specular, 0.2f, 0.2f, 0.2f, 1.0f);
GLint otcMatrix = glGetUniformLocation(g_program, "objectToClipMatrix");
GLint oteMatrix = glGetUniformLocation(g_program, "objectToEyeMatrix");
GLfloat modelView[16], proj[16], mvp[16];
glGetFloatv(GL_MODELVIEW_MATRIX, modelView);
glGetFloatv(GL_PROJECTION_MATRIX, proj);
multMatrix(mvp, modelView, proj);
glProgramUniformMatrix4fv(g_program, otcMatrix, 1, false, mvp);
glProgramUniformMatrix4fv(g_program, oteMatrix, 1, false, modelView);
}
GLuint bVertex = g_vertexBuffer->GetGpuBuffer();
glEnableVertexAttribArray(0);
glEnableVertexAttribArray(1);
glBindBuffer(GL_ARRAY_BUFFER, bVertex);
glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, 0);
glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, (float*)12);
glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, g_elementArrayBuffer->GetGlBuffer());
glPolygonMode(GL_FRONT_AND_BACK, g_wire==0 ? GL_LINE : GL_FILL);
// glPatchParameteri(GL_PATCH_VERTICES, 4);
// glDrawElements(GL_PATCHES, g_numIndices, GL_UNSIGNED_INT, 0);
glDrawElements(GL_LINES_ADJACENCY, g_elementArrayBuffer->GetNumIndices(), GL_UNSIGNED_INT, 0);
glUseProgram(0);
if (g_drawNormals)
drawNormals();
glDisableVertexAttribArray(0);
glDisableVertexAttribArray(1);
glBindBuffer(GL_ARRAY_BUFFER, 0);
glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, 0);
if (g_ptexDebug)
drawPtexLayout(g_ptexDebug-1);
glColor3f(1, 1, 1);
drawFmtString(10, 10, "LEVEL = %d", g_level);
drawFmtString(10, 30, "# of Vertices = %d", g_osdmesh->GetFarMesh()->GetNumVertices());
drawFmtString(10, 50, "KERNEL = %s", getKernelName(g_kernel));
drawFmtString(10, 70, "CPU TIME = %.3f ms", g_cpuTime);
drawFmtString(10, 90, "GPU TIME = %.3f ms", g_gpuTime);
g_fpsTimer.Stop();
drawFmtString(10, 110, "FPS = %3.1f", 1.0/g_fpsTimer.GetElapsed());
g_fpsTimer.Start();
drawFmtString(10, 130, "SUBDIVISION = %s", g_scheme==0 ? "CATMARK" : "BILINEAR");
drawString(10, g_height-10, "a: ambient occlusion on/off");
drawString(10, g_height-30, "c: color on/off");
drawString(10, g_height-50, "d: displacement on/off");
drawString(10, g_height-70, "e: show normal vector");
drawString(10, g_height-90, "f: fit frame");
drawString(10, g_height-110, "w: toggle wireframe");
drawString(10, g_height-130, "m: toggle vertex moving");
drawString(10, g_height-150, "s: bilinear / catmark");
drawString(10, g_height-170, "1-7: subdivision level");
glFinish();
glutSwapBuffers();
}
//------------------------------------------------------------------------------
int main(int argc, char ** argv) {
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA |GLUT_DOUBLE | GLUT_DEPTH);
glutInitWindowSize(1024, 1024);
glutCreateWindow("OpenSubdiv test");
std::string str;
if (argc > 1) {
std::ifstream ifs(argv[1]);
if (ifs) {
std::stringstream ss;
ss << ifs.rdbuf();
ifs.close();
str = ss.str();
g_defaultShapes.push_back(SimpleShape(str.c_str(), argv[1], kCatmark));
}
}
initializeShapes();
int smenu = glutCreateMenu(modelMenu);
for(int i = 0; i < (int)g_defaultShapes.size(); ++i){
glutAddMenuEntry( g_defaultShapes[i].name.c_str(), i);
}
int lmenu = glutCreateMenu(levelMenu);
for(int i = 1; i < 8; ++i){
char level[16];
sprintf(level, "Level %d\n", i);
glutAddMenuEntry(level, i);
}
// Register Osd compute kernels
OpenSubdiv::OsdCpuKernelDispatcher::Register();
OpenSubdiv::OsdGlslKernelDispatcher::Register();
#if OPENSUBDIV_HAS_OPENCL
OpenSubdiv::OsdClKernelDispatcher::Register();
#endif
#if OPENSUBDIV_HAS_CUDA
OpenSubdiv::OsdCudaKernelDispatcher::Register();
// Note: This function randomly crashes with linux 5.0-dev driver.
// cudaGetDeviceProperties overrun stack..?
cudaGLSetGLDevice( cutGetMaxGflopsDeviceId() );
#endif
int kmenu = glutCreateMenu(kernelMenu);
int nKernels = OpenSubdiv::OsdKernelDispatcher::kMAX;
for(int i = 0; i < nKernels; ++i)
if(OpenSubdiv::OsdKernelDispatcher::HasKernelType(
OpenSubdiv::OsdKernelDispatcher::KernelType(i)))
glutAddMenuEntry(getKernelName(i), i);
glutCreateMenu(menu);
glutAddSubMenu("Level", lmenu);
glutAddSubMenu("Model", smenu);
glutAddSubMenu("Kernel", kmenu);
glutAttachMenu(GLUT_RIGHT_BUTTON);
glutDisplayFunc(display);
glutReshapeFunc(reshape);
glutMouseFunc(mouse);
glutKeyboardFunc(keyboard);
glutMotionFunc(motion);
glewInit();
initGL();
const char *filename = NULL;
for (int i = 1; i < argc; ++i) {
if (!strcmp(argv[i], "-d"))
g_level = atoi(argv[++i]);
else if (!strcmp(argv[i], "-c"))
g_repeatCount = atoi(argv[++i]);
else
filename = argv[i];
}
modelMenu(0);
glutIdleFunc(idle);
glutMainLoop();
quit();
}
//------------------------------------------------------------------------------
void
display() {
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glViewport(0, 0, g_width, g_height);
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
glDisable(GL_LIGHTING);
glColor3f(1, 1, 1);
glBegin(GL_QUADS);
glColor3f(1, 1, 1);
glVertex3f(-1, -1, 1);
glVertex3f( 1, -1, 1);
glVertex3f( 1, 1, 1);
glVertex3f(-1, 1, 1);
glEnd();
double aspect = g_width/(double)g_height;
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
gluPerspective(45.0, aspect, 0.01, 500.0);
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
glTranslatef(-g_pan[0], -g_pan[1], -g_dolly);
glRotatef(g_rotate[1], 1, 0, 0);
glRotatef(g_rotate[0], 0, 1, 0);
glTranslatef(-g_center[0], -g_center[1], -g_center[2]);
glRotatef(-90, 1, 0, 0); // z-up model
GLuint bVertex = g_vertexBuffer->GetGpuBuffer();
glEnableVertexAttribArray(0);
glEnableVertexAttribArray(1);
glBindBuffer(GL_ARRAY_BUFFER, bVertex);
glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, 0);
glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, sizeof (GLfloat) * 6, (float*)12);
glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, g_elementArrayBuffer->GetGlBuffer());
GLenum primType = GL_LINES_ADJACENCY;
if (g_scheme == kLoop) {
primType = GL_TRIANGLES;
bindProgram(g_triFillProgram);
} else {
bindProgram(g_quadFillProgram);
}
if (g_wire > 0) {
glDrawElements(primType, g_elementArrayBuffer->GetNumIndices(), GL_UNSIGNED_INT, NULL);
}
if (g_wire == 0 || g_wire == 2) {
GLuint lineProgram = g_scheme == kLoop ? g_triLineProgram : g_quadLineProgram;
bindProgram(lineProgram);
GLuint fragColor = glGetUniformLocation(lineProgram, "fragColor");
if (g_wire == 2) {
glProgramUniform4f(lineProgram, fragColor, 0, 0, 0.5, 1);
} else {
glProgramUniform4f(lineProgram, fragColor, 1, 1, 1, 1);
}
glDrawElements(primType, g_elementArrayBuffer->GetNumIndices(), GL_UNSIGNED_INT, NULL);
}
glDisableVertexAttribArray(0);
glDisableVertexAttribArray(1);
glUseProgram(0);
if (g_drawNormals)
drawNormals();
if (g_drawCoarseMesh)
drawCoarseMesh(g_drawCoarseMesh);
glBindBuffer(GL_ARRAY_BUFFER, 0);
glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, 0);
glDisableClientState(GL_VERTEX_ARRAY);
if (g_drawHUD) {
glColor3f(1, 1, 1);
drawString(10, 10, "LEVEL = %d", g_level);
drawString(10, 30, "# of Vertices = %d", g_osdmesh->GetFarMesh()->GetNumVertices());
drawString(10, 50, "KERNEL = %s", getKernelName(g_kernel));
drawString(10, 70, "CPU TIME = %.3f ms", g_cpuTime);
drawString(10, 90, "GPU TIME = %.3f ms", g_gpuTime);
drawString(10, 110, "SUBDIVISION = %s", g_scheme==kBilinear ? "BILINEAR" : (g_scheme == kLoop ? "LOOP" : "CATMARK"));
drawString(10, g_height-30, "w: toggle wireframe");
drawString(10, g_height-50, "e: display normal vector");
drawString(10, g_height-70, "m: toggle vertex deforming");
drawString(10, g_height-90, "h: display control cage");
drawString(10, g_height-110, "n/p: change model");
drawString(10, g_height-130, "1-7: subdivision level");
drawString(10, g_height-150, "space: freeze/unfreeze time");
}
glFinish();
glutSwapBuffers();
}
bool OpenCLParser::convert(std::string fileNameIN, std::string fileNameOUT) {
if ( access( fileNameIN.c_str(), F_OK ) == -1 ) {
LOG(ERROR) << "kernel source file = '" << fileNameIN.c_str() << "' doesn't exist";
return false;
}
if ( access( fileNameIN.c_str(), R_OK ) == -1 ) {
LOG(ERROR) << "kernel source file = '" << fileNameIN.c_str() << "' isn't readable";
return false;
}
if ( access( fileNameOUT.c_str(), F_OK ) == 0 ) {
struct stat statIN;
if (stat(fileNameIN.c_str(), &statIN) == -1) {
perror(fileNameIN.c_str());
return false;
}
struct stat statOUT;
if (stat(fileNameOUT.c_str(), &statOUT) == -1) {
perror(fileNameOUT.c_str());
return false;
}
if ( statOUT.st_mtime > statIN.st_mtime ) {
DLOG(INFO) << "kernel source file = '" << fileNameOUT.c_str() << "' up-to-date";
return true;
}
}
std::ifstream file;
file.open(fileNameIN.c_str(), std::ifstream::in );
if ( ! file.is_open() ) {
LOG(ERROR) << "failed to open file = '" << fileNameIN.c_str() << "' for reading";
return false;
}
std::string line;
std::string kernel_buffer;
std::string kernel_name;
std::string kernel_type;
std::string kernel_name_typed;
std::string kernel_line_typed;
std::string kernel_modified;
std::string type_replace;
std::string stdOpenCL;
stdOpenCL += "// This file was auto-generated from file '" + fileNameIN + "' to conform to standard OpenCL\n";
bool recording = false;
while (std::getline(file, line)) {
if ( isAttributeLine(line) ) {
if ( recording ) {
recording = false;
}
kernel_name_typed = getTypedKernelName(line);
kernel_line_typed = "__kernel void " + kernel_name_typed + getTypedKernelLine(line) + " {";
if ( isFloatType(kernel_name_typed) ) {
type_replace = "float";
}
if ( isDoubleType(kernel_name_typed) ) {
type_replace = "double";
}
kernel_modified = kernel_line_typed + "\n" + kernel_buffer;
boost::regex re;
re = boost::regex("\\sT\\s", boost::regex::perl);
kernel_modified = boost::regex_replace(kernel_modified, re, " "+type_replace+" ");
re = boost::regex("\\sT\\*\\s", boost::regex::perl);
kernel_modified = boost::regex_replace(kernel_modified, re, " "+type_replace+"* ");
stdOpenCL += kernel_modified;
continue;
}
if ( isTemplateKernelLine(line) ) {
kernel_name = getKernelName(line);
kernel_type = getKernelType(line);
DLOG(INFO)<<"found template kernel '"<<kernel_name<<"' with type '"<<kernel_type<<"'";
if ( recording == false ) {
recording = true;
} else {
LOG(ERROR) << "error parsing kernel source file = '" << fileNameIN.c_str() << "'";
return false;
}
continue;
}
if ( recording ) {
kernel_buffer += line + "\n";
} else {
kernel_buffer = "";
stdOpenCL += line + "\n";
}
}
std::ofstream out(fileNameOUT.c_str());
out << stdOpenCL;
out.close();
DLOG(INFO) << "convert AMD OpenCL '"<<fileNameIN.c_str()<<"' to standard OpenCL '"<<fileNameOUT.c_str()<<"'";
return true;
}
