Archives Discussions

[twiig]

After calling clBuildProgram, my program crashes.  I have traced it down to trying to index a write-only global portion of memory with an index that has been involved in a division.  That make sense?

For example, in the attached code, the program will compile successfully if I remove the normalize call.  Furthermore, if I write my own normal function it will still fail.  I have tracked it down to the index variable(s) being involved in any way with a divide.  This only happens when it is compiled for the GPU by the way.

First, I think this must be a compiler bug.  Second, does anyone know how to normalize a vector without using a division?  I tried multiplying the vector by the length calculated using rsqrt to no avail.

__kernel void planar_hough(const int n, __global const float *points, const int NRANGE, __global int *out) { size_t gid = get_global_id(0); float4 v0 = (float4)(points[gid*3+0], points[gid*3+1], points[gid*3+2], 1.0); const int k = 3; int a[3]; for(int i = 0; i < k; i++) a = gid+i; while(true) { float4 v1 = (float4)(points[a[1]*3+0], points[a[1]*3+1], points[a[1]*3+2], 1.0); float4 v2 = (float4)(points[a[2]*3+0], points[a[2]*3+1], points[a[2]*3+2], 1.0); float4 v1_v0 = v1 - v0; float4 v2_v0 = v2 - v0; //distance constraint? float3 norm = normalize(cross(v1_v0, v2_v0)).xyz; out[(int)(norm.x*NRANGE*NRANGE) + (int)(norm.y*NRANGE)]++; ... ...

[omkaranathan]

Originally posted by: twiig  

For example, in the attached code, the program will compile successfully if I remove the normalize call.  

You are calling the normalize function on a float3 variable which is not supported as per OpenCL Specification.

Furthermore, if I write my own normal function it will still fail.  I have tracked it down to the index variable(s) being involved in any way with a divide.  This only happens when it is compiled for the GPU by the way.

Could you provide a test case for the above, as it will enable our developers to track down the problem easily.

[twiig]

I am actually not calling normalize on a float3 if you take another look.

My code for the kernel is posted below.  It compiles and runs correctly on the CPU, but crashes on clBuildProgram if compiled for the GPU.  If you need anything else to track it down, let me know.

*Note that the code is just for testing at the moment and does nothing terribly productive.

float dot(float3 v0, float3 v1) { return (v0.x*v1.x + v0.y*v1.y + v0.z+v1.z); } __kernel void planar_hough(const int n, __global const float *points, const int NRANGE, __global float *out) { size_t gid = get_global_id(0); float4 v0 = (float4)(points[gid*3+0], points[gid*3+1], points[gid*3+2], 1.0f); const int k = 3; int a[3]; for(int i = 0; i < k; i++) a = gid+i; while(true) { //compute index and vote float4 v1 = (float4)(points[a[1]*3+0], points[a[1]*3+1], points[a[1]*3+2], 1.0f); float4 v2 = (float4)(points[a[2]*3+0], points[a[2]*3+1], points[a[2]*3+2], 1.0f); float4 v1_v0 = v1 - v0; float4 v2_v0 = v2 - v0; //distance constraint? float3 norm = normalize(cross(v1_v0, v2_v0)).xyz; float d = dot(v0.xyz, norm); norm = norm*NRANGE * 0.5f; int index = ((int)norm.x)*NRANGE*NRANGE + ((int)norm.y)*NRANGE; out[index] = max(out[index], fabs(d)); //next combination please ... int j = k - 1; if(a < n - 1) { ++a; continue; } while(a - a[j-1] == 1) --j; if(j <= 1) break; int z = ++a[j-1]; while(j < k) { a = ++z; ++j; } } }

[nou]

is float3 even supported? i mean in OpenCL specification there is nothing aboutn float3 exept that floatX is reserved data type.

[twiig]

very good point.  It seems floatN where N = {2^x | x = 1,2,3,4} is supported.  I seem to be taking too much for granted from the shader languages.  It did however work on the CPU.

However, if I modify the program to use the substituted chunk of code listed below, it still fails.

float2 norm = normalize(cross(v1_v0, v2_v0)).xy; norm *= 0.5f; norm *= NRANGE; norm += NRANGE/2; int index = ((int)norm.x)*NRANGE + ((int)norm.y); out[index]++;

[omkaranathan]

The source code seems to be incomplete. It will be good if you can provide a small standalone test case. Or the complete source code(cpp, headers and kernel code).

Originally posted by: twiig  

However, if I modify the program to use the substituted chunk of code listed below, it still fails.

The code snippet (when tried as a standalone test case) is not giving any error at my end.

[twiig]

Ok, here is everything copied from multiple files (files are separated with comments below).  I also included a test file "ball.obj".  I hope that helps.

/*********cl_lib.h***************/ #ifndef _CL_LIB_H_ #define _CL_LIB_H_ #include <CL/cl.h> #ifdef X86 #pragma comment(lib, "../lib/OpenCL_x86.lib") #else #pragma comment(lib, "../lib/OpenCL_x64.lib") #endif using namespace std; #define null 0 cl_int status; bool CheckError(const cl_int error, const char *name, const bool exitOnFail) { if(error != CL_SUCCESS) { cerr << "CL Error: " << name << " (" << error << ")" << endl; if(exitOnFail) exit(EXIT_FAILURE); } return (error != CL_SUCCESS); } char* ReadKernelSource(const char *filename) { char source[65536]; ifstream infile(filename); if(!infile.is_open()) return null; int pos = 0; while(!infile.eof()) source[pos++] = infile.get(); source[pos-1] = 0; infile.close(); char *ret = new char[pos]; strcpy_s(ret, pos, source); return ret; } void SetupCLFirstDeviceOfType(cl_context &context, cl_device_id &deviceID, const cl_device_type type) { context = clCreateContextFromType(0, type, null, null, &status); CheckError(status, "clCreateContextFromType", true); status = clGetDeviceIDs(null, type, 1, &deviceID, null); CheckError(status, "clGetDeviceIDs", true); } cl_command_queue CreateCommandQueue(const cl_context context, const cl_device_id deviceID) { cl_command_queue cmdQueue = clCreateCommandQueue(context, deviceID, CL_QUEUE_PROFILING_ENABLE, &status); CheckError(status, "clCreateCommandQueue", true); return cmdQueue; } cl_program CreateProgramFromFile(const cl_context context, const int numDevices, const cl_device_id *deviceIDs, const char *filename) { const char *source = ReadKernelSource(filename); cl_program program = clCreateProgramWithSource(context, 1, &source, null, &status); CheckError(status, "clCreateProgramWithSource", true); status = clBuildProgram(program, numDevices, deviceIDs, null, null, null); bool fail = CheckError(status, "clBuildProgram", false); if(fail) { char *log = new char[65536]; status = clGetProgramBuildInfo(program, deviceIDs[0], CL_PROGRAM_BUILD_LOG, sizeof(char)*65536, log, null); //log[size-2] = 0;//FIX!! Should not be needed. cerr << endl << "Build Log:" << endl << log << endl; delete[] log; } delete[] source; if(fail) exit(EXIT_FAILURE); return program; } cl_kernel CreateKernel(const cl_program program, const char *name) { cl_kernel kernel = clCreateKernel(program, name, &status); CheckError(status, "clCreateKernel", true); return kernel; } cl_mem SetupKernelArg_global(const int pos, const cl_context context, const cl_kernel kernel, void* ptr, const size_t size, cl_mem_flags flags) { cl_mem buffer = clCreateBuffer(context, flags, size, ptr, &status); CheckError(status, "clCreateBuffer", true); status = clSetKernelArg(kernel, pos, sizeof(cl_mem), &buffer); CheckError(status, "clSetKernelArg", true); return buffer; } void SetupKernelArg(const int pos, const cl_kernel kernel, void* ptr, size_t size) { status = clSetKernelArg(kernel, pos, size, ptr); CheckError(status, "clSetKernelArg", true); } #endif /*************main.cpp***************/ #include <iostream> #include <utility> #include <fstream> #include <math.h> #include "../lib/cl_lib.h" #define ISDIGIT(x) ((int)(x) >= 48 && (int)(x) <= 57) #define null 0 void ReadPointCloud(const char *filename, float **points, int &numPoints); void ScalePointCloud(float *points, const int numPoints, const float sx, const float sy, const float sz); int main() { int NRANGE = 200; cl_context context; cl_device_id deviceID; float *points; int numPoints; ReadPointCloud("peace.obj", &points, numPoints); ScalePointCloud(points, numPoints, 1, 1, 1); float *votes = new float[NRANGE*NRANGE]; memset(votes, 0, sizeof(int)*NRANGE*NRANGE); SetupCLFirstDeviceOfType(context, deviceID, CL_DEVICE_TYPE_GPU); cl_command_queue cmdQueue = CreateCommandQueue(context, deviceID); cl_program program = CreateProgramFromFile(context, 1, &deviceID, "hough.cl"); cl_kernel voteKernel = CreateKernel(program, "planar_hough"); SetupKernelArg(0, voteKernel, &numPoints, sizeof(int)); cl_mem pts_mem = SetupKernelArg_global(1, context, voteKernel, points, sizeof(float)*numPoints*3, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR); SetupKernelArg(2, voteKernel, &NRANGE, sizeof(int)); cl_mem votes_mem = SetupKernelArg_global(3, context, voteKernel, votes, sizeof(int)*NRANGE*NRANGE, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR); size_t global_size[] = {numPoints}; size_t local_size[] = {1}; status = clEnqueueNDRangeKernel(cmdQueue, voteKernel, 1, null, global_size, local_size, 0, null, null); CheckError(status, "clEnqueueNDRangeKernel", true); clFinish(cmdQueue); status = clEnqueueReadBuffer(cmdQueue, votes_mem, CL_TRUE, 0, sizeof(int)*NRANGE*NRANGE, votes, 0, null, null); CheckError(status, "clEnqueueReadBuffer", true); std::ofstream outfile("votes.dat", std::ios_base::out); for(int i = 0; i < NRANGE; i++) { for(int j = 0; j < NRANGE; j++) outfile << votes[i*NRANGE + j] << " "; outfile << endl; } outfile.flush(); outfile.close(); clReleaseKernel(voteKernel); clReleaseProgram(program); clReleaseCommandQueue(cmdQueue); clReleaseContext(context); delete[] points; delete[] votes; } void ReadPointCloud(const char *filename, float **points, int &numPoints) { char line[1024]; std::ifstream *infile = new std::ifstream(filename, std::ios_base::in); if(!infile->is_open()) { cerr << "Error opening '" << filename << "'." << endl; exit(EXIT_FAILURE); } numPoints = 0; while(!infile->eof()) { infile->getline(line, 1024); if(line[0] == 'v') numPoints++; } infile->close(); delete infile; infile = new std::ifstream(filename, std::ios_base::in); *points = new float[numPoints*3]; numPoints = 0; while(!infile->eof()) { infile->getline(line, 1024); char *context, *token; token = strtok_s(line, " \r\n\t", &context); if(token != null && strcmp(token, "v") == 0) { token = strtok_s(null, " \r\n\t", &context); (*points)[numPoints++] = (float)atof(token); token = strtok_s(null, " \r\n\t", &context); (*points)[numPoints++] = (float)atof(token); token = strtok_s(null, " \r\n\t", &context); (*points)[numPoints++] = (float)atof(token); } } numPoints /= 3; infile->close(); } void ScalePointCloud(float *points, const int numPoints, const float sx, const float sy, const float sz) { double maxX, maxY, maxZ; maxX = maxY = maxZ = 0.000000000001; //avoid divide-by-zero for(int i = 0; i < numPoints; i++) { double x = fabs(points[i*3]+0); double y = fabs(points[i*3]+1); double z = fabs(points[i*3]+2); if(x > maxX) maxX = x; if(y > maxY) maxY = y; if(z > maxZ) maxZ = z; } double factorX = sx/maxX; double factorY = sy/maxY; double factorZ = sz/maxZ; for(int i = 0; i < numPoints; i++) { points[i*3+0] = (float)(factorX*points[i*3+0]); points[i*3+1] = (float)(factorY*points[i*3+1]); points[i*3+2] = (float)(factorZ*points[i*3+2]); } } /**********************hough.cl****************/ __kernel void planar_hough(const int n, __global const float *points, const int NRANGE, __global float *out) { const int NRANGE_2 = NRANGE/2; size_t gid = get_global_id(0); float4 v0 = (float4)(points[gid*3+0], points[gid*3+1], points[gid*3+2], 0.0f); const int k = 3; int a[3]; for(int i = 0; i < k; i++) a = gid+i; while(true) { //compute index and vote float4 v1 = (float4)(points[a[1]*3+0], points[a[1]*3+1], points[a[1]*3+2], 0.0f); float4 v2 = (float4)(points[a[2]*3+0], points[a[2]*3+1], points[a[2]*3+2], 0.0f); float4 v1_v0 = v1 - v0; float4 v2_v0 = v2 - v0; //distance constraint? float4 norm = (float4)(normalize(cross(v1_v0, v2_v0)).xyz, 0.0); float2 norm2 = norm.xy * (NRANGE_2) + (NRANGE_2); float d = dot(norm, v0) * (NRANGE_2) + (NRANGE_2); int index = ((int)norm2.x)*NRANGE + ((int)norm2.y);// + (int)d; out[index] = max(out[index], norm2.x*NRANGE*NRANGE); //next combination please ... int j = k - 1; if(a < n - 1) { ++a; continue; } while(a - a[j-1] == 1) --j; if(j <= 1) break; int z = ++a[j-1]; while(j < k) { a = ++z; ++j; } } } /****************ball.obj****************/ # Alias OBJ Model File # Exported from SketchUp, (c) 2000-2006 Google, Inc. # File units = inches mtllib ball.mtl g Mesh1 Default1 skp8DA8_1 Model usemtl Default1 v -0.185824 1.10438e-017 -0.202462 vt 0.0917398 0.0649468 vn 0.181871 -0.132137 0.974404 v -0.366188 3.74915e-033 -0.168797 vt -0.0917398 0.0649468 v -0.310453 -0.171537 -0.202462 vt -0.0307734 -0.108108 f 1/1/1 2/2/1 3/3/1 vt -0.0917398 -0.0649468 vn 0.181871 0.132137 0.974404 vt 0.0917398 -0.0649468 v -0.310453 0.171537 -0.202462 vt -0.0307734 0.108108 f 2/4/2 1/5/2 4/6/2 vt -0.103166 0.0586877 vn 0.324476 0.235746 0.916046 vt 0.0143112 -0.117824 v -0.120627 0.178411 -0.27147 vt 0.0988075 0.065761 f 4/7/3 1/8/3 5/9/3 vt 4.79192e-008 0.113406 vn 0.520872 0.135641 0.842789 vt -0.0917394 -0.0666693 v -0.0297472 2.06004e-017 -0.298922 vt 0.0917395 -0.0666693 f 5/10/4 1/11/4 6/12/4 vt -0.0917394 0.0666693 vn 0.520872 -0.135641 0.842789 v -0.120627 -0.178411 -0.27147 vt 4.79192e-008 -0.113406 vt 0.0917395 0.0666693 f 1/13/5 7/14/5 6/15/5 vt 0.0143112 0.117824 vn 0.324476 -0.235746 0.916046 vt -0.103166 -0.0586877 vt 0.0988075 -0.065761 f 1/16/6 3/17/6 7/18/6 vt 0.10294 0.0475855 vn 0.289961 -0.453463 0.842789 vt -0.0990103 0.0552979 v -0.262222 -0.319975 -0.298922 vt -0.022022 -0.111247 f 7/19/7 3/20/7 8/21/7 vt 0.0380265 0.106841 vn 0.0319559 -0.537294 0.842789 v -0.459985 -0.288675 -0.27147 vt -0.108783 -0.0320482 vt 0.089877 -0.0691596 f 3/22/8 9/23/8 8/24/8 vt 0.117757 0.0148568 vn -0.123939 -0.381444 0.916046 v -0.512107 -0.106016 -0.202462 vt -0.0820764 0.0857373 vt -0.0396773 -0.111862 f 3/25/9 10/26/9 9/27/9 vt 0.0887571 -0.0689671 vn -0.0694685 -0.213802 0.974404 vt 0.0355564 0.10663 vt -0.112386 -0.00189484 f 3/28/10 2/29/10 10/30/10 vt -0.0373489 -0.106016 vn -0.224805 0 0.974404 vt 0.112402 3.74915e-033 v -0.512107 0.106016 -0.202462 vt -0.0373489 0.106016 f 10/31/11 2/32/11 11/33/11 vt 0.0355564 -0.10663 vn -0.0694685 0.213802 0.974404 vt 0.0887571 0.0689671 vt -0.112386 0.00189484 f 2/34/12 4/35/12 11/36/12 vt -0.0820764 -0.0857373 vn -0.123939 0.381444 0.916046 vt 0.117757 -0.0148568 v -0.459985 0.288675 -0.27147 vt -0.0396773 0.111862 f 11/37/13 4/38/13 12/39/13 vt -0.108783 0.0320482 vn 0.0319559 0.537294 0.842789 vt 0.0380265 -0.106841 v -0.262222 0.319975 -0.298922 vt 0.089877 0.0691596 f 12/40/14 4/41/14 13/42/14 vt -0.0990103 -0.0552979 vn 0.289961 0.453463 0.842789 vt 0.10294 -0.0475855 vt -0.022022 0.111247 f 4/43/15 5/44/15 13/45/15 vt -0.106016 0.0533674 vn 0.432555 0.568378 0.699888 vt -8.97816e-009 -0.11869 v -0.0818576 0.319975 -0.410394 vt 0.106016 0.0533674 f 13/46/16 5/47/16 14/48/16 vt -0.102939 -0.0475855 vn 0.624139 0.453464 0.636256 v 0.0259886 0.171537 -0.410394 vt 0.0990104 -0.055298 vt 0.0220221 0.111247 f 5/49/17 15/50/17 14/51/17 vt -0.0988075 0.0657609 vn 0.674226 0.235746 0.699888 vt -0.0143112 -0.117825 vt 0.103166 0.0586876 f 5/52/18 6/53/18 15/54/18 vt 0.0307734 0.108108 vn 0.790198 0.132137 0.598437 vt -0.0917399 -0.0649468 v 0.0810251 2.7383e-017 -0.44519 vt 0.0917398 -0.0649468 f 15/55/19 6/56/19 16/57/19 vt -0.0917399 0.0649468 vn 0.790198 -0.132137 0.598437 v 0.0259886 -0.171537 -0.410394 vt 0.0307734 -0.108108 vt 0.0917398 0.0649468 f 6/58/20 17/59/20 16/60/20 vt -0.0143112 0.117825 vn 0.674226 -0.235746 0.699888 vt -0.0988075 -0.0657609 vt 0.103166 -0.0586876 f 6/61/21 7/62/21 17/63/21 vt 0.0990104 0.055298 vn 0.624139 -0.453464 0.636256 vt -0.102939 0.0475855 v -0.0818576 -0.319975 -0.410394 vt 0.0220221 -0.111247 f 17/64/22 7/65/22 18/66/22 vt -8.97816e-009 0.11869 vn 0.432555 -0.568378 0.699888 vt -0.106016 -0.0533674 vt 0.106016 -0.0533674 f 7/67/23 8/68/23 18/69/23 vt 0.106016 0.037349 vn 0.369854 -0.710691 0.598437 vt -0.106016 0.037349 v -0.227992 -0.425325 -0.44519 vt 1.21844e-009 -0.112402 f 18/70/24 8/71/24 19/72/24 vt 0.0301305 0.108289 vn 0.118515 -0.792356 0.598437 v -0.408141 -0.42599 -0.410394 vt -0.0913523 -0.0654907 vt 0.0921241 -0.0644005 f 8/73/25 20/74/25 19/75/25 vt 0.112319 0.0383651 vn -0.0158601 -0.714076 0.699888 vt -0.0847707 0.0830744 vt -0.0360872 -0.113071 f 8/76/26 9/77/26 20/78/26 vt 0.051381 -0.101098 vn -0.2384 -0.733719 0.636256 vt 0.0515773 0.100999 v -0.582639 -0.369292 -0.410394 vt -0.112024 -0.0176518 f 20/79/27 9/80/27 21/81/27 vt 0.118609 -0.0043894 vn -0.406892 -0.587022 0.699888 v -0.638375 -0.197755 -0.298922 vt -0.0494102 0.107917 vt -0.0572515 -0.10397 f 9/82/28 22/83/28 21/84/28 vt 0.0623513 -0.0947272 vn -0.341667 -0.415897 0.842789 vt 0.0399741 0.106127 vt -0.113284 0.00524942 f 9/85/29 10/86/29 22/87/29 vt 0.112912 -0.0105688 vn -0.501122 -0.196425 0.842789 v -0.669719 3.09805e-012 -0.27147 vt -0.0578296 0.0975534 vt -0.0449186 -0.104131 f 10/88/30 23/89/30 22/90/30 vt 0.0533673 -0.106016 vn -0.401074 -2.00652e-017 0.916046 vt 0.0533673 0.106016 vt -0.11869 3.09806e-012 f 10/91/31 11/92/31 23/93/31 vt -0.0578296 -0.0975534 vn -0.501122 0.196425 0.842789 vt 0.112912 0.0105688 v -0.638375 0.197755 -0.298922 vt -0.0449186 0.104131 f 23/94/32 11/95/32 24/96/32 vt 0.0399741 -0.106127 vn -0.341667 0.415897 0.842789 vt 0.0623513 0.0947272 vt -0.113284 -0.00524942 f 11/97/33 12/98/33 24/99/33 vt -0.0494102 -0.107917 vn -0.406892 0.587022 0.699888 vt 0.118609 0.0043894 v -0.582639 0.369292 -0.410394 vt -0.0572515 0.10397 f 24/100/34 12/101/34 25/102/34 vt 0.0515773 -0.100999 vn -0.2384 0.733719 0.636256 v -0.408141 0.42599 -0.410394 vt 0.051381 0.101098 vt -0.112024 0.0176518 f 12/103/35 26/104/35 25/105/35 vt -0.0847707 -0.0830744 vn -0.0158601 0.714076 0.699888 vt 0.112319 -0.0383651 vt -0.0360872 0.113071 f 12/106/36 13/107/36 26/108/36 vt -0.0913523 0.0654907 vn 0.118515 0.792356 0.598437 vt 0.0301305 -0.108289 v -0.227992 0.425325 -0.44519 vt 0.0921241 0.0644005 f 26/109/37 13/110/37 27/111/37 vt -0.106016 -0.037349 vn 0.369854 0.710691 0.598437 vt 0.106016 -0.037349 vt 1.21844e-009 0.112402 f 13/112/38 14/113/38 27/114/38 vt -0.0921241 0.0644006 vn 0.482257 0.792356 0.373632 vt -0.0301305 -0.108289 v -0.116304 0.42599 -0.590759 vt 0.0913523 0.0654907 f 27/115/39 14/116/39 28/117/39 vt -0.112319 -0.0383651 vn 0.633092 0.714076 0.298814 v 0.0311388 0.288675 -0.575001 vt 0.0847707 -0.0830744 vt 0.0360872 0.113071 f 14/118/40 29/119/40 28/120/40 vt -0.089877 0.0691595 vn 0.739522 0.537294 0.405489 vt -0.0380265 -0.106841 vt 0.108783 0.0320481 f 14/121/41 15/122/41 29/123/41 vt 0.0396773 0.111862 vn 0.874763 0.381444 0.298814 vt -0.117757 -0.0148569 v 0.116171 0.106016 -0.590759 vt 0.0820763 -0.0857374 f 29/124/42 15/125/42 30/126/42 vt -0.0887571 0.0689672 vn 0.902601 0.213802 0.373632 vt -0.0355564 -0.10663 vt 0.112386 0.00189491 f 15/127/43 16/128/43 30/129/43 vt -0.112402 3.13258e-017 vn 0.972069 -8.0926e-018 0.234695 v 0.116171 -0.106016 -0.590759 vt 0.0373489 -0.106016 vt 0.0373489 0.106016 f 16/130/44 31/131/44 30/132/44 vt -0.0355564 0.10663 vn 0.902601 -0.213802 0.373632 vt -0.0887571 -0.0689672 vt 0.112386 -0.00189491 f 16/133/45 17/134/45 31/135/45 vt -0.117757 0.0148569 vn 0.874763 -0.381444 0.298814 v 0.0311388 -0.288675 -0.575001 vt 0.0396773 -0.111862 vt 0.0820763 0.0857374 f 17/136/46 32/137/46 31/138/46 vt -0.0380265 0.106841 vn 0.739522 -0.537294 0.405489 vt -0.089877 -0.0691595 vt 0.108783 -0.0320481 f 17/139/47 18/140/47 32/141/47 vt 0.0847707 0.0830744 vn 0.633092 -0.714076 0.298814 vt -0.112319 0.0383651 v -0.116304 -0.42599 -0.590759 vt 0.0360872 -0.113071 f 32/142/48 18/143/48 33/144/48 vt -0.0301305 0.108289 vn 0.482257 -0.792356 0.373632 vt -0.0921241 -0.0644006 vt 0.0913523 -0.0654907 f 18/145/49 19/146/49 33/147/49 vt 0.0923533 0.0640715 vn 0.300386 -0.924493 0.234695 vt -0.0911181 0.0658162 v -0.317958 -0.491512 -0.590759 vt -0.0318 -0.10781 f 33/148/50 19/149/50 34/150/50 vt 0.0911181 0.0658162 vn 0.0755812 -0.924493 0.373632 vt -0.0923533 0.0640715 vt 0.0318 -0.10781 f 19/151/51 20/152/51 34/153/51 vt 0.0363044 0.113002 vn -0.0924585 -0.949822 0.298814 v -0.517954 -0.467086 -0.575001 vt -0.117258 -0.018382 vt 0.0691426 -0.0964714 f 20/154/52 35/155/52 34/156/52 vt 0.11328 -0.00533913 vn -0.282472 -0.869361 0.405489 vt -0.0299016 0.109393 vt -0.0709145 -0.0884989 f 20/157/53 21/158/53 35/159/53 vt 2.28081e-009 -0.11869 vn -0.483491 -0.822768 0.298814 vt 0.106016 0.0533674 v -0.694111 -0.369292 -0.590759 vt -0.106016 0.0533674 f 35/160/54 21/161/54 36/162/54 vt 0.106016 -0.037349 vn -0.60455 -0.703505 0.373632 v -0.727992 -0.262866 -0.44519 vt -1.15564e-008 0.112402 vt -0.106016 -0.037349 f 21/163/55 37/164/55 36/165/55 vt 0.0189973 -0.110785 vn -0.561616 -0.571368 0.598437 vt 0.0546377 0.0982294 vt -0.110803 0.0188937 f 21/166/56 22/167/56 37/168/56 vt 0.109538 -0.0252155 vn -0.716952 -0.357566 0.598437 v -0.784293 -0.0917395 -0.410394 vt -0.0695445 0.0883053 vt -0.0601798 -0.0949351 f 22/169/57 38/170/57 37/171/57 vt 0.0638686 -0.100041 vn -0.684028 -0.205578 0.699888 vt 0.0606405 0.10203 vt -0.118401 0.00828846 f 22/172/58 23/173/58 38/174/58 vt -0.0666693 -0.0917395 vn -0.771478 1.82267e-017 0.636256 vt 0.113406 3.09804e-012 v -0.784293 0.0917395 -0.410394 vt -0.0666693 0.0917395 f 38/175/59 23/176/59 39/177/59 vt 0.0606405 -0.10203 vn -0.684028 0.205578 0.699888 vt 0.0638686 0.100041 vt -0.118401 -0.00828846 f 23/178/60 24/179/60 39/180/60 vt -0.0695445 -0.0883053 vn -0.716952 0.357566 0.598437 vt 0.109538 0.0252155 v -0.727992 0.262866 -0.44519 vt -0.0601798 0.0949351 f 39/181/61 24/182/61 40/183/61 vt 0.0546377 -0.0982294 vn -0.561616 0.571368 0.598437 vt 0.0189973 0.110785 vt -0.110803 -0.0188937 f 24/184/62 25/185/62 40/186/62 vt -1.15564e-008 -0.112402 vn -0.60455 0.703505 0.373632 vt 0.106016 0.037349 v -0.694111 0.369292 -0.590759 vt -0.106016 0.037349 f 40/187/63 25/188/63 41/189/63 vt 0.106016 -0.0533674 vn -0.483491 0.822768 0.298814 v -0.517954 0.467086 -0.575001 vt 2.28081e-009 0.11869 vt -0.106016 -0.0533674 f 25/190/64 42/191/64 41/192/64 vt -0.0299016 -0.109393 vn -0.282472 0.869361 0.405489 vt 0.11328 0.00533913 vt -0.0709145 0.0884989 f 25/193/65 26/194/65 42/195/65 vt -0.117258 0.018382 vn -0.0924585 0.949822 0.298814 vt 0.0363044 -0.113002 v -0.317958 0.491512 -0.590759 vt 0.0691426 0.0964714 f 42/196/66 26/197/66 43/198/66 vt -0.0923533 -0.0640715 vn 0.0755812 0.924493 0.373632 vt 0.0911181 -0.0658162 vt 0.0318 0.10781 f 26/199/67 27/200/67 43/201/67 vt -0.0911181 -0.0658162 vn 0.300386 0.924493 0.234695 vt 0.0923533 -0.0640715 vt -0.0318 0.10781 f 27/202/68 28/203/68 43/204/68 vt -0.0691426 0.0964714 vn 0.308616 0.949822 0.0509362 vt -0.0363044 -0.113002 v -0.214423 0.467086 -0.762593 vt 0.117258 0.0183821 f 43/205/69 28/206/69 44/207/69 vt -0.11328 0.00533913 vn 0.489006 0.869361 -0.0713108 v -0.038266 0.369292 -0.746835 vt 0.0299016 -0.109393 vt 0.0709145 0.0884989 f 28/208/70 45/209/70 44/210/70 vt -0.051381 0.101098 vn 0.6757 0.733719 0.0713109 vt -0.0515773 -0.100999 vt 0.112024 0.0176518 f 28/211/71 29/212/71 45/213/71 vt -0.118609 0.00438942 vn 0.807967 0.587022 -0.0509362 v 0.0863631 0.197755 -0.746835 vt 0.0494102 -0.107917 vt 0.0572515 0.10397 f 29/214/72 46/215/72 45/216/72 vt -0.0623513 0.0947272 vn 0.906611 0.415897 0.0713108 vt -0.0399741 -0.106127 vt 0.113284 -0.00524943 f 29/217/73 30/218/73 46/219/73 vt 0.0449186 0.104131 vn 0.977922 0.196425 -0.0713109 vt -0.112912 0.0105688 v 0.124935 -3.26274e-012 -0.762593 vt 0.0578296 -0.0975534 f 46/220/74 30/221/74 47/222/74 vt -0.0533674 0.106016 vn 0.998702 -4.34726e-019 0.0509362 vt -0.0533674 -0.106016 vt 0.11869 -3.26274e-012 f 30/223/75 31/224/75 47/225/75 vt -0.112912 -0.0105688 vn 0.977922 -0.196425 -0.0713109 v 0.0863631 -0.197755 -0.746835 vt 0.0449186 -0.104131 vt 0.0578296 0.0975534 f 31/226/76 48/227/76 47/228/76 vt -0.0399741 0.106127 vn 0.906611 -0.415897 0.0713108 vt -0.0623513 -0.0947272 vt 0.113284 0.00524943 f 31/229/77 32/230/77 48/231/77 vt 0.0494102 0.107917 vn 0.807967 -0.587022 -0.0509362 vt -0.118609 -0.00438942 v -0.038266 -0.369292 -0.746835 vt 0.0572515 -0.10397 f 48/232/78 32/233/78 49/234/78 vt -0.0515773 0.100999 vn 0.6757 -0.733719 0.0713109 vt -0.051381 -0.101098 vt 0.112024 -0.0176518 f 32/235/79 33/236/79 49/237/79 vt 0.0299016 0.109393 vn 0.489006 -0.869361 -0.0713108 vt -0.11328 -0.00533913 v -0.214423 -0.467086 -0.762593 vt 0.0709145 -0.0884989 f 49/238/80 33/239/80 50/240/80 vt -0.0363044 0.113002 vn 0.308616 -0.949822 0.0509362 vt -0.0691426 -0.0964714 vt 0.117258 -0.0183821 f 33/241/81 34/242/81 50/243/81 vt -0.0917395 -0.0666693 vn 0.115383 -0.990758 -0.0713109 v -0.414419 -0.491512 -0.746835 vt 0.0917395 -0.0666693 vt 4.07047e-009 0.113406 f 34/244/82 51/245/82 50/246/82 vt 0.0917395 -0.0666693 vn -0.115383 -0.990758 0.0713109 vt 4.07047e-009 0.113406 vt -0.0917395 -0.0666693 f 34/247/83 35/248/83 51/249/83 vt -0.0691426 -0.0964714 vn -0.308616 -0.949822 -0.0509362 vt 0.117258 -0.0183821 v -0.616073 -0.42599 -0.746835 f 51/250/84 35/251/84 52/241/84 vn -0.489006 -0.869361 0.0713108 f 35/240/85 36/238/85 52/239/85 vt -0.051381 -0.101098 vn -0.6757 -0.733719 -0.0713109 vt 0.112024 -0.0176518 v -0.763516 -0.288675 -0.762593 vt -0.0515773 0.100999 f 52/252/86 36/253/86 53/254/86 vt 0.0572515 -0.10397 vn -0.807967 -0.587022 0.0509362 v -0.81874 -0.197755 -0.590759 vt 0.0494102 0.107917 f 36/255/87 54/256/87 53/233/87 vt -0.0189973 -0.110785 vn -0.78642 -0.571368 0.234695 vt 0.110803 0.0188937 vt -0.0546377 0.0982294 f 36/257/88 37/258/88 54/259/88 vt 0.0601798 -0.0949352 vn -0.855889 -0.357566 0.373632 vt 0.0695446 0.0883053 vt -0.109538 -0.0252155 f 37/260/89 38/261/89 54/262/89 vt 0.118401 0.00828847 vn -0.931906 -0.205578 0.298814 v -0.857312 5.01275e-012 -0.575001 vt -0.0606405 0.10203 vt -0.0638686 -0.100041 f 38/263/90 55/264/90 54/265/90 vt 0.0666693 -0.0917395 vn -0.9141 1.86786e-018 0.405489 vt 0.0666693 0.0917395 vt -0.113406 5.01275e-012 f 38/266/91 39/267/91 55/268/91 vt -0.0606405 -0.10203 vn -0.931906 0.205578 0.298814 vt 0.118401 -0.00828847 v -0.81874 0.197755 -0.590759 vt -0.0638686 0.100041 f 55/269/92 39/270/92 56/271/92 vt 0.0695446 -0.0883053 vn -0.855889 0.357566 0.373632 vt 0.0601798 0.0949352 vt -0.109538 0.0252155 f 39/272/93 40/273/93 56/274/93 vt 0.110803 -0.0188937 vn -0.78642 0.571368 0.234695 vt -0.0189973 0.110785 vt -0.0546377 -0.0982294 f 40/275/94 41/276/94 56/277/94 vt 0.0494102 -0.107917 vn -0.807967 0.587022 0.0509362 vt 0.0572515 0.10397 v -0.763516 0.288675 -0.762593 f 56/278/95 41/279/95 57/214/95 vt 0.112024 0.0176518 vn -0.6757 0.733719 -0.0713109 v -0.616073 0.42599 -0.746835 vt -0.051381 0.101098 vt -0.0515773 -0.100999 f 41/280/96 58/281/96 57/282/96 vn -0.489006 0.869361 0.0713108 f 41/209/97 42/210/97 58/208/97 vt 0.117258 0.0183821 vn -0.308616 0.949822 -0.0509362 v -0.414419 0.491512 -0.746835 vt -0.0691426 0.0964714 f 42/283/98 59/284/98 58/206/98 vt 4.07047e-009 -0.113406 vn -0.115383 0.990758 0.0713109 vt 0.0917395 0.0666693 vt -0.0917395 0.0666693 f 42/285/99 43/286/99 59/287/99 vt 0.0917395 0.0666693 vn 0.115383 0.990758 -0.0713109 vt -0.0917395 0.0666693 vt 4.07047e-009 -0.113406 f 59/288/100 43/289/100 44/290/100 vt 0.0691426 0.0964714 vn 0.0924585 0.949822 -0.298814 vt -0.117258 0.018382 v -0.324236 0.42599 -0.9272 vt 0.0363044 -0.113002 f 59/291/101 44/292/101 60/293/101 vn 0.282472 0.869361 -0.405489 v -0.149738 0.369292 -0.9272 f 44/195/102 61/193/102 60/194/102 vt 2.28081e-009 0.11869 vn 0.483491 0.822768 -0.298814 vt -0.106016 -0.0533674 vt 0.106016 -0.0533674 f 44/294/103 45/295/103 61/296/103 vn 0.60455 0.703505 -0.373632 v -0.00438506 0.262866 -0.892404 f 45/189/104 62/187/104 61/188/104 vn 0.78642 0.571368 -0.234695 f 45/276/105 46/277/105 62/275/105 vt 0.0601798 0.0949352 vn 0.855889 0.357566 -0.373632 vt -0.109538 0.0252155 v 0.0519164 0.0917395 -0.9272 vt 0.0695446 -0.0883053 f 62/297/106 46/298/106 63/299/106 vt -0.0638686 0.100041 vn 0.931906 0.205578 -0.298814 vt -0.0606405 -0.10203 vt 0.118401 -0.00828847 f 46/300/107 47/301/107 63/302/107 vt -0.113406 -3.26275e-012 vn 0.9141 3.56536e-018 -0.405489 v 0.0519164 -0.0917395 -0.9272 f 47/303/108 64/266/108 63/267/108 vt -0.0606405 0.10203 vn 0.931906 -0.205578 -0.298814 vt -0.0638686 -0.100041 vt 0.118401 0.00828847 f 47/304/109 48/305/109 64/306/109 vt -0.109538 -0.0252155 vn 0.855889 -0.357566 -0.373632 v -0.00438506 -0.262866 -0.892404 vt 0.0601798 -0.0949352 vt 0.0695446 0.0883053 f 48/307/110 65/308/110 64/309/110 vn 0.78642 -0.571368 -0.234695 f 48/259/111 49/257/111 65/258/111 vn 0.60455 -0.703505 -0.373632 v -0.149738 -0.369292 -0.9272 f 65/164/112 49/165/112 66/163/112 vt -0.106016 0.0533674 vn 0.483491 -0.822768 -0.298814 vt 2.28081e-009 -0.11869 vt 0.106016 0.0533674 f 49/310/113 50/311/113 66/312/113 vn 0.282472 -0.869361 -0.405489 v -0.324236 -0.42599 -0.9272 f 66/158/114 50/159/114 67/157/114 vt -0.117258 -0.018382 vn 0.0924585 -0.949822 -0.298814 vt 0.0691426 -0.0964714 vt 0.0363044 0.113002 f 50/313/115 51/314/115 67/315/115 vn -0.0755812 -0.924493 -0.373632 v -0.504385 -0.425325 -0.892404 f 51/153/116 68/151/116 67/152/116 vn -0.300386 -0.924493 -0.234695 f 51/150/117 52/148/117 68/149/117 vt -0.0921241 -0.0644006 vn -0.482257 -0.792356 -0.373632 vt 0.0913523 -0.0654907 v -0.650519 -0.319975 -0.9272 vt -0.0301305 0.108289 f 68/316/118 52/317/118 69/318/118 vt 0.0360872 -0.113071 vn -0.633092 -0.714076 -0.298814 vt 0.0847707 0.0830744 vt -0.112319 0.0383651 f 52/319/119 53/320/119 69/321/119 vn -0.739522 -0.537294 -0.405489 v -0.758366 -0.171537 -0.9272 f 69/140/120 53/141/120 70/139/120 vt 0.0396773 -0.111862 vn -0.874763 -0.381444 -0.298814 v -0.848548 -0.106016 -0.746835 vt 0.0820763 0.0857374 vt -0.117757 0.0148569 f 53/322/121 71/323/121 70/324/121 vn -0.906611 -0.415897 -0.0713108 f 54/231/122 71/229/122 53/230/122 vt 0.0449186 -0.104131 vn -0.977922 -0.196425 0.0713109 vt 0.0578296 0.0975534 vt -0.112912 -0.0105688 f 54/325/123 55/326/123 71/327/123 vt 0.11869 5.01275e-012 vn -0.998702 4.69867e-019 -0.0509362 v -0.848548 0.106016 -0.746835 vt -0.0533674 0.106016 vt -0.0533674 -0.106016 f 55/328/124 72/329/124 71/330/124 vt 0.0578296 -0.0975534 vn -0.977922 0.196425 0.0713109 vt 0.0449186 0.104131 vt -0.112912 0.0105688 f 55/331/125 56/332/125 72/333/125 vn -0.906611 0.415897 -0.0713108 f 72/218/126 56/219/126 57/217/126 vt 0.0820763 -0.0857374 vn -0.874763 0.381444 -0.298814 vt 0.0396773 0.111862 v -0.758366 0.171537 -0.9272 vt -0.117757 -0.0148569 f 72/334/127 57/335/127 73/336/127 vn -0.739522 0.537294 -0.405489 v -0.650519 0.319975 -0.9272 f 57/123/128 74/121/128 73/122/128 vt 0.0847707 -0.0830744 vn -0.633092 0.714076 -0.298814 vt 0.0360872 0.113071 vt -0.112319 -0.0383651 f 57/337/129 58/338/129 74/339/129 vt 0.0913523 0.0654907 vn -0.482257 0.792356 -0.373632 v -0.504385 0.425325 -0.892404 vt -0.0921241 0.0644006 vt -0.0301305 -0.108289 f 58/340/130 75/341/130 74/342/130 vn -0.300386 0.924493 -0.234695 f 58/203/131 59/204/131 75/202/131 vn -0.0755812 0.924493 -0.373632 f 75/200/132 59/201/132 60/199/132 vt 0.0921241 0.0644005 vn -0.118515 0.792356 -0.598437 vt -0.0913523 0.0654907 v -0.470155 0.319975 -1.03867 vt 0.0301305 -0.108289 f 75/343/133 60/344/133 76/345/133 vn 0.0158601 0.714076 -0.699888 v -0.272392 0.288675 -1.06612 f 76/107/134 60/108/134 77/106/134 vn 0.2384 0.733719 -0.636256 f 60/104/135 61/105/135 77/103/135 vn 0.406892 0.587022 -0.699888 v -0.0940017 0.197755 -1.03867 f 61/102/136 78/100/136 77/101/136 vt 0.0189973 0.110785 vn 0.561616 0.571368 -0.598437 vt 0.0546377 -0.0982294 f 61/346/137 62/186/137 78/347/137 vt -0.0601798 0.0949351 vn 0.716952 0.357566 -0.598437 vt -0.0695445 -0.0883053 vt 0.109538 0.0252155 f 62/348/138 63/349/138 78/350/138 vt 0.0638686 0.100041 vn 0.684028 0.205578 -0.699888 vt -0.118401 -0.00828846 v -0.0626575 -3.09814e-012 -1.06612 vt 0.0606405 -0.10203 f 78/351/139 63/352/139 79/353/139 vt -0.0666693 0.0917395 vn 0.771478 5.80655e-018 -0.636256 vt -0.0666693 -0.0917395 vt 0.113406 -3.09814e-012 f 63/354/140 64/355/140 79/356/140 vt -0.118401 0.00828846 vn 0.684028 -0.205578 -0.699888 v -0.0940017 -0.197755 -1.03867 vt 0.0638686 -0.100041 vt 0.0606405 0.10203 f 64/357/141 80/358/141 79/359/141 vt -0.0695445 0.0883053 vn 0.716952 -0.357566 -0.598437 vt -0.0601798 -0.0949351 vt 0.109538 -0.0252155 f 64/360/142 65/361/142 80/362/142 vn 0.561616 -0.571368 -0.598437 vt 0.0189973 -0.110785 vt 0.0546377 0.0982294 f 65/168/143 66/363/143 80/364/143 vn 0.406892 -0.587022 -0.699888 v -0.272392 -0.288675 -1.06612 f 80/83/144 66/84/144 81/82/144 vn 0.2384 -0.733719 -0.636256 f 66/81/145 67/79/145 81/80/145 vn 0.0158601 -0.714076 -0.699888 v -0.470155 -0.319975 -1.03867 f 67/78/146 82/76/146 81/77/146 vt -0.0913523 -0.0654907 vn -0.118515 -0.792356 -0.598437 vt 0.0921241 -0.0644005 vt 0.0301305 0.108289 f 67/365/147 68/366/147 82/367/147 vn -0.369854 -0.710691 -0.598437 f 68/72/148 69/70/148 82/71/148 vt -0.106016 -0.0533674 vn -0.432555 -0.568378 -0.699888 vt 0.106016 -0.0533674 v -0.61175 -0.178411 -1.06612 vt -8.97816e-009 0.11869 f 82/368/149 69/369/149 83/370/149 vn -0.624139 -0.453464 -0.636256 f 69/66/150 70/64/150 83/65/150 vt 0.103166 -0.0586876 vn -0.674226 -0.235746 -0.699888 v -0.70263 3.43398e-012 -1.03867 vt -0.0143112 0.117825 vt -0.0988075 -0.0657609 f 70/371/151 84/372/151 83/373/151 vt 0.0307734 -0.108108 vn -0.790198 -0.132137 -0.598437 v -0.813402 4.5646e-012 -0.892404 vt 0.0917398 0.0649468 vt -0.0917399 0.0649468 f 70/374/152 85/375/152 84/376/152 vt 0.112386 -0.00189491 vn -0.902601 -0.213802 -0.373632 vt -0.0355564 0.10663 vt -0.0887571 -0.0689672 f 71/377/153 85/378/153 70/379/153 vn -0.972069 6.70917e-018 -0.234695 vt -0.112402 4.5646e-012 f 71/131/154 72/132/154 85/380/154 vt -0.0355564 -0.10663 vn -0.902601 0.213802 -0.373632 vt 0.112386 0.00189491 vt -0.0887571 0.0689672 f 85/381/155 72/382/155 73/383/155 vt 0.0917398 -0.0649468 vn -0.790198 0.132137 -0.598437 vt 0.0307734 0.108108 vt -0.0917399 -0.0649468 f 85/384/156 73/385/156 84/386/156 vt -0.0143112 -0.117825 vn -0.674226 0.235746 -0.699888 vt 0.103166 0.0586876 v -0.61175 0.178411 -1.06612 vt -0.0988075 0.0657609 f 84/387/157 73/388/157 86/389/157 vn -0.624139 0.453464 -0.636256 f 73/50/158 74/51/158 86/49/158 vt 0.106016 0.0533674 vn -0.432555 0.568378 -0.699888 vt -0.106016 0.0533674 vt -8.97816e-009 -0.11869 f 74/390/159 76/391/159 86/392/159 vn -0.369854 0.710691 -0.598437 f 74/113/160 75/114/160 76/112/160 vt 0.10294 -0.0475855 vn -0.289961 0.453463 -0.842789 vt -0.022022 0.111247 v -0.421924 0.171537 -1.13513 vt -0.0990103 -0.0552979 f 86/393/161 76/394/161 87/395/161 vn -0.0319559 0.537294 -0.842789 f 76/42/162 77/40/162 87/41/162 vn 0.123939 0.381444 -0.916046 v -0.22027 0.106016 -1.13513 f 77/39/163 88/37/163 87/38/163 vt 0.0623513 0.0947272 vn 0.341667 0.415897 -0.842789 vt 0.0399741 -0.106127 f 77/396/164 78/99/164 88/397/164 vt -0.0449186 0.104131 vn 0.501122 0.196425 -0.842789 vt -0.0578296 -0.0975534 vt 0.112912 0.0105688 f 78/398/165 79/399/165 88/400/165 vn 0.401074 -2.58794e-017 -0.916046 vt -0.11869 -3.09813e-012 v -0.22027 -0.106016 -1.13513 f 88/92/166 79/401/166 89/91/166 vt -0.0578296 0.0975534 vn 0.501122 -0.196425 -0.842789 vt -0.0449186 -0.104131 vt 0.112912 -0.0105688 f 79/402/167 80/403/167 89/404/167 vn 0.341667 -0.415897 -0.842789 vt 0.0623513 -0.0947272 vt 0.0399741 0.106127 f 80/87/168 81/405/168 89/406/168 vn 0.123939 -0.381444 -0.916046 v -0.421924 -0.171537 -1.13513 f 89/26/169 81/27/169 90/25/169 vn -0.0319559 -0.537294 -0.842789 f 81/23/170 82/24/170 90/22/170 vt -0.022022 -0.111247 vn -0.289961 -0.453463 -0.842789 vt 0.10294 0.0475855 vt -0.0990103 0.0552979 f 82/407/171 83/408/171 90/409/171 vt -0.103166 -0.0586877 vn -0.324476 -0.235746 -0.916046 vt 0.0988075 -0.065761 v -0.546553 1.84094e-012 -1.13513 vt 0.0143112 0.117824 f 90/410/172 83/411/172 91/412/172 vt 4.79186e-008 -0.113406 vn -0.520872 -0.135641 -0.842789 vt 0.0917395 0.0666693 vt -0.0917394 0.0666693 f 83/413/173 84/414/173 91/415/173 vt 0.0917395 -0.0666693 vn -0.520872 0.135641 -0.842789 vt 4.79197e-008 0.113406 vt -0.0917394 -0.0666693 f 84/416/174 86/417/174 91/418/174 vt 0.0988075 0.065761 vn -0.324476 0.235746 -0.916046 vt -0.103166 0.0586877 vt 0.0143112 -0.117824 f 86/419/175 87/420/175 91/421/175 vt 0.0917398 -0.0649468 vn -0.181871 0.132137 -0.974404 vt -0.0307734 0.108108 v -0.366188 -3.74915e-033 -1.1688 vt -0.0917398 -0.0649468 f 91/422/176 87/423/176 92/424/176 vt 0.0887571 0.0689671 vn 0.0694685 0.213802 -0.974404 vt -0.112386 0.00189484 f 87/425/177 88/426/177 92/34/177 vt -0.0373489 0.106016 vn 0.224805 -5.1013e-049 -0.974404 vt -0.0373489 -0.106016 vt 0.112402 -3.74915e-033 f 88/427/178 89/428/178 92/429/178 vt -0.112386 -0.00189484 vn 0.0694685 -0.213802 -0.974404 vt 0.0887571 -0.0689671 f 89/430/179 90/431/179 92/29/179 vt -0.0307734 -0.108108 vn -0.181871 -0.132137 -0.974404 vt 0.0917398 0.0649468 vt -0.0917398 0.0649468 f 90/432/180 91/433/180 92/434/180

[omkaranathan]

Originally posted by: twiig I am actually not calling normalize on a float3 if you take another look.

My bad, I noticed it later.

My code for the kernel is posted below.  It compiles and runs correctly on the CPU, but crashes on clBuildProgram if compiled for the GPU.  If you need anything else to track it down, let me know. 

It will be helpful to have the host side code too.

It seems floatN where N = {2^x | x = 1,2,3,4} is supported.

As per OpenCL Specification, in floatn, supported values of n are 2, 4, 8, and 16.

[twiig]

Ok, here is the host code.  It is a bit messy as I'm working out the best way to break it into a few function calls.

void SetupCLFirstDeviceOfType(cl_context &context, cl_device_id &deviceID, const cl_device_type type) { context = clCreateContextFromType(0, type, null, null, &status); CheckError(status, "clCreateContextFromType", true); status = clGetDeviceIDs(null, type, 1, &deviceID, null); CheckError(status, "clGetDeviceIDs", true); } cl_command_queue CreateCommandQueue(const cl_context context, const cl_device_id deviceID) { cl_command_queue cmdQueue = clCreateCommandQueue(context, deviceID, CL_QUEUE_PROFILING_ENABLE, &status); CheckError(status, "clCreateCommandQueue", true); return cmdQueue; } cl_program CreateProgramFromFile(const cl_context context, const int numDevices, const cl_device_id *deviceIDs, const char *filename) { const char *source = ReadKernelSource(filename); cl_program program = clCreateProgramWithSource(context, 1, &source, null, &status); CheckError(status, "clCreateProgramWithSource", true); status = clBuildProgram(program, numDevices, deviceIDs, null, null, null); bool fail = CheckError(status, "clBuildProgram", false); if(fail) { char *log = new char[65536]; status = clGetProgramBuildInfo(program, deviceIDs[0], CL_PROGRAM_BUILD_LOG, sizeof(char)*65536, log, null); cerr << endl << "Build Log:" << endl << log << endl; delete[] log; } delete[] source; if(fail) exit(EXIT_FAILURE); return program; } cl_kernel CreateKernel(const cl_program program, const char *name) { cl_kernel kernel = clCreateKernel(program, name, &status); CheckError(status, "clCreateKernel", true); return kernel; } cl_mem SetupKernelArg_global(const int pos, const cl_context context, const cl_kernel kernel, void* ptr, const size_t size, cl_mem_flags flags) { cl_mem buffer = clCreateBuffer(context, flags, size, ptr, &status); CheckError(status, "clCreateBuffer", true); status = clSetKernelArg(kernel, pos, sizeof(cl_mem), &buffer); CheckError(status, "clSetKernelArg", true); return buffer; } void SetupKernelArg(const int pos, const cl_kernel kernel, void* ptr, size_t size) { status = clSetKernelArg(kernel, pos, size, ptr); CheckError(status, "clSetKernelArg", true); } int main() { int NRANGE = 200; cl_context context; cl_device_id deviceID; float *points; int numPoints; ReadPointCloud("peace.obj", &points, numPoints); ScalePointCloud(points, numPoints, 1, 1, 1); float *votes = new float[NRANGE*NRANGE]; memset(votes, 0, sizeof(int)*NRANGE*NRANGE); SetupCLFirstDeviceOfType(context, deviceID, CL_DEVICE_TYPE_GPU); cl_command_queue cmdQueue = CreateCommandQueue(context, deviceID); cl_program program = CreateProgramFromFile(context, 1, &deviceID, "hough.cl"); cl_kernel voteKernel = CreateKernel(program, "planar_hough"); SetupKernelArg(0, voteKernel, &numPoints, sizeof(int)); cl_mem pts_mem = SetupKernelArg_global(1, context, voteKernel, points, sizeof(float)*numPoints*3, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR); SetupKernelArg(2, voteKernel, &NRANGE, sizeof(int)); cl_mem votes_mem = SetupKernelArg_global(3, context, voteKernel, votes, sizeof(int)*NRANGE*NRANGE, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR); size_t global_size[] = {numPoints}; size_t local_size[] = {1}; status = clEnqueueNDRangeKernel(cmdQueue, voteKernel, 1, null, global_size, local_size, 0, null, null); CheckError(status, "clEnqueueNDRangeKernel", true); status = clEnqueueReadBuffer(cmdQueue, votes_mem, CL_TRUE, 0, sizeof(int)*NRANGE*NRANGE, votes, 0, null, null); CheckError(status, "clEnqueueReadBuffer", true); clReleaseKernel(voteKernel); clReleaseProgram(program); clReleaseCommandQueue(cmdQueue); clReleaseContext(context); delete[] points; delete[] votes; }