{"action":"created","issue":{"url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers/issues/40","repository_url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers","labels_url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers/issues/40/labels{/name}","comments_url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers/issues/40/comments","events_url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers/issues/40/events","html_url":"https://github.com/eyalroz/cuda-api-wrappers/issues/40","id":296712681,"number":40,"title":"Kernel argument order for cooperative launches is reversed","user":{"login":"codecircuit","id":4549533,"avatar_url":"https://avatars3.githubusercontent.com/u/4549533?v=4","gravatar_id":"","url":"https://api.github.com/users/codecircuit","html_url":"https://github.com/codecircuit","followers_url":"https://api.github.com/users/codecircuit/followers","following_url":"https://api.github.com/users/codecircuit/following{/other_user}","gists_url":"https://api.github.com/users/codecircuit/gists{/gist_id}","starred_url":"https://api.github.com/users/codecircuit/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/codecircuit/subscriptions","organizations_url":"https://api.github.com/users/codecircuit/orgs","repos_url":"https://api.github.com/users/codecircuit/repos","events_url":"https://api.github.com/users/codecircuit/events{/privacy}","received_events_url":"https://api.github.com/users/codecircuit/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":7,"created_at":"2018-02-13T12:00:46Z","updated_at":"2018-02-13T23:29:12Z","closed_at":null,"author_association":"NONE","body":"Compile and execute:\r\n\r\n```cpp\r\n#include <iostream>\r\n\r\n#include \"cuda/api_wrappers.h\"\r\n\r\n__global__ void foo(int a, int b, int c) {\r\n\tif (threadIdx.x == 0) {\r\n\t\tprintf(\"a = %d, b = %d, c = %d\\n\", a, b, c);\r\n\t}\r\n\t__syncthreads();\r\n}\r\n\r\n\r\nint main() {\r\n\tconstexpr cuda::grid_block_dimension_t block_dim = 32;\r\n\tconstexpr cuda::grid_dimension_t grid_dim = 1;\r\n\tint a = 1;\r\n\tint b = 2;\r\n\tint c = 3;\r\n\tstd::cout << \"Launching non-cooperative\" << std::endl;\r\n\tauto config = cuda::make_launch_config(grid_dim, block_dim);\r\n\tcuda::enqueue_launch(cuda::thread_blocks_may_not_cooperate,\r\n\t                     foo,\r\n\t                     cuda::stream::default_stream_id,\r\n\t                     config,\r\n\t                     a, b, c);\r\n\tauto device = cuda::device::current::get();\r\n\tdevice.synchronize();\r\n\r\n\tstd::cout << \"Launching cooperative\" << std::endl;\r\n\tcuda::enqueue_launch(cuda::thread_blocks_may_cooperate,\r\n\t                     foo,\r\n\t                     cuda::stream::default_stream_id,\r\n\t                     config,\r\n\t                     a, b, c);\r\n\tdevice.synchronize();\r\n}\r\n```\r\n\r\nExpected output:\r\n\r\n```\r\nLaunching non-cooperative\r\na = 1, b = 2, c = 3\r\nLaunching cooperative\r\na = 1, b = 2, c = 3\r\n```\r\n\r\nActual output:\r\n\r\n```\r\nLaunching non-cooperative\r\na = 1, b = 2, c = 3\r\nLaunching cooperative\r\na = 3, b = 2, c = 1\r\n```"},"comment":{"url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers/issues/comments/365441359","html_url":"https://github.com/eyalroz/cuda-api-wrappers/issues/40#issuecomment-365441359","issue_url":"https://api.github.com/repos/eyalroz/cuda-api-wrappers/issues/40","id":365441359,"user":{"login":"lahwaacz","id":1289205,"avatar_url":"https://avatars0.githubusercontent.com/u/1289205?v=4","gravatar_id":"","url":"https://api.github.com/users/lahwaacz","html_url":"https://github.com/lahwaacz","followers_url":"https://api.github.com/users/lahwaacz/followers","following_url":"https://api.github.com/users/lahwaacz/following{/other_user}","gists_url":"https://api.github.com/users/lahwaacz/gists{/gist_id}","starred_url":"https://api.github.com/users/lahwaacz/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/lahwaacz/subscriptions","organizations_url":"https://api.github.com/users/lahwaacz/orgs","repos_url":"https://api.github.com/users/lahwaacz/repos","events_url":"https://api.github.com/users/lahwaacz/events{/privacy}","received_events_url":"https://api.github.com/users/lahwaacz/received_events","type":"User","site_admin":false},"created_at":"2018-02-13T23:29:12Z","updated_at":"2018-02-13T23:29:12Z","author_association":"NONE","body":"I bet that after inlining all the lambdas the result will be essentially identical to your original version (which executed in the correct order). As for compilation time, I don't see how you could do less than _O(N)_ instantiations when you actually _need_ to call `f` with each of the _N_ arguments. As a bonus to readability, my version features a reasonably small constant factor due to lack of almost any wrapping code."}}
{"actor":{"display_login":"lahwaacz"}}