{"action":"created","issue":{"url":"https://api.github.com/repos/Syncleus/aparapi/issues/10","repository_url":"https://api.github.com/repos/Syncleus/aparapi","labels_url":"https://api.github.com/repos/Syncleus/aparapi/issues/10/labels{/name}","comments_url":"https://api.github.com/repos/Syncleus/aparapi/issues/10/comments","events_url":"https://api.github.com/repos/Syncleus/aparapi/issues/10/events","html_url":"https://github.com/Syncleus/aparapi/issues/10","id":193239832,"number":10,"title":"Kernel error for MonteCarlo on NVIDIA and AMD GPUs","user":{"login":"freemo","id":3033346,"avatar_url":"https://avatars.githubusercontent.com/u/3033346?v=3","gravatar_id":"","url":"https://api.github.com/users/freemo","html_url":"https://github.com/freemo","followers_url":"https://api.github.com/users/freemo/followers","following_url":"https://api.github.com/users/freemo/following{/other_user}","gists_url":"https://api.github.com/users/freemo/gists{/gist_id}","starred_url":"https://api.github.com/users/freemo/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/freemo/subscriptions","organizations_url":"https://api.github.com/users/freemo/orgs","repos_url":"https://api.github.com/users/freemo/repos","events_url":"https://api.github.com/users/freemo/events{/privacy}","received_events_url":"https://api.github.com/users/freemo/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":1,"created_at":"2016-12-02T23:05:41Z","updated_at":"2016-12-02T23:05:42Z","closed_at":null,"body":"_From @jjfumero on February 17, 2016 10:23_\n\nI am running MonteCarlo simulation within Aparapi. For testing I am using Intel OpenCL locally. I am running with JDK 1.8_65. The kernel that Aparapi generates is correct and the result when I compare to the sequential code is correct as well. However if I use the GPU, NVidia GPU or AMD GPU, the kernel is not correct. One declaration type is missing. \n\nMy understanding is, Aparapi generates the OpenCL kernel indendently of the architecture behind. Bytecodes -> C OpenCL. Is that correct? or is there any communication during the code generation? \n\nHere the details,  this is the Aparapi Kernel: \n\n```\npublic static class MonteCarloKernel extends Kernel {\n\n        private int size;\n        private float[] result;\n\n        public MonteCarloKernel(int size) {\n            this.size = size;\n            result = new float[size];\n        }\n\n        @Override\n        public void run() {\n            int idx = getGlobalId();\n            int iter = 25000;\n\n            long seed = idx;\n            float sum = 0.0f;\n\n            for (int j = 0; j < iter; ++j) {\n                // generate a pseudo random number (you do need it twice)\n                seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n                seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n\n                // this generates a number between 0 and 1 (with an awful entropy)\n                float x = ((float) (seed & 0x0FFFFFFF)) / 268435455f;\n\n                // repeat for y\n                seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n                seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n                float y = ((float) (seed & 0x0FFFFFFF)) / 268435455f;\n\n                float dist = (float) Math.sqrt(x * x + y * y);\n                if (dist <= 1.0f)\n                    sum += 1.0f;\n            }\n            sum *= 4;\n            result[idx] = (float) sum / (float) iter;\n        }\n\n        public boolean checkResult(float[] seq) {\n            for (int i = 0; i < seq.length; i++) {\n                if (Math.abs( (float)(result[i] - seq[i])) > 0.001) {\n                    return false;\n                }\n            }\n            return true;\n        }\n\n        public float[] getResult() {\n            return result;\n        }\n\n        public int getSize() {\n            return size;\n        }\n    }\n\n\n```\n\nIf I use Intel OpenCL: \n\nNAME:         Intel(R) Core(TM) i5-3470 CPU @ 3.20GHz\nVENDOR: Intel(R) Corporation\nTYPE: CPU\nDRIVER: 1.2.0.57\n\nThis is the kernel Aparapi generates (the correct kernel): \n\n```\n\n#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n\ntypedef struct This_s{\n   __global float *result;\n   int passid;\n}This;\nint get_pass_id(This *this){\n   return this->passid;\n}\n__kernel void run(\n   __global float *result, \n   int passid\n){\n   This thisStruct;\n   This* this=&thisStruct;\n   this->result = result;\n   this->passid = passid;\n   {\n      int idx = get_global_id(0);\n      int iter = 25000;\n      long seed = (long)idx;\n      float sum = 0.0f;\n      for (int j = 0; j<iter; j++){\n         seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n         seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n         float x = (float)(seed & 268435455L) / 2.68435456E8f;\n         seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n         seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n         float y = (float)(seed & 268435455L) / 2.68435456E8f;\n         float dist = (float)sqrt((double)((x * x) + (y * y)));\n         if (dist<=1.0f){\n            sum = sum + 1.0f;\n         }\n      }\n      sum = sum * 4.0f;\n      this->result[idx]  = sum / (float)iter;\n      return;\n   }\n}\n\n```\n\nWhen I use Aparapi on NVIDIA or AMD GPUs (same JVM - JDK 1.8.65, but different driver), I get this kernel: \n\n```\n#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n\ntypedef struct This_s{\n   __global float *result;\n   int passid;\n}This;\nint get_pass_id(This *this){\n   return this->passid;\n}\n__kernel void run(\n   __global float *result, \n   int passid\n){\n   This thisStruct;\n   This* this=&thisStruct;\n   this->result = result;\n   this->passid = passid;\n   {\n      int i_1 = get_global_id(0);\n      int i_2 = 25000;\n       l_3 = (long)i_1;\n      float f_5 = 0.0f;\n      int i_6 = 0;\n      for (; i_6<i_2; i_6++){\n         l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n         l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n         float f_7 = (float)(l_3 & 268435455L) / 2.68435456E8f;\n         l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n         l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n         float f_8 = (float)(l_3 & 268435455L) / 2.68435456E8f;\n         float f_9 = (float)sqrt((double)((f_7 * f_7) + (f_8 * f_8)));\n         if (f_9<=1.0f){\n            f_5 = f_5 + 1.0f;\n         }\n      }\n      f_5 = f_5 * 4.0f;\n      this->result[i_1]  = f_5 / (float)i_2;\n      return;\n   }\n}\n```\n\nThere is an error: \n\n```\nclBuildProgram failed\n************************************************\n:21:8: error: use of undeclared identifier 'l_3'\n       l_3 = (long)i_1;\n       ^\n```\n\nNote: \nNVIDIA-SMI 331.79     Driver Version: 331.79 \n\nAMD: \n\n  Name:                                          Hawaii\n  Vendor:                                        Advanced Micro Devices, Inc.\n  Device OpenCL C version:                       OpenCL C 1.2 \n  Driver version:                                1598.5 (VM)\n\n\n_Copied from original issue: aparapi/aparapi#25_"},"comment":{"url":"https://api.github.com/repos/Syncleus/aparapi/issues/comments/264586680","html_url":"https://github.com/Syncleus/aparapi/issues/10#issuecomment-264586680","issue_url":"https://api.github.com/repos/Syncleus/aparapi/issues/10","id":264586680,"user":{"login":"freemo","id":3033346,"avatar_url":"https://avatars.githubusercontent.com/u/3033346?v=3","gravatar_id":"","url":"https://api.github.com/users/freemo","html_url":"https://github.com/freemo","followers_url":"https://api.github.com/users/freemo/followers","following_url":"https://api.github.com/users/freemo/following{/other_user}","gists_url":"https://api.github.com/users/freemo/gists{/gist_id}","starred_url":"https://api.github.com/users/freemo/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/freemo/subscriptions","organizations_url":"https://api.github.com/users/freemo/orgs","repos_url":"https://api.github.com/users/freemo/repos","events_url":"https://api.github.com/users/freemo/events{/privacy}","received_events_url":"https://api.github.com/users/freemo/received_events","type":"User","site_admin":false},"created_at":"2016-12-02T23:05:42Z","updated_at":"2016-12-02T23:05:42Z","body":"_From @SubaruWRC on February 17, 2016 12:53_\n\nHmm, so it's not declaring the variable as a long. Does this also occur if you declare the value as a double?\n\nSent from my iPhone --- Please excuse any typos or autocorrect mistakes\n\n> On Feb 17, 2016, at 2:23 AM, Juan Fumero notifications@github.com wrote:\n> \n> I am running MonteCarlo simulation within Aparapi. For testing I am using Intel OpenCL locally. I am running with JDK 1.8_65. The kernel that Aparapi generates is correct and the result when I compare to the sequential code is correct as well. However if I use the GPU, NVidia GPU or AMD GPU, the kernel is not correct. One declaration type is missing.\n> \n> My understanding is, Aparapi generates the OpenCL kernel indendently of the architecture behind. Bytecodes -> C OpenCL. Is that correct? or is there any communication during the code generation?\n> \n> Here the details, this is the Aparapi Kernel:\n> \n> public static class MonteCarloKernel extends Kernel {\n> \n> ```\n>     private int size;\n>     private float[] result;\n> \n>     public MonteCarloKernel(int size) {\n>         this.size = size;\n>         result = new float[size];\n>     }\n> \n>     @Override\n>     public void run() {\n>         int idx = getGlobalId();\n>         int iter = 25000;\n> \n>         long seed = idx;\n>         float sum = 0.0f;\n> \n>         for (int j = 0; j < iter; ++j) {\n>             // generate a pseudo random number (you do need it twice)\n>             seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n>             seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n> \n>             // this generates a number between 0 and 1 (with an awful entropy)\n>             float x = ((float) (seed & 0x0FFFFFFF)) / 268435455f;\n> \n>             // repeat for y\n>             seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n>             seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);\n>             float y = ((float) (seed & 0x0FFFFFFF)) / 268435455f;\n> \n>             float dist = (float) Math.sqrt(x * x + y * y);\n>             if (dist <= 1.0f)\n>                 sum += 1.0f;\n>         }\n>         sum *= 4;\n>         result[idx] = (float) sum / (float) iter;\n>     }\n> \n>     public boolean checkResult(float[] seq) {\n>         for (int i = 0; i < seq.length; i++) {\n>             if (Math.abs( (float)(result[i] - seq[i])) > 0.001) {\n>                 return false;\n>             }\n>         }\n>         return true;\n>     }\n> \n>     public float[] getResult() {\n>         return result;\n>     }\n> \n>     public int getSize() {\n>         return size;\n>     }\n> }\n> ```\n> \n> If I use Intel OpenCL:\n> \n> NAME: Intel(R) Core(TM) i5-3470 CPU @ 3.20GHz\n> VENDOR: Intel(R) Corporation\n> TYPE: CPU\n> DRIVER: 1.2.0.57\n> \n> This is the kernel Aparapi generates (the correct kernel):\n> \n> #pragma OPENCL EXTENSION cl_khr_fp64 : enable\n> \n> typedef struct This_s{\n>    __global float _result;\n>    int passid;\n> }This;\n> int get_pass_id(This *this){\n>    return this->passid;\n> }\n> __kernel void run(\n>    __global float *result, \n>    int passid\n> ){\n>    This thisStruct;\n>    This_ this=&thisStruct;\n>    this->result = result;\n>    this->passid = passid;\n>    {\n>       int idx = get_global_id(0);\n>       int iter = 25000;\n>       long seed = (long)idx;\n>       float sum = 0.0f;\n>       for (int j = 0; j<iter; j++){\n>          seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n>          seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n>          float x = (float)(seed & 268435455L) / 2.68435456E8f;\n>          seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n>          seed = ((seed * 25214903917L) + 11L) & 281474976710655L;\n>          float y = (float)(seed & 268435455L) / 2.68435456E8f;\n>          float dist = (float)sqrt((double)((x * x) + (y * y)));\n>          if (dist<=1.0f){\n>             sum = sum + 1.0f;\n>          }\n>       }\n>       sum = sum * 4.0f;\n>       this->result[idx]  = sum / (float)iter;\n>       return;\n>    }\n> }\n> \n> When I use Aparapi on NVIDIA or AMD GPUs (same JVM - JDK 1.8.65, but different driver), I get this kernel:\n> \n> #pragma OPENCL EXTENSION cl_khr_fp64 : enable\n> \n> typedef struct This_s{\n>    __global float _result;\n>    int passid;\n> }This;\n> int get_pass_id(This *this){\n>    return this->passid;\n> }\n> __kernel void run(\n>    __global float *result, \n>    int passid\n> ){\n>    This thisStruct;\n>    This_ this=&thisStruct;\n>    this->result = result;\n>    this->passid = passid;\n>    {\n>       int i_1 = get_global_id(0);\n>       int i_2 = 25000;\n>        l_3 = (long)i_1;\n>       float f_5 = 0.0f;\n>       int i_6 = 0;\n>       for (; i_6<i_2; i_6++){\n>          l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n>          l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n>          float f_7 = (float)(l_3 & 268435455L) / 2.68435456E8f;\n>          l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n>          l_3 = ((l_3 * 25214903917L) + 11L) & 281474976710655L;\n>          float f_8 = (float)(l_3 & 268435455L) / 2.68435456E8f;\n>          float f_9 = (float)sqrt((double)((f_7 * f_7) + (f_8 * f_8)));\n>          if (f_9<=1.0f){\n>             f_5 = f_5 + 1.0f;\n>          }\n>       }\n>       f_5 = f_5 * 4.0f;\n>       this->result[i_1]  = f_5 / (float)i_2;\n>       return;\n>    }\n> }\n> There is an error:\n> \n> clBuildProgram failed\n> \n> ---\n> \n> :21:8: error: use of undeclared identifier 'l_3'\n>        l_3 = (long)i_1;\n>        ^\n> Note: \n> NVIDIA-SMI 331.79 Driver Version: 331.79\n> \n> AMD:\n> \n> Name: Hawaii\n> Vendor: Advanced Micro Devices, Inc.\n> Device OpenCL C version: OpenCL C 1.2 \n> Driver version: 1598.5 (VM)\n> \n> —\n> Reply to this email directly or view it on GitHub.\n"}}
{"actor":{"display_login":"freemo"}}