{"action":"opened","issue":{"url":"https://api.github.com/repos/ClearVolume/ClearVolume/issues/32","labels_url":"https://api.github.com/repos/ClearVolume/ClearVolume/issues/32/labels{/name}","comments_url":"https://api.github.com/repos/ClearVolume/ClearVolume/issues/32/comments","events_url":"https://api.github.com/repos/ClearVolume/ClearVolume/issues/32/events","html_url":"https://github.com/ClearVolume/ClearVolume/issues/32","id":98362354,"number":32,"title":"Denoising with default settings throws OpenCL error","user":{"login":"skalarproduktraum","id":586495,"avatar_url":"https://avatars.githubusercontent.com/u/586495?v=3","gravatar_id":"","url":"https://api.github.com/users/skalarproduktraum","html_url":"https://github.com/skalarproduktraum","followers_url":"https://api.github.com/users/skalarproduktraum/followers","following_url":"https://api.github.com/users/skalarproduktraum/following{/other_user}","gists_url":"https://api.github.com/users/skalarproduktraum/gists{/gist_id}","starred_url":"https://api.github.com/users/skalarproduktraum/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/skalarproduktraum/subscriptions","organizations_url":"https://api.github.com/users/skalarproduktraum/orgs","repos_url":"https://api.github.com/users/skalarproduktraum/repos","events_url":"https://api.github.com/users/skalarproduktraum/events{/privacy}","received_events_url":"https://api.github.com/users/skalarproduktraum/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"milestone":null,"comments":0,"created_at":"2015-07-31T10:32:07Z","updated_at":"2015-07-31T10:32:07Z","closed_at":null,"body":"When using `LocalMeans` mode for denoising, I get the following OpenCL-related exception:\r\n```\r\n>>> ) (make sure to log all errors with environment variable CL_LOG_ERRORS=stdout)\r\n\tat sun.reflect.NativeConstructorAccessorImpl.newInstance0(Native Method)\r\n\tat sun.reflect.NativeConstructorAccessorImpl.newInstance(NativeConstructorAccessorImpl.java:62)\r\n\tat sun.reflect.DelegatingConstructorAccessorImpl.newInstance(DelegatingConstructorAccessorImpl.java:45)\r\n\tat java.lang.reflect.Constructor.newInstance(Constructor.java:422)\r\n\tat java.lang.Class.newInstance(Class.java:442)\r\n\tat com.nativelibs4java.opencl.CLException.error(CLException.java:302)\r\n\tat com.nativelibs4java.opencl.CLKernel.setKernelArg(CLKernel.java:273)\r\n\tat com.nativelibs4java.opencl.CLKernel.setArg(CLKernel.java:357)\r\n\tat com.nativelibs4java.opencl.CLKernel.setObjectArg(CLKernel.java:204)\r\n\tat com.nativelibs4java.opencl.CLKernel.setArgs(CLKernel.java:191)\r\n\tat clearvolume.renderer.processors.impl.OpenCLDenoise.process_nlm_fast(OpenCLDenoise.java:293)\r\n\tat clearvolume.renderer.processors.impl.OpenCLDenoise.process(OpenCLDenoise.java:203)\r\n\tat clearvolume.renderer.opencl.OpenCLVolumeRenderer.runProcessorHook(OpenCLVolumeRenderer.java:544)\r\n\tat clearvolume.renderer.opencl.OpenCLVolumeRenderer.updateBufferAndRunKernel(OpenCLVolumeRenderer.java:381)\r\n\tat clearvolume.renderer.opencl.OpenCLVolumeRenderer.renderVolume(OpenCLVolumeRenderer.java:289)\r\n\tat clearvolume.renderer.cleargl.ClearGLVolumeRenderer.displayInternal(ClearGLVolumeRenderer.java:762)\r\n\tat clearvolume.renderer.cleargl.ClearGLVolumeRenderer.display(ClearGLVolumeRenderer.java:699)\r\n\tat jogamp.opengl.GLDrawableHelper.displayImpl(GLDrawableHelper.java:691)\r\n\tat jogamp.opengl.GLDrawableHelper.display(GLDrawableHelper.java:673)\r\n\tat jogamp.opengl.GLAutoDrawableBase$2.run(GLAutoDrawableBase.java:442)\r\n\tat jogamp.opengl.GLDrawableHelper.invokeGLImpl(GLDrawableHelper.java:1277)\r\n\tat jogamp.opengl.GLDrawableHelper.invokeGL(GLDrawableHelper.java:1131)\r\n\tat com.jogamp.newt.opengl.GLWindow.display(GLWindow.java:680)\r\n\tat com.jogamp.opengl.util.AWTAnimatorImpl.display(AWTAnimatorImpl.java:77)\r\n\tat com.jogamp.opengl.util.AnimatorBase.display(AnimatorBase.java:451)\r\n\tat com.jogamp.opengl.util.FPSAnimator$MainTask.run(FPSAnimator.java:178)\r\n\tat java.util.TimerThread.mainLoop(Timer.java:555)\r\n\tat java.util.TimerThread.run(Timer.java:505)\r\ncom.nativelibs4java.opencl.CLException$InvalidArgIndex: InvalidArgIndex (kernel name = dist, num args = 5, arg index = 5, source = <<<\r\n\r\n\t#ifndef FS\r\n\t#define FS 1\r\n\t#endif\r\n\r\n\t#ifndef BS\r\n\t#define BS 2\r\n\t#endif /* BS */\r\n\r\n\t#define NPATCH ((2.f*FS+1)*(2.f*FS+1)*(2.f*FS+1))\r\n\r\n\r\n\r\n\r\n\t__kernel void dist(__read_only image3d_t input,__global float * output, const int dx,const int dy,const int dz){\r\n\r\n\r\n\tconst sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |\tCLK_ADDRESS_CLAMP_TO_EDGE |\tCLK_FILTER_NEAREST ;\r\n\r\n\tuint i = get_global_id(0);\r\n\tuint j = get_global_id(1);\r\n\tuint k = get_global_id(2);\r\n\r\n\tint Nx = get_global_size(0);\r\n\tint Ny = get_global_size(1);\r\n\r\n\tfloat pix1  = read_imagef(input,sampler,(int4)(i,j,k,0)).x;\r\n\tfloat pix2  = read_imagef(input,sampler,(int4)(i+dx,j+dy,k+dz,0)).x;\r\n\r\n\tfloat d = (pix1-pix2);\r\n\r\n\r\n\td = d*d/NPATCH;\r\n\r\n\r\n\toutput[i+Nx*j+Nx*Ny*k] = d;\r\n\r\n\r\n\t}\r\n\r\n\r\n\t__kernel void convolve(\t__global float * \tinput,\r\n\t\t\t\t\t\t\t__global float * \toutput,\r\n\t\t\t\t\t\t\tconst int \t\t\tflag)\r\n\t{\r\n\r\n\t  // flag = 1 -> in x axis\r\n\t  // flag = 2 -> in y axis\r\n\t  // flag = 4 -> in z axis\r\n\r\n\t  int i = get_global_id(0);\r\n\t  int j = get_global_id(1);\r\n\t  int k = get_global_id(2);\r\n\r\n\r\n\t  int Nx = get_global_size(0);\r\n\t  int Ny = get_global_size(1);\r\n\t  int Nz = get_global_size(2);\r\n\r\n\t  const int dx = flag & 1;\r\n\t  const int dy = (flag&2)/2;\r\n\t  const int dz = (flag&4)/4;\r\n\r\n\t\tif (flag==0){\r\n\t\t\toutput[i+Nx*j+Nx*Ny*k] = input[i+Nx*j+Nx*Ny*k];\r\n\t\t\treturn;\r\n\t\t}\r\n\r\n\r\n\t  const int Nh = 2*FS+1;\r\n\r\n\t  float res = 0.f;\r\n\t  int delta = (Nh-1)/2;\r\n\r\n\r\n\t  for (int p = 0; p < Nh; ++p){\r\n\r\n\t\tint i1 = i+dx*(p-delta);\r\n\t\tint j1 = j+dy*(p-delta);\r\n\t\tint k1 = k+dz*(p-delta);\r\n\r\n\t\tfloat val = 0.f;\r\n\r\n\r\n\t\t//clamp to border\r\n\t\ti1 = clamp(i1,0,Nx-1);\r\n\t\tj1 = clamp(j1,0,Ny-1);\r\n\t\tk1 = clamp(k1,0,Nz-1);\r\n\r\n\t\tval = input[i1+Nx*j1+Nx*Ny*k1];\r\n\r\n\t\tres += val;\r\n\r\n\t  }\r\n\r\n\t   output[i+Nx*j+Nx*Ny*k] = res;\r\n\r\n\t }\r\n\r\n\t__kernel void computePlus(\t__read_only image3d_t input,\r\n\t\t\t\t\t\t\t\t__global float * \tdistImg,\r\n\t\t\t\t\t\t\t  __global float* accBuf,__global float* weightBuf,\r\n\t\t\t\t\t\t\t  const int dx,const int dy,const int dz, const float sigma){\r\n\r\n\r\n\t    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |\tCLK_ADDRESS_CLAMP_TO_EDGE |\tCLK_FILTER_NEAREST ;\r\n\r\n\r\n\t  uint i0 = get_global_id(0);\r\n\t  uint j0 = get_global_id(1);\r\n\t  uint k0 = get_global_id(2);\r\n\r\n\t  int Nx = get_global_size(0);\r\n\t  int Ny = get_global_size(1);\r\n\t  int Nz = get_global_size(2);\r\n\r\n\t  float dist  = distImg[i0+Nx*j0+Nx*Ny*k0];\r\n\r\n\t  float pix  = pix  = read_imagef(input,sampler,(int4)(i0+dx,j0+dy,k0+dz,0)).x;\r\n\r\n\t  float weight = exp(-1.f*dist/sigma/sigma);\r\n\r\n\t  accBuf[i0+Nx*j0+Nx*Ny*k0] += (float)(weight*pix);\r\n\t  weightBuf[i0+Nx*j0+Nx*Ny*k0] += (float)(weight);\r\n\r\n\r\n\t}\r\n\r\n\t__kernel void computeMinus(\t__read_only image3d_t input,\r\n\t\t\t\t\t\t\t\t__global float * \tdistImg,\r\n\t\t\t\t\t\t\t  __global float* accBuf,__global float* weightBuf,\r\n\t\t\t\t\t\t\t  const int dx,const int dy,const int dz, const float sigma){\r\n\r\n\r\n\t    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |\tCLK_ADDRESS_CLAMP_TO_EDGE |\tCLK_FILTER_NEAREST ;\r\n\r\n\r\n\t  uint i0 = get_global_id(0);\r\n\t  uint j0 = get_global_id(1);\r\n\t  uint k0 = get_global_id(2);\r\n\r\n\t  int Nx = get_global_size(0);\r\n\t  int Ny = get_global_size(1);\r\n\t  int Nz = get_global_size(2);\r\n\r\n\t  int i1 = clamp((int)(i0-dx),(int)0,(int)(Nx-1));\r\n\t  int j1 = clamp((int)(j0-dy),(int)0,(int)(Ny-1));\r\n\t  int k1 = clamp((int)(k0-dz),(int)0,(int)(Nz-1));\r\n\r\n\t  float dist  = distImg[i1+Nx*j1+Nx*Ny*k1];\r\n\r\n\t  float pix  = pix  = read_imagef(input,sampler,(int4)(i0-dx,j0-dy,k0-dz,0)).x;\r\n\r\n\t  float weight = exp(-1.f*dist/sigma/sigma);\r\n\r\n\t  accBuf[i0+Nx*j0+Nx*Ny*k0] += (float)(weight*pix);\r\n\t  weightBuf[i0+Nx*j0+Nx*Ny*k0] += (float)(weight);\r\n\r\n\r\n\t}\r\n\r\n\t__kernel void assemble(\t__global float * accBuf,\r\n\t\t\t\t\t\t\t__global float* weightBuf,\r\n\t\t\t\t\t\t\t__global float* outBuf){\r\n\r\n\r\n\r\n\t  uint i0 = get_global_id(0);\r\n\t  uint j0 = get_global_id(1);\r\n\t  uint k0 = get_global_id(2);\r\n\r\n\t  int Nx = get_global_size(0);\r\n\t  int Ny = get_global_size(1);\r\n\r\n\t  outBuf[i0+Nx*j0+Nx*Ny*k0] = accBuf[i0+Nx*j0+Nx*Ny*k0]/weightBuf[i0+Nx*j0+Nx*Ny*k0];\r\n\r\n\r\n\r\n\t}\r\n\r\n\r\n\t__kernel void startup(\t__global float * accBuf,\r\n\t\t\t\t\t\t\t__global float* weightBuf){\r\n\r\n\r\n\r\n\t  uint i0 = get_global_id(0);\r\n\t  uint j0 = get_global_id(1);\r\n\t  uint k0 = get_global_id(2);\r\n\r\n\t  int Nx = get_global_size(0);\r\n\t  int Ny = get_global_size(1);\r\n\r\n\t  accBuf[i0+Nx*j0+Nx*Ny*k0] = 0.f;\r\n\t  weightBuf[i0+Nx*j0+Nx*Ny*k0] = 0.f;\r\n\r\n\r\n\t}\r\n```\r\n\r\nThe error does not occur when using `BilateralFiltering` instead of `LocalMeans`.\r\nPlatform is Java 1.8.0_45 on MacOS X 10.11/DP5 with OpenCL running on an Intel Iris Pro."}}