Optimizing kernel shuffled keys code - OpenCL -


i have started getting opencl , going through basics of writing kernel code. have written kernel code calculating shuffled keys points array. so, number of points n, shuffled keys calculated in 3-bit fashion, x-bit @ depth d (0 

xd = 0 if p.x < cd.x   xd = 1, otherwise

the shuffled xyz key given as:

x1y1z1x2y2z2...xdydzd

the kernel code written given below. point inputted in column major format. 

__constant float3 boundsoffsettable[8] = {               {-0.5,-0.5,-0.5},               {+0.5,-0.5,-0.5},               {-0.5,+0.5,-0.5},               {-0.5,-0.5,+0.5},               {+0.5,+0.5,-0.5},               {+0.5,-0.5,+0.5},               {-0.5,+0.5,+0.5},               {+0.5,+0.5,+0.5} }; uint setbit(uint x,unsigned char position) { uint mask = 1<<position; return x|mask; }  __kernel void morton_code(__global float* point,__global uint*code,int level, float3          center,float radius,int size){ // index of current element processed int = get_global_id(0); float3 pt;  pt.x = point[i];pt.y = point[size+i]; pt.z = point[2*size+i]; code[i] = 0; float3 newcenter; float newradius; if(pt.x>center.x) code = setbit(code,0); if(pt.y>center.y) code = setbit(code,1); if(pt.z>center.z) code = setbit(code,2); for(int l = 1;l<level;l++) {     for(int i=0;i<8;i++)     {         newradius = radius *0.5;         newcenter = center + boundoffsettable[i]*radius;         if(newcenter.x-newradius<pt.x && newcenter.x+newradius>pt.x && newcenter.y-newradius<pt.y && newcenter.y+newradius>pt.y && newcenter.z-newradius<pt.z && newcenter.z+newradius>pt.z)         {             if(pt.x>newcenter.x) code = setbit(code,3*l);             if(pt.y>newcenter.y) code = setbit(code,3*l+1);             if(pt.z>newcenter.z) code = setbit(code,3*l+2);         }     } } }

it works wanted ask if missing in code , if there way optimize code.

try kernel:

__kernel void morton_code(__global float* point,__global uint*code,int level, float3          center,float radius,int size){ // index of current element processed int = get_global_id(0); float3 pt;  pt.x = point[i];pt.y = point[size+i]; pt.z = point[2*size+i]; uint res; res = 0; float3 newcenter; float newradius; if(pt.x>center.x) res = setbit(res,0); if(pt.y>center.y) res = setbit(res,1); if(pt.z>center.z) res = setbit(res,2); for(int l = 1;l<level;l++) {     for(int i=0;i<8;i++)     {         newradius = radius *0.5;         newcenter = center + boundoffsettable[i]*radius;         if(newcenter.x-newradius<pt.x && newcenter.x+newradius>pt.x && newcenter.y-newradius<pt.y && newcenter.y+newradius>pt.y && newcenter.z-newradius<pt.z && newcenter.z+newradius>pt.z)         {             if(pt.x>newcenter.x) res = setbit(res,3*l);             if(pt.y>newcenter.y) res = setbit(res,3*l+1);             if(pt.z>newcenter.z) res = setbit(res,3*l+2);         }     } } //save result code[i] = res; }

rules optimize:

  1. avoid global memory (you using "code" directly global memory, changed that), should see 3x increase in performance now.
  2. avoid ifs, use "select" instead if possible. (see opencl documentation)
  3. use more memory inside kernel. don't need operate @ bit level. operation @ int level better , avoid huge amount of calls "setbit". can construct result @ end.

another interesting thing. if operating @ 3d level, can use float3 variables , compute distances opencl operators. can increase performance quite lot. requires complete rewrite of kernel.


Comments

Post a Comment

Popular posts from this blog

jquery - How can I dynamically add a browser tab? -

since using iframes house external pages are, accounts "not preferred method" , disallowed sites (flickr, one), want dynamically generate browser tabs. e.g., when user submits form, want dynamically generate tab on browser in response submitted. e.g., might want add new tab url like: http://www.bigsurgarrapata.com/contact how in jquery? you want form target attribute action attribute. http://www.w3schools.com/tags/att_form_target.asp browsers have ultimate control on result of target attribute open new tab value of _blank. <form action="contact_submit.php" target="_blank" method="post">
Read more

keyboard - C++ GetAsyncKeyState alternative -

i using windows.h library's getasynckeystate function define key pushed on keyboard . here program returns int value of pushed key example. #include <iostream> #include <windows.h> using namespace std; int main (){ char i; for(i=8;i<190;i++){ if(getasynckeystate(i)== -32767){ cout<<int (i)<<endl; } } return 0; } but trying make opengl simple game , function appeared slow . there function not need cycle 180 times , if statement . thanks you're trying make opengl game, you're going use glfw. if that's case can away using glfwgetkey() such: if (glfwgetkey(glfw_key_up) == glfw_press) {...} if (glfwgetkey(glfw_key_down) == glfw_press) {...} that method nicely explained in tutorial 6 opengl-tutorial.org . should start first tutorial of beginners tutorials , later intermediate tutorials . if you're going make win32 window can listen wm_keydown or wm_keyup messages , virtua...
Read more

android - java.net.UnknownHostException(Unable to resolve host “URL”: No address associated with hostname) -

i developing app testing asp.net webservice in android.for simple webservice adding 2 numbers used.where numbers passed parameters.the result in dialogue.the webserver works in local. when application runs following result exception.i provided internet permission in manifest. java.net.unknownhostexception(unable resolve host “ http://url.com/ ”: no address associated hostname) my code given below main activity public static string rslt=""; /** called when activity first created. */ @override public void oncreate(bundle savedinstancestate) { super.oncreate(savedinstancestate); setcontentview(r.layout.activity_main); button b1=(button)findviewbyid(r.id.button1); final alertdialog ad=new alertdialog.builder(this).create(); b1.setonclicklistener(new onclicklistener() { @override public void onclick(view arg0) { // todo auto-generated method stub try { ...
Read more