{"action":"closed","number":2338,"pull_request":{"url":"https://api.github.com/repos/dmlc/tvm/pulls/2338","id":240994520,"node_id":"MDExOlB1bGxSZXF1ZXN0MjQwOTk0NTIw","html_url":"https://github.com/dmlc/tvm/pull/2338","diff_url":"https://github.com/dmlc/tvm/pull/2338.diff","patch_url":"https://github.com/dmlc/tvm/pull/2338.patch","issue_url":"https://api.github.com/repos/dmlc/tvm/issues/2338","number":2338,"state":"closed","locked":false,"title":"[TOPI, CUDA] Update cuda softmax schedule for spatial inputs","user":{"login":"masahi","id":1776403,"node_id":"MDQ6VXNlcjE3NzY0MDM=","avatar_url":"https://avatars3.githubusercontent.com/u/1776403?v=4","gravatar_id":"","url":"https://api.github.com/users/masahi","html_url":"https://github.com/masahi","followers_url":"https://api.github.com/users/masahi/followers","following_url":"https://api.github.com/users/masahi/following{/other_user}","gists_url":"https://api.github.com/users/masahi/gists{/gist_id}","starred_url":"https://api.github.com/users/masahi/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/masahi/subscriptions","organizations_url":"https://api.github.com/users/masahi/orgs","repos_url":"https://api.github.com/users/masahi/repos","events_url":"https://api.github.com/users/masahi/events{/privacy}","received_events_url":"https://api.github.com/users/masahi/received_events","type":"User","site_admin":false},"body":"As pointed out in https://discuss.tvm.ai/t/softmax-is-really-slow/1365/4,\r\n\r\nThe current cuda softmax schedule is hard coded for 2D inputs (after fully connected layer), so it is super slow on spatial inputs (image segmentation, object detection etc). For (1, 16, 256, 256) input, it generates the following schedule.\r\n\r\n\r\n```\r\n// attr [compute] storage_scope = \"global\"\r\nallocate compute[float32 * 1 * 256 * 256]\r\n// attr [compute] storage_scope = \"global\"\r\nallocate compute[float32 * 1 * 256 * 256]\r\nproduce compute {\r\n  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 1\r\n  for (i1, 0, 256) {\r\n    for (i2, 0, 256) {\r\n      compute[((i1*256) + i2)] = -340282346638528859811704183484516925440.000000f\r\n      for (k, 0, 16) {\r\n        compute[((i1*256) + i2)] = max(compute[((i1*256) + i2)], A[(((i1*256) + i2) + (k*65536))])\r\n      }\r\n    }\r\n  }\r\n}\r\nproduce compute {\r\n  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 1\r\n  // attr [compute.rf] storage_scope = \"local\"\r\n  allocate compute.rf[float32 * 1 * 1 * 1 * 1]\r\n  // attr [reduce_temp0] storage_scope = \"local\"\r\n  allocate reduce_temp0[float32 * 1]\r\n  for (ax1, 0, 256) {\r\n    for (ax2, 0, 256) {\r\n      // attr [iter_var(threadIdx.x, Range(min=0, extent=64), threadIdx.x)] thread_extent = 64\r\n      produce compute.rf {\r\n        compute.rf[0] = 0.000000f\r\n        if ((threadIdx.x < 16)) {\r\n          compute.rf[0] = (compute.rf[0] + exp((A[(((ax1*256) + ax2) + (threadIdx.x*65536))] - compute[((ax1*256) + ax2)])))\r\n        }\r\n      }\r\n      // attr [comm_reducer(result=[(x + y)], lhs=[x], rhs=[y], identity_element=[0.000000f])] reduce_scope = reinterpret((uint64)0)\r\n      tvm_thread_allreduce((uint32)1, compute.rf[0], (uint1)1, reduce_temp0, threadIdx.x)\r\n      if ((threadIdx.x == 0)) {\r\n        compute[((ax1*256) + ax2)] = reduce_temp0[0]\r\n      }\r\n    }\r\n  }\r\n}\r\nproduce compute {\r\n  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 1\r\n  // attr [iter_var(threadIdx.x, Range(min=0, extent=64), threadIdx.x)] thread_extent = 64\r\n  for (i2, 0, 256) {\r\n    for (i3, 0, 256) {\r\n      if (likely((threadIdx.x < 16))) {\r\n        compute[(((i2 + (threadIdx.x*256))*256) + i3)] = (exp((A[(((i2 + (threadIdx.x*256))*256) + i3)] - compute[((i2*256) + i3)]))/compute[((i2*256) + i3)])\r\n      }\r\n    }\r\n  }\r\n}\r\n\r\n```\r\nI updated the softmax schedule so that it works well for 4D and 5D inputs. Since the number of classes is \r\n typically small for these inputs, I simply use injective schedule (i.e., without using shared memory reduce).  Below is the updated schedule with this PR. @vinx13  @merrymercy @eqy please review.\r\n\r\n```\r\n// attr [compute] storage_scope = \"global\"\r\nallocate compute[float32 * 1 * 256 * 256]\r\n// attr [compute] storage_scope = \"global\"\r\nallocate compute[float32 * 1 * 256 * 256]\r\nproduce compute {\r\n  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 128\r\n  // attr [iter_var(threadIdx.x, , threadIdx.x)] thread_extent = 512\r\n  compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))] = -340282346638528859811704183484516925440.000000f\r\n  for (k, 0, 16) {\r\n    compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))] = max(compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))], A[(((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256)) + (k*65536))])\r\n  }\r\n}\r\nproduce compute {\r\n  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 128\r\n  // attr [iter_var(threadIdx.x, , threadIdx.x)] thread_extent = 512\r\n  compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))] = 0.000000f\r\n  for (k, 0, 16) {\r\n    compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))] = (compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))] + exp((A[(((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256)) + (k*65536))] - compute[((((blockIdx.x*2) + (threadIdx.x/256))*256) + (threadIdx.x % 256))])))\r\n  }\r\n}\r\nproduce compute {\r\n  // attr [iter_var(blockIdx.x, , blockIdx.x)] thread_extent = 256\r\n  // attr [iter_var(threadIdx.x, , threadIdx.x)] thread_extent = 512\r\n  for (i0.i1.fused.i2.fused.i3.fused.outer, 0, 8) {\r\n    compute[((((((blockIdx.x*2) + (threadIdx.x/256))/256)*65536) + (((((blockIdx.x*2) + (threadIdx.x/256)) % 256)*256) + (threadIdx.x % 256))) + (i0.i1.fused.i2.fused.i3.fused.outer*131072))] = (exp((A[((((((blockIdx.x*2) + (threadIdx.x/256))/256)*65536) + (((((blockIdx.x*2) + (threadIdx.x/256)) % 256)*256) + (threadIdx.x % 256))) + (i0.i1.fused.i2.fused.i3.fused.outer*131072))] - compute[(((((blockIdx.x*2) + (threadIdx.x/256)) % 256)*256) + (threadIdx.x % 256))]))/compute[(((((blockIdx.x*2) + (threadIdx.x/256)) % 256)*256) + (threadIdx.x % 256))])\r\n  }\r\n}\r\n```","created_at":"2018-12-26T15:08:12Z","updated_at":"2018-12-29T18:09:36Z","closed_at":"2018-12-29T18:09:36Z","merged_at":"2018-12-29T18:09:36Z","merge_commit_sha":"3708b311be20f2a0ffb8731d1a40ec545fd8eba6","assignee":null,"assignees":[],"requested_reviewers":[],"requested_teams":[],"labels":[{"id":858780915,"node_id":"MDU6TGFiZWw4NTg3ODA5MTU=","url":"https://api.github.com/repos/dmlc/tvm/labels/status:%20need%20review","name":"status: need review","color":"fbca04","default":false}],"milestone":null,"commits_url":"https://api.github.com/repos/dmlc/tvm/pulls/2338/commits","review_comments_url":"https://api.github.com/repos/dmlc/tvm/pulls/2338/comments","review_comment_url":"https://api.github.com/repos/dmlc/tvm/pulls/comments{/number}","comments_url":"https://api.github.com/repos/dmlc/tvm/issues/2338/comments","statuses_url":"https://api.github.com/repos/dmlc/tvm/statuses/1d5a29526a58f2674c1fdaf80a99aa566aa86d03","head":{"label":"masahi:cuda-softmax","ref":"cuda-softmax","sha":"1d5a29526a58f2674c1fdaf80a99aa566aa86d03","user":{"login":"masahi","id":1776403,"node_id":"MDQ6VXNlcjE3NzY0MDM=","avatar_url":"https://avatars3.githubusercontent.com/u/1776403?v=4","gravatar_id":"","url":"https://api.github.com/users/masahi","html_url":"https://github.com/masahi","followers_url":"https://api.github.com/users/masahi/followers","following_url":"https://api.github.com/users/masahi/following{/other_user}","gists_url":"https://api.github.com/users/masahi/gists{/gist_id}","starred_url":"https://api.github.com/users/masahi/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/masahi/subscriptions","organizations_url":"https://api.github.com/users/masahi/orgs","repos_url":"https://api.github.com/users/masahi/repos","events_url":"https://api.github.com/users/masahi/events{/privacy}","received_events_url":"https://api.github.com/users/masahi/received_events","type":"User","site_admin":false},"repo":{"id":138661036,"node_id":"MDEwOlJlcG9zaXRvcnkxMzg2NjEwMzY=","name":"tvm","full_name":"masahi/tvm","private":false,"owner":{"login":"masahi","id":1776403,"node_id":"MDQ6VXNlcjE3NzY0MDM=","avatar_url":"https://avatars3.githubusercontent.com/u/1776403?v=4","gravatar_id":"","url":"https://api.github.com/users/masahi","html_url":"https://github.com/masahi","followers_url":"https://api.github.com/users/masahi/followers","following_url":"https://api.github.com/users/masahi/following{/other_user}","gists_url":"https://api.github.com/users/masahi/gists{/gist_id}","starred_url":"https://api.github.com/users/masahi/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/masahi/subscriptions","organizations_url":"https://api.github.com/users/masahi/orgs","repos_url":"https://api.github.com/users/masahi/repos","events_url":"https://api.github.com/users/masahi/events{/privacy}","received_events_url":"https://api.github.com/users/masahi/received_events","type":"User","site_admin":false},"html_url":"https://github.com/masahi/tvm","description":"Open deep learning compiler stack for cpu, gpu and specialized accelerators","fork":true,"url":"https://api.github.com/repos/masahi/tvm","forks_url":"https://api.github.com/repos/masahi/tvm/forks","keys_url":"https://api.github.com/repos/masahi/tvm/keys{/key_id}","collaborators_url":"https://api.github.com/repos/masahi/tvm/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/masahi/tvm/teams","hooks_url":"https://api.github.com/repos/masahi/tvm/hooks","issue_events_url":"https://api.github.com/repos/masahi/tvm/issues/events{/number}","events_url":"https://api.github.com/repos/masahi/tvm/events","assignees_url":"https://api.github.com/repos/masahi/tvm/assignees{/user}","branches_url":"https://api.github.com/repos/masahi/tvm/branches{/branch}","tags_url":"https://api.github.com/repos/masahi/tvm/tags","blobs_url":"https://api.github.com/repos/masahi/tvm/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/masahi/tvm/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/masahi/tvm/git/refs{/sha}","trees_url":"https://api.github.com/repos/masahi/tvm/git/trees{/sha}","statuses_url":"https://api.github.com/repos/masahi/tvm/statuses/{sha}","languages_url":"https://api.github.com/repos/masahi/tvm/languages","stargazers_url":"https://api.github.com/repos/masahi/tvm/stargazers","contributors_url":"https://api.github.com/repos/masahi/tvm/contributors","subscribers_url":"https://api.github.com/repos/masahi/tvm/subscribers","subscription_url":"https://api.github.com/repos/masahi/tvm/subscription","commits_url":"https://api.github.com/repos/masahi/tvm/commits{/sha}","git_commits_url":"https://api.github.com/repos/masahi/tvm/git/commits{/sha}","comments_url":"https://api.github.com/repos/masahi/tvm/comments{/number}","issue_comment_url":"https://api.github.com/repos/masahi/tvm/issues/comments{/number}","contents_url":"https://api.github.com/repos/masahi/tvm/contents/{+path}","compare_url":"https://api.github.com/repos/masahi/tvm/compare/{base}...{head}","merges_url":"https://api.github.com/repos/masahi/tvm/merges","archive_url":"https://api.github.com/repos/masahi/tvm/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/masahi/tvm/downloads","issues_url":"https://api.github.com/repos/masahi/tvm/issues{/number}","pulls_url":"https://api.github.com/repos/masahi/tvm/pulls{/number}","milestones_url":"https://api.github.com/repos/masahi/tvm/milestones{/number}","notifications_url":"https://api.github.com/repos/masahi/tvm/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/masahi/tvm/labels{/name}","releases_url":"https://api.github.com/repos/masahi/tvm/releases{/id}","deployments_url":"https://api.github.com/repos/masahi/tvm/deployments","created_at":"2018-06-25T23:39:51Z","updated_at":"2018-06-25T23:39:55Z","pushed_at":"2018-12-26T15:06:31Z","git_url":"git://github.com/masahi/tvm.git","ssh_url":"git@github.com:masahi/tvm.git","clone_url":"https://github.com/masahi/tvm.git","svn_url":"https://github.com/masahi/tvm","homepage":"https://tvm.ai","size":9650,"stargazers_count":0,"watchers_count":0,"language":"C++","has_issues":false,"has_projects":true,"has_downloads":true,"has_wiki":false,"has_pages":false,"forks_count":0,"mirror_url":null,"archived":false,"open_issues_count":0,"license":{"key":"apache-2.0","name":"Apache License 2.0","spdx_id":"Apache-2.0","url":"https://api.github.com/licenses/apache-2.0","node_id":"MDc6TGljZW5zZTI="},"forks":0,"open_issues":0,"watchers":0,"default_branch":"master"}},"base":{"label":"dmlc:master","ref":"master","sha":"b076cad542524cb3744149d953c341b5815f6474","user":{"login":"dmlc","id":11508361,"node_id":"MDEyOk9yZ2FuaXphdGlvbjExNTA4MzYx","avatar_url":"https://avatars2.githubusercontent.com/u/11508361?v=4","gravatar_id":"","url":"https://api.github.com/users/dmlc","html_url":"https://github.com/dmlc","followers_url":"https://api.github.com/users/dmlc/followers","following_url":"https://api.github.com/users/dmlc/following{/other_user}","gists_url":"https://api.github.com/users/dmlc/gists{/gist_id}","starred_url":"https://api.github.com/users/dmlc/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/dmlc/subscriptions","organizations_url":"https://api.github.com/users/dmlc/orgs","repos_url":"https://api.github.com/users/dmlc/repos","events_url":"https://api.github.com/users/dmlc/events{/privacy}","received_events_url":"https://api.github.com/users/dmlc/received_events","type":"Organization","site_admin":false},"repo":{"id":70746484,"node_id":"MDEwOlJlcG9zaXRvcnk3MDc0NjQ4NA==","name":"tvm","full_name":"dmlc/tvm","private":false,"owner":{"login":"dmlc","id":11508361,"node_id":"MDEyOk9yZ2FuaXphdGlvbjExNTA4MzYx","avatar_url":"https://avatars2.githubusercontent.com/u/11508361?v=4","gravatar_id":"","url":"https://api.github.com/users/dmlc","html_url":"https://github.com/dmlc","followers_url":"https://api.github.com/users/dmlc/followers","following_url":"https://api.github.com/users/dmlc/following{/other_user}","gists_url":"https://api.github.com/users/dmlc/gists{/gist_id}","starred_url":"https://api.github.com/users/dmlc/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/dmlc/subscriptions","organizations_url":"https://api.github.com/users/dmlc/orgs","repos_url":"https://api.github.com/users/dmlc/repos","events_url":"https://api.github.com/users/dmlc/events{/privacy}","received_events_url":"https://api.github.com/users/dmlc/received_events","type":"Organization","site_admin":false},"html_url":"https://github.com/dmlc/tvm","description":"Open deep learning compiler stack for cpu, gpu and specialized accelerators","fork":false,"url":"https://api.github.com/repos/dmlc/tvm","forks_url":"https://api.github.com/repos/dmlc/tvm/forks","keys_url":"https://api.github.com/repos/dmlc/tvm/keys{/key_id}","collaborators_url":"https://api.github.com/repos/dmlc/tvm/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/dmlc/tvm/teams","hooks_url":"https://api.github.com/repos/dmlc/tvm/hooks","issue_events_url":"https://api.github.com/repos/dmlc/tvm/issues/events{/number}","events_url":"https://api.github.com/repos/dmlc/tvm/events","assignees_url":"https://api.github.com/repos/dmlc/tvm/assignees{/user}","branches_url":"https://api.github.com/repos/dmlc/tvm/branches{/branch}","tags_url":"https://api.github.com/repos/dmlc/tvm/tags","blobs_url":"https://api.github.com/repos/dmlc/tvm/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/dmlc/tvm/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/dmlc/tvm/git/refs{/sha}","trees_url":"https://api.github.com/repos/dmlc/tvm/git/trees{/sha}","statuses_url":"https://api.github.com/repos/dmlc/tvm/statuses/{sha}","languages_url":"https://api.github.com/repos/dmlc/tvm/languages","stargazers_url":"https://api.github.com/repos/dmlc/tvm/stargazers","contributors_url":"https://api.github.com/repos/dmlc/tvm/contributors","subscribers_url":"https://api.github.com/repos/dmlc/tvm/subscribers","subscription_url":"https://api.github.com/repos/dmlc/tvm/subscription","commits_url":"https://api.github.com/repos/dmlc/tvm/commits{/sha}","git_commits_url":"https://api.github.com/repos/dmlc/tvm/git/commits{/sha}","comments_url":"https://api.github.com/repos/dmlc/tvm/comments{/number}","issue_comment_url":"https://api.github.com/repos/dmlc/tvm/issues/comments{/number}","contents_url":"https://api.github.com/repos/dmlc/tvm/contents/{+path}","compare_url":"https://api.github.com/repos/dmlc/tvm/compare/{base}...{head}","merges_url":"https://api.github.com/repos/dmlc/tvm/merges","archive_url":"https://api.github.com/repos/dmlc/tvm/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/dmlc/tvm/downloads","issues_url":"https://api.github.com/repos/dmlc/tvm/issues{/number}","pulls_url":"https://api.github.com/repos/dmlc/tvm/pulls{/number}","milestones_url":"https://api.github.com/repos/dmlc/tvm/milestones{/number}","notifications_url":"https://api.github.com/repos/dmlc/tvm/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/dmlc/tvm/labels{/name}","releases_url":"https://api.github.com/repos/dmlc/tvm/releases{/id}","deployments_url":"https://api.github.com/repos/dmlc/tvm/deployments","created_at":"2016-10-12T22:20:28Z","updated_at":"2018-12-29T10:47:08Z","pushed_at":"2018-12-29T18:09:36Z","git_url":"git://github.com/dmlc/tvm.git","ssh_url":"git@github.com:dmlc/tvm.git","clone_url":"https://github.com/dmlc/tvm.git","svn_url":"https://github.com/dmlc/tvm","homepage":"https://tvm.ai","size":9282,"stargazers_count":2582,"watchers_count":2582,"language":"Python","has_issues":true,"has_projects":true,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":621,"mirror_url":null,"archived":false,"open_issues_count":81,"license":{"key":"apache-2.0","name":"Apache License 2.0","spdx_id":"Apache-2.0","url":"https://api.github.com/licenses/apache-2.0","node_id":"MDc6TGljZW5zZTI="},"forks":621,"open_issues":81,"watchers":2582,"default_branch":"master"}},"_links":{"self":{"href":"https://api.github.com/repos/dmlc/tvm/pulls/2338"},"html":{"href":"https://github.com/dmlc/tvm/pull/2338"},"issue":{"href":"https://api.github.com/repos/dmlc/tvm/issues/2338"},"comments":{"href":"https://api.github.com/repos/dmlc/tvm/issues/2338/comments"},"review_comments":{"href":"https://api.github.com/repos/dmlc/tvm/pulls/2338/comments"},"review_comment":{"href":"https://api.github.com/repos/dmlc/tvm/pulls/comments{/number}"},"commits":{"href":"https://api.github.com/repos/dmlc/tvm/pulls/2338/commits"},"statuses":{"href":"https://api.github.com/repos/dmlc/tvm/statuses/1d5a29526a58f2674c1fdaf80a99aa566aa86d03"}},"author_association":"CONTRIBUTOR","merged":true,"mergeable":null,"rebaseable":null,"mergeable_state":"unknown","merged_by":{"login":"tqchen","id":2577440,"node_id":"MDQ6VXNlcjI1Nzc0NDA=","avatar_url":"https://avatars3.githubusercontent.com/u/2577440?v=4","gravatar_id":"","url":"https://api.github.com/users/tqchen","html_url":"https://github.com/tqchen","followers_url":"https://api.github.com/users/tqchen/followers","following_url":"https://api.github.com/users/tqchen/following{/other_user}","gists_url":"https://api.github.com/users/tqchen/gists{/gist_id}","starred_url":"https://api.github.com/users/tqchen/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/tqchen/subscriptions","organizations_url":"https://api.github.com/users/tqchen/orgs","repos_url":"https://api.github.com/users/tqchen/repos","events_url":"https://api.github.com/users/tqchen/events{/privacy}","received_events_url":"https://api.github.com/users/tqchen/received_events","type":"User","site_admin":false},"comments":11,"review_comments":0,"maintainer_can_modify":false,"commits":1,"additions":47,"deletions":44,"changed_files":2}}
{"actor":{"display_login":"tqchen"}}