{"action":"closed","issue":{"url":"https://api.github.com/repos/RadeonOpenCompute/ROCm/issues/347","repository_url":"https://api.github.com/repos/RadeonOpenCompute/ROCm","labels_url":"https://api.github.com/repos/RadeonOpenCompute/ROCm/issues/347/labels{/name}","comments_url":"https://api.github.com/repos/RadeonOpenCompute/ROCm/issues/347/comments","events_url":"https://api.github.com/repos/RadeonOpenCompute/ROCm/issues/347/events","html_url":"https://github.com/RadeonOpenCompute/ROCm/issues/347","id":300994348,"node_id":"MDU6SXNzdWUzMDA5OTQzNDg=","number":347,"title":"ROCM 1.7 - clBuildProgram segfaults","user":{"login":"sampie","id":619607,"node_id":"MDQ6VXNlcjYxOTYwNw==","avatar_url":"https://avatars2.githubusercontent.com/u/619607?v=4","gravatar_id":"","url":"https://api.github.com/users/sampie","html_url":"https://github.com/sampie","followers_url":"https://api.github.com/users/sampie/followers","following_url":"https://api.github.com/users/sampie/following{/other_user}","gists_url":"https://api.github.com/users/sampie/gists{/gist_id}","starred_url":"https://api.github.com/users/sampie/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/sampie/subscriptions","organizations_url":"https://api.github.com/users/sampie/orgs","repos_url":"https://api.github.com/users/sampie/repos","events_url":"https://api.github.com/users/sampie/events{/privacy}","received_events_url":"https://api.github.com/users/sampie/received_events","type":"User","site_admin":false},"labels":[{"id":854942820,"node_id":"MDU6TGFiZWw4NTQ5NDI4MjA=","url":"https://api.github.com/repos/RadeonOpenCompute/ROCm/labels/Compiler%20Functional%20Bug","name":"Compiler Functional Bug","color":"d847b6","default":false}],"state":"closed","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":8,"created_at":"2018-02-28T11:48:10Z","updated_at":"2018-06-03T14:39:36Z","closed_at":"2018-06-03T14:39:36Z","author_association":"NONE","body":"Hi,\r\n\r\nROCM seems to give a segfault when trying to compile a kernel. I got segfault after adding enqueue_kernel() to the code.  \r\n\r\nThread 1 \"program\" received signal SIGSEGV, Segmentation fault.                                                             |        uint64_t a;\r\n0x00007ffff3294db7 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                                |        float b;\r\n(gdb) bt                                                                                                                   |} MType[256];\r\n#0  0x00007ffff3294db7 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#1  0x00007ffff31c8601 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |typedef struct {\r\n#2  0x00007ffff3254f64 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |           uint64_t offset; //  item offset\r\n#3  0x00007ffff3192908 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |           uint64_t len; // item length\r\n#4  0x00007ffff2e58124 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |           uint64_t id; // item id\r\n#5  0x00007ffff3341742 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |           uint64_t thread; // thread id\r\n#6  0x00007ffff33427a0 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |} seqItem;\r\n#7  0x00007ffff317c131 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#8  0x00007ffff318b47b in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#9  0x00007ffff319decb in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |void my_func(__global volatile uint64_t *result) {\r\n#10 0x00007ffff319e44b in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |     int gid = get_global_id(0);\r\n#11 0x00007ffff319e700 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |     result[gid] = 1;\r\n#12 0x00007ffff31a1178 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |}\r\n#13 0x00007ffff319de88 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#14 0x00007ffff319fbf1 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |__kernel void gpu_dot_kernel(\r\n#15 0x00007ffff319e316 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global const uint64_t *lSeqs,\r\n#16 0x00007ffff319e44b in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global const float *fpfSeqs,\r\n#17 0x00007ffff31bccfe in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global const seqItem *seqItems,\r\n#18 0x00007ffff31c68da in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    const uint64_t iCount,\r\n#19 0x00007ffff31331f3 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global const MType *m_amiType,\r\n#20 0x00007ffff3142c98 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global const uint64_t *m_amiTypeLen,\r\n#21 0x00007ffff3143a15 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global const uint64_t *amiTypeThreadOffset,\r\n#22 0x00007ffff3144479 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global float *fResult,\r\n#23 0x00007ffff32696d3 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                    __global uint64_t *lResult) {\r\n#24 0x00007ffff30fef41 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#25 0x00007ffff35240c6 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |    for (uint64_t k = 0; k < iCount; k++) {\r\n#26 0x00007ffff30fd5a4 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#27 0x00007ffff33a5d6e in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |        void (^block)(void) = ^{my_func(\r\n#28 0x00007ffff339647d in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                                        &lResult[k]\r\n#29 0x00007ffff30f6151 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                                        );};\r\n#30 0x00007ffff30f6760 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#31 0x00007ffff2fb58cf in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |        size_t n = 42;\r\n#32 0x00007ffff2fe58eb in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |\r\n#33 0x00007ffff2fb46b6 in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |        enqueue_kernel(get_default_queue(),\r\n#34 0x00007ffff2fbd7fd in ?? () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                            |                   CLK_ENQUEUE_FLAGS_WAIT_KERNEL,\r\n#35 0x00007ffff2f947c9 in clBuildProgram () from /opt/rocm/opencl/lib/x86_64/libamdocl64.so                                |                   ndrange_1D(n),\r\n#36 0x000055555556340c in cl::Program::build (this=0x5555558536d8 <GPU::getInstance()::INSTANCE+24>,                       |                   block);\r\n    options=0x55555560b089 \"-cl-std=CL2.0\", notifyFptr=0x0, data=0x0) at /usr/include/CL/cl2.hpp:6326\r\n\r\n\r\n--- KERNEL CODE ---\r\n\r\n`\r\n#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\r\n\r\ntypedef char int8_t;\r\ntypedef unsigned char uint8_t;\r\n\r\ntypedef short int16_t;\r\ntypedef unsigned short uint16_t;\r\n\r\ntypedef int int32_t;\r\ntypedef unsigned int uint32_t;\r\n\r\ntypedef long int64_t;\r\ntypedef unsigned long uint64_t;\r\n\r\ntypedef struct {\r\n        uint64_t a; \r\n        float b; \r\n} MType[256];\r\n\r\ntypedef struct {\r\n           uint64_t offset; //  item offset\r\n           uint64_t len; // item length\r\n           uint64_t id; // item id\r\n           uint64_t thread; // thread id\r\n} seqItem;\r\n\r\n\r\nvoid my_func(__global volatile uint64_t *result) {\r\n     int gid = get_global_id(0);\r\n     result[gid] = 1;\r\n}\r\n\r\n__kernel void gpu_dot_kernel(\r\n                    __global const uint64_t *lSeqs,\r\n                    __global const float *fpfSeqs,\r\n                    __global const seqItem *seqItems,\r\n                    const uint64_t iCount,\r\n                    __global const MType *m_amiType,\r\n                    __global const uint64_t *m_amiTypeLen,\r\n                    __global const uint64_t *amiTypeThreadOffset,\r\n                    __global float *fResult,\r\n                    __global uint64_t *lResult) {\r\n\r\n    for (uint64_t k = 0; k < iCount; k++) {\r\n\r\n        void (^block)(void) = ^{my_func( \r\n                                        &lResult[k]\r\n                                        );};\r\n\r\n        size_t n = 42;\r\n\r\n        enqueue_kernel(get_default_queue(),\r\n                   CLK_ENQUEUE_FLAGS_WAIT_KERNEL,\r\n                   ndrange_1D(n), \r\n                   block);\r\n\r\n    }\r\n\r\n}\r\n`\r\n\r\n"}}
{"actor":{"display_login":"gstoner"}}