{"action":"created","issue":{"url":"https://api.github.com/repos/boostorg/compute/issues/768","repository_url":"https://api.github.com/repos/boostorg/compute","labels_url":"https://api.github.com/repos/boostorg/compute/issues/768/labels{/name}","comments_url":"https://api.github.com/repos/boostorg/compute/issues/768/comments","events_url":"https://api.github.com/repos/boostorg/compute/issues/768/events","html_url":"https://github.com/boostorg/compute/issues/768","id":300187572,"number":768,"title":"meta_kernel::compile is very slow, even for cached source","user":{"login":"Ulfgard","id":15045382,"avatar_url":"https://avatars1.githubusercontent.com/u/15045382?v=4","gravatar_id":"","url":"https://api.github.com/users/Ulfgard","html_url":"https://github.com/Ulfgard","followers_url":"https://api.github.com/users/Ulfgard/followers","following_url":"https://api.github.com/users/Ulfgard/following{/other_user}","gists_url":"https://api.github.com/users/Ulfgard/gists{/gist_id}","starred_url":"https://api.github.com/users/Ulfgard/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/Ulfgard/subscriptions","organizations_url":"https://api.github.com/users/Ulfgard/orgs","repos_url":"https://api.github.com/users/Ulfgard/repos","events_url":"https://api.github.com/users/Ulfgard/events{/privacy}","received_events_url":"https://api.github.com/users/Ulfgard/received_events","type":"User","site_admin":false},"labels":[{"id":161634003,"url":"https://api.github.com/repos/boostorg/compute/labels/design","name":"design","color":"0052cc","default":false},{"id":71700026,"url":"https://api.github.com/repos/boostorg/compute/labels/performance","name":"performance","color":"006b75","default":false}],"state":"open","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":2,"created_at":"2018-02-26T10:40:18Z","updated_at":"2018-05-29T09:56:12Z","closed_at":null,"author_association":"NONE","body":"Hi,\r\n\r\nI am currently benchmarking the overhead of boost::compute. My test-program is computing nothing big and the equivalent cpu program finishes in no measurable runtime. The OpenCL version takes several seconds, so i decided to check with valgrind-callgrind where this time is lost. My program utilizes 11 different kernels and the runtime is the Intel OpenCL cpu implementation. I am using callgrind with O1 optimizations.\r\n\r\n71% of total run-time is spent in 8149 calls to meta_kernel::compile.\r\nInside, we have over all 8149 calls:\r\ntotal-runtime% function\r\n43.21% clCreateKernel\r\n14.5% compute::detail::sha1::sha1\r\n8.7% program::build_with_source\r\n\r\nAs a sanity check: \r\n1.6% meta_kernel::source\r\n1.2% clEnqueueNDRangeKernel\r\n\r\nso, roughly 58% of total running time is spent on figuring out whether the program is already build and creating the kernel afterwards. Compiling and caching works fine, we have exactly 11 calls to build_with_source.\r\n\r\nThe problem is that compute assumes that kernels are throw-away objects. This is unfortunately not the case. e.g. on NVIDIA, clCreateKernel is reported to have an overhead of roughly 1ms. My runtime seems to be in the same order. The programs built with meta_kernel have only one kernel function, so one easy workaround would be to cache the kernel together with the program.\r\n\r\nAlso, sha1 is terribly slow. The hash function should not take an order of magnitude longer than creating the source in the first place. \r\n"},"comment":{"url":"https://api.github.com/repos/boostorg/compute/issues/comments/392720960","html_url":"https://github.com/boostorg/compute/issues/768#issuecomment-392720960","issue_url":"https://api.github.com/repos/boostorg/compute/issues/768","id":392720960,"user":{"login":"Ulfgard","id":15045382,"avatar_url":"https://avatars1.githubusercontent.com/u/15045382?v=4","gravatar_id":"","url":"https://api.github.com/users/Ulfgard","html_url":"https://github.com/Ulfgard","followers_url":"https://api.github.com/users/Ulfgard/followers","following_url":"https://api.github.com/users/Ulfgard/following{/other_user}","gists_url":"https://api.github.com/users/Ulfgard/gists{/gist_id}","starred_url":"https://api.github.com/users/Ulfgard/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/Ulfgard/subscriptions","organizations_url":"https://api.github.com/users/Ulfgard/orgs","repos_url":"https://api.github.com/users/Ulfgard/repos","events_url":"https://api.github.com/users/Ulfgard/events{/privacy}","received_events_url":"https://api.github.com/users/Ulfgard/received_events","type":"User","site_admin":false},"created_at":"2018-05-29T09:56:12Z","updated_at":"2018-05-29T09:56:12Z","author_association":"NONE","body":"I am trying to get a round this issue. I have tuned my meta_kernel so that each expression is always mapping to the same source(i.e. all constant are transformed to kernel variables). Could i then use static thread_local kernels like this?\r\n\r\n`static thread_local compute::kernel k = create_my_kernel(...);`\r\n\r\nWhere create_my_kernel creates the source with meta_kernel, compiles the program and returns the kernel. Would this fix the caching issue as every thread always has a complete list of all kernels? Could I run into problems with artifical driver limitations, e.g. maximum number of programs upported by the OpenCL implementation? I have not found any hint towards that limitation. \r\n"}}
{"actor":{"display_login":"Ulfgard"}}