{"action":"created","comment":{"url":"https://api.github.com/repos/torch/cutorch/pulls/comments/27021518","id":27021518,"diff_hunk":"@@ -2270,3 +2189,127 @@ void THCudaTensor_indexSelect(THCState *state, THCudaTensor *res_, THCudaTensor\n   THCudaCheck(cudaFree(stride_));\n   THCudaTensor_free(state, indices_);\n }\n+\n+struct TensorMaskedFillOp {\n+  TensorMaskedFillOp(float v) : value(v) {}\n+  __device__ __forceinline__ void operator()(float* t, float* mask) {\n+    // Really mask should be `0` or `1` but we can't propagate errors here.\n+    const float maskVal = *mask;\n+    if (maskVal != 0.0f) {\n+      *t = value;\n+    }\n+  }\n+\n+  float value;\n+};\n+\n+void THCudaTensor_maskedFill(THCState* state,\n+                             THCudaTensor *tensor, THCudaTensor *mask, float value)\n+{\n+  THArgCheck(THCudaTensor_nElement(state, tensor) ==\n+             THCudaTensor_nElement(state, mask),\n+             2, \"sizes do not match\");\n+\n+  if (!THCudaTensor_pointwiseApply2(state, tensor, mask, TensorMaskedFillOp(value))) {\n+    THArgCheck(false, 2, CUTORCH_DIM_WARNING);\n+  }\n+}\n+\n+struct TensorMaskedCopyOp {\n+  __device__ __forceinline__ void operator()(float* out, float* mask, float* in) {\n+    // Really mask should be `0` or `1` but we can't propagate errors here.\n+    if (*mask != 0.0f) {\n+      *out = *in;\n+    }\n+  }\n+};\n+\n+void THCudaTensor_maskedCopy(THCState* state,\n+                             THCudaTensor *tensor, THCudaTensor *mask, THCudaTensor *src)\n+{\n+  THArgCheck(THCudaTensor_nElement(state, mask) ==\n+             THCudaTensor_nElement(state, src),\n+             2, \"sizes do not match\");\n+\n+  THCudaTensor_resizeAs(state, tensor, src);\n+\n+  if (!THCudaTensor_pointwiseApply3(state, tensor, mask, src, TensorMaskedCopyOp())) {\n+    THArgCheck(false, 2, CUTORCH_DIM_WARNING);\n+  }\n+}\n+\n+struct TensorMaskedSelectOp {\n+  TensorMaskedSelectOp(float* t) : out(t) {}\n+  __device__ __forceinline__ void operator()(float* mask, float* maskPrefixSum, float* in) {\n+    // Really mask should be `0` or `1` but we can't propagate errors here.\n+    if (*mask != 0.0f) {\n+      out[(int) *maskPrefixSum] = *in;\n+    }\n+  }\n+\n+  float* out;\n+};\n+\n+void THCudaTensor_maskedSelect(THCState* state,\n+                               THCudaTensor *tensor, THCudaTensor *src, THCudaTensor *mask)\n+{\n+  THArgCheck(THCudaTensor_nElement(state, mask) == THCudaTensor_nElement(state, src),\n+             2, \"sizes do not match\");\n+\n+  // Determine our output size\n+  THCudaTensor* contigMask = THCudaTensor_newContiguous(state, mask);\n+  int totalElements = (int) THCudaTensor_sumall(state, contigMask);\n+  THCudaTensor_resize1d(state, tensor, totalElements);\n+\n+  // Use a prefix sum to determine the output locations of the masked elements\n+  THCudaTensor* maskPrefixSum = THCudaTensor_new(state);\n+  THCudaTensor_resizeAs(state, maskPrefixSum, mask);\n+\n+  thrust::device_ptr<float>\n+    maskData(THCudaTensor_data(state, contigMask));\n+  thrust::device_ptr<float>\n+    maskPrefixSumData(THCudaTensor_data(state, maskPrefixSum));\n+  thrust::exclusive_scan(maskData,\n+                         maskData + THCudaTensor_nElement(state, contigMask),\n+                         maskPrefixSumData);\n+\n+  // Then copy over the masked elements at their desired output index\n+  if (!THCudaTensor_pointwiseApply3(\n+        state, contigMask, maskPrefixSum,\n+        src, TensorMaskedSelectOp(THCudaTensor_data(state, tensor)))) {\n+    THArgCheck(false, 2, CUTORCH_DIM_WARNING);\n+  }\n+\n+  THCudaTensor_free(state, contigMask);\n+  THCudaTensor_free(state, maskPrefixSum);\n+}\n+\n+void THCudaTensor_maskedFillByte(THCState* state, THCudaTensor *tensor, THByteTensor *mask, float value)\n+{\n+  THLongStorage* maskSize = THByteTensor_newSizeOf(mask);\n+  THCudaTensor* maskCuda = THCudaTensor_newWithSize(state, maskSize, NULL);\n+  THLongStorage_free(maskSize);\n+","path":"lib/THC/THCTensorMath.cu","position":1600,"original_position":1600,"commit_id":"00460c067020379da7c55383839fac95d425a00a","original_commit_id":"00460c067020379da7c55383839fac95d425a00a","user":{"login":"dominikgrewe","id":1619348,"avatar_url":"https://avatars.githubusercontent.com/u/1619348?v=3","gravatar_id":"","url":"https://api.github.com/users/dominikgrewe","html_url":"https://github.com/dominikgrewe","followers_url":"https://api.github.com/users/dominikgrewe/followers","following_url":"https://api.github.com/users/dominikgrewe/following{/other_user}","gists_url":"https://api.github.com/users/dominikgrewe/gists{/gist_id}","starred_url":"https://api.github.com/users/dominikgrewe/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/dominikgrewe/subscriptions","organizations_url":"https://api.github.com/users/dominikgrewe/orgs","repos_url":"https://api.github.com/users/dominikgrewe/repos","events_url":"https://api.github.com/users/dominikgrewe/events{/privacy}","received_events_url":"https://api.github.com/users/dominikgrewe/received_events","type":"User","site_admin":false},"body":"You need to copy `mask` into `maskCuda`, otherwise it will be uninitialized.\r\nSame in the two functions below.","created_at":"2015-03-24T12:12:06Z","updated_at":"2015-03-24T12:12:06Z","html_url":"https://github.com/torch/cutorch/pull/120#discussion_r27021518","pull_request_url":"https://api.github.com/repos/torch/cutorch/pulls/120","_links":{"self":{"href":"https://api.github.com/repos/torch/cutorch/pulls/comments/27021518"},"html":{"href":"https://github.com/torch/cutorch/pull/120#discussion_r27021518"},"pull_request":{"href":"https://api.github.com/repos/torch/cutorch/pulls/120"}}},"pull_request":{"url":"https://api.github.com/repos/torch/cutorch/pulls/120","id":31579895,"html_url":"https://github.com/torch/cutorch/pull/120","diff_url":"https://github.com/torch/cutorch/pull/120.diff","patch_url":"https://github.com/torch/cutorch/pull/120.patch","issue_url":"https://api.github.com/repos/torch/cutorch/issues/120","number":120,"state":"open","locked":false,"title":"adds Sort, revamps TensorMath, adds masked* operations","user":{"login":"soumith","id":1310570,"avatar_url":"https://avatars.githubusercontent.com/u/1310570?v=3","gravatar_id":"","url":"https://api.github.com/users/soumith","html_url":"https://github.com/soumith","followers_url":"https://api.github.com/users/soumith/followers","following_url":"https://api.github.com/users/soumith/following{/other_user}","gists_url":"https://api.github.com/users/soumith/gists{/gist_id}","starred_url":"https://api.github.com/users/soumith/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/soumith/subscriptions","organizations_url":"https://api.github.com/users/soumith/orgs","repos_url":"https://api.github.com/users/soumith/repos","events_url":"https://api.github.com/users/soumith/events{/privacy}","received_events_url":"https://api.github.com/users/soumith/received_events","type":"User","site_admin":false},"body":"Implements for #70 :\r\n* maskedSelect\r\n* maskedCopy\r\n* maskedFill\r\n* sort\r\n\r\nAdds a TensorApply and TensorReduce kernels that work on both contiguous and non-contiguous tensors, and maintains equal performance for contiguous cases.\r\n\r\nThis PR makes performance with non-contiguous kernels a huge huge improvement.\r\n\r\nImplemented completely by @wickedfoo\r\n\r\nHaving some build issues, looking into it","created_at":"2015-03-20T00:42:29Z","updated_at":"2015-03-24T12:12:06Z","closed_at":null,"merged_at":null,"merge_commit_sha":"d73e6fb47383a7f3cfe2d8c3c5a3ffe1abd5addc","assignee":null,"milestone":null,"commits_url":"https://api.github.com/repos/torch/cutorch/pulls/120/commits","review_comments_url":"https://api.github.com/repos/torch/cutorch/pulls/120/comments","review_comment_url":"https://api.github.com/repos/torch/cutorch/pulls/comments{/number}","comments_url":"https://api.github.com/repos/torch/cutorch/issues/120/comments","statuses_url":"https://api.github.com/repos/torch/cutorch/statuses/00460c067020379da7c55383839fac95d425a00a","head":{"label":"torch:goodies","ref":"goodies","sha":"00460c067020379da7c55383839fac95d425a00a","user":{"login":"torch","id":4978924,"avatar_url":"https://avatars.githubusercontent.com/u/4978924?v=3","gravatar_id":"","url":"https://api.github.com/users/torch","html_url":"https://github.com/torch","followers_url":"https://api.github.com/users/torch/followers","following_url":"https://api.github.com/users/torch/following{/other_user}","gists_url":"https://api.github.com/users/torch/gists{/gist_id}","starred_url":"https://api.github.com/users/torch/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/torch/subscriptions","organizations_url":"https://api.github.com/users/torch/orgs","repos_url":"https://api.github.com/users/torch/repos","events_url":"https://api.github.com/users/torch/events{/privacy}","received_events_url":"https://api.github.com/users/torch/received_events","type":"Organization","site_admin":false},"repo":{"id":13677478,"name":"cutorch","full_name":"torch/cutorch","owner":{"login":"torch","id":4978924,"avatar_url":"https://avatars.githubusercontent.com/u/4978924?v=3","gravatar_id":"","url":"https://api.github.com/users/torch","html_url":"https://github.com/torch","followers_url":"https://api.github.com/users/torch/followers","following_url":"https://api.github.com/users/torch/following{/other_user}","gists_url":"https://api.github.com/users/torch/gists{/gist_id}","starred_url":"https://api.github.com/users/torch/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/torch/subscriptions","organizations_url":"https://api.github.com/users/torch/orgs","repos_url":"https://api.github.com/users/torch/repos","events_url":"https://api.github.com/users/torch/events{/privacy}","received_events_url":"https://api.github.com/users/torch/received_events","type":"Organization","site_admin":false},"private":false,"html_url":"https://github.com/torch/cutorch","description":"A CUDA backend for Torch7","fork":false,"url":"https://api.github.com/repos/torch/cutorch","forks_url":"https://api.github.com/repos/torch/cutorch/forks","keys_url":"https://api.github.com/repos/torch/cutorch/keys{/key_id}","collaborators_url":"https://api.github.com/repos/torch/cutorch/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/torch/cutorch/teams","hooks_url":"https://api.github.com/repos/torch/cutorch/hooks","issue_events_url":"https://api.github.com/repos/torch/cutorch/issues/events{/number}","events_url":"https://api.github.com/repos/torch/cutorch/events","assignees_url":"https://api.github.com/repos/torch/cutorch/assignees{/user}","branches_url":"https://api.github.com/repos/torch/cutorch/branches{/branch}","tags_url":"https://api.github.com/repos/torch/cutorch/tags","blobs_url":"https://api.github.com/repos/torch/cutorch/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/torch/cutorch/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/torch/cutorch/git/refs{/sha}","trees_url":"https://api.github.com/repos/torch/cutorch/git/trees{/sha}","statuses_url":"https://api.github.com/repos/torch/cutorch/statuses/{sha}","languages_url":"https://api.github.com/repos/torch/cutorch/languages","stargazers_url":"https://api.github.com/repos/torch/cutorch/stargazers","contributors_url":"https://api.github.com/repos/torch/cutorch/contributors","subscribers_url":"https://api.github.com/repos/torch/cutorch/subscribers","subscription_url":"https://api.github.com/repos/torch/cutorch/subscription","commits_url":"https://api.github.com/repos/torch/cutorch/commits{/sha}","git_commits_url":"https://api.github.com/repos/torch/cutorch/git/commits{/sha}","comments_url":"https://api.github.com/repos/torch/cutorch/comments{/number}","issue_comment_url":"https://api.github.com/repos/torch/cutorch/issues/comments{/number}","contents_url":"https://api.github.com/repos/torch/cutorch/contents/{+path}","compare_url":"https://api.github.com/repos/torch/cutorch/compare/{base}...{head}","merges_url":"https://api.github.com/repos/torch/cutorch/merges","archive_url":"https://api.github.com/repos/torch/cutorch/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/torch/cutorch/downloads","issues_url":"https://api.github.com/repos/torch/cutorch/issues{/number}","pulls_url":"https://api.github.com/repos/torch/cutorch/pulls{/number}","milestones_url":"https://api.github.com/repos/torch/cutorch/milestones{/number}","notifications_url":"https://api.github.com/repos/torch/cutorch/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/torch/cutorch/labels{/name}","releases_url":"https://api.github.com/repos/torch/cutorch/releases{/id}","created_at":"2013-10-18T12:29:23Z","updated_at":"2015-03-19T16:40:42Z","pushed_at":"2015-03-20T00:42:29Z","git_url":"git://github.com/torch/cutorch.git","ssh_url":"git@github.com:torch/cutorch.git","clone_url":"https://github.com/torch/cutorch.git","svn_url":"https://github.com/torch/cutorch","homepage":"","size":1633,"stargazers_count":32,"watchers_count":32,"language":"Cuda","has_issues":true,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":29,"mirror_url":null,"open_issues_count":8,"forks":29,"open_issues":8,"watchers":32,"default_branch":"master"}},"base":{"label":"torch:master","ref":"master","sha":"37d636d6fcf22f78fc87a7b887127c2742bb0aac","user":{"login":"torch","id":4978924,"avatar_url":"https://avatars.githubusercontent.com/u/4978924?v=3","gravatar_id":"","url":"https://api.github.com/users/torch","html_url":"https://github.com/torch","followers_url":"https://api.github.com/users/torch/followers","following_url":"https://api.github.com/users/torch/following{/other_user}","gists_url":"https://api.github.com/users/torch/gists{/gist_id}","starred_url":"https://api.github.com/users/torch/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/torch/subscriptions","organizations_url":"https://api.github.com/users/torch/orgs","repos_url":"https://api.github.com/users/torch/repos","events_url":"https://api.github.com/users/torch/events{/privacy}","received_events_url":"https://api.github.com/users/torch/received_events","type":"Organization","site_admin":false},"repo":{"id":13677478,"name":"cutorch","full_name":"torch/cutorch","owner":{"login":"torch","id":4978924,"avatar_url":"https://avatars.githubusercontent.com/u/4978924?v=3","gravatar_id":"","url":"https://api.github.com/users/torch","html_url":"https://github.com/torch","followers_url":"https://api.github.com/users/torch/followers","following_url":"https://api.github.com/users/torch/following{/other_user}","gists_url":"https://api.github.com/users/torch/gists{/gist_id}","starred_url":"https://api.github.com/users/torch/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/torch/subscriptions","organizations_url":"https://api.github.com/users/torch/orgs","repos_url":"https://api.github.com/users/torch/repos","events_url":"https://api.github.com/users/torch/events{/privacy}","received_events_url":"https://api.github.com/users/torch/received_events","type":"Organization","site_admin":false},"private":false,"html_url":"https://github.com/torch/cutorch","description":"A CUDA backend for Torch7","fork":false,"url":"https://api.github.com/repos/torch/cutorch","forks_url":"https://api.github.com/repos/torch/cutorch/forks","keys_url":"https://api.github.com/repos/torch/cutorch/keys{/key_id}","collaborators_url":"https://api.github.com/repos/torch/cutorch/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/torch/cutorch/teams","hooks_url":"https://api.github.com/repos/torch/cutorch/hooks","issue_events_url":"https://api.github.com/repos/torch/cutorch/issues/events{/number}","events_url":"https://api.github.com/repos/torch/cutorch/events","assignees_url":"https://api.github.com/repos/torch/cutorch/assignees{/user}","branches_url":"https://api.github.com/repos/torch/cutorch/branches{/branch}","tags_url":"https://api.github.com/repos/torch/cutorch/tags","blobs_url":"https://api.github.com/repos/torch/cutorch/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/torch/cutorch/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/torch/cutorch/git/refs{/sha}","trees_url":"https://api.github.com/repos/torch/cutorch/git/trees{/sha}","statuses_url":"https://api.github.com/repos/torch/cutorch/statuses/{sha}","languages_url":"https://api.github.com/repos/torch/cutorch/languages","stargazers_url":"https://api.github.com/repos/torch/cutorch/stargazers","contributors_url":"https://api.github.com/repos/torch/cutorch/contributors","subscribers_url":"https://api.github.com/repos/torch/cutorch/subscribers","subscription_url":"https://api.github.com/repos/torch/cutorch/subscription","commits_url":"https://api.github.com/repos/torch/cutorch/commits{/sha}","git_commits_url":"https://api.github.com/repos/torch/cutorch/git/commits{/sha}","comments_url":"https://api.github.com/repos/torch/cutorch/comments{/number}","issue_comment_url":"https://api.github.com/repos/torch/cutorch/issues/comments{/number}","contents_url":"https://api.github.com/repos/torch/cutorch/contents/{+path}","compare_url":"https://api.github.com/repos/torch/cutorch/compare/{base}...{head}","merges_url":"https://api.github.com/repos/torch/cutorch/merges","archive_url":"https://api.github.com/repos/torch/cutorch/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/torch/cutorch/downloads","issues_url":"https://api.github.com/repos/torch/cutorch/issues{/number}","pulls_url":"https://api.github.com/repos/torch/cutorch/pulls{/number}","milestones_url":"https://api.github.com/repos/torch/cutorch/milestones{/number}","notifications_url":"https://api.github.com/repos/torch/cutorch/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/torch/cutorch/labels{/name}","releases_url":"https://api.github.com/repos/torch/cutorch/releases{/id}","created_at":"2013-10-18T12:29:23Z","updated_at":"2015-03-19T16:40:42Z","pushed_at":"2015-03-20T00:42:29Z","git_url":"git://github.com/torch/cutorch.git","ssh_url":"git@github.com:torch/cutorch.git","clone_url":"https://github.com/torch/cutorch.git","svn_url":"https://github.com/torch/cutorch","homepage":"","size":1633,"stargazers_count":32,"watchers_count":32,"language":"Cuda","has_issues":true,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":29,"mirror_url":null,"open_issues_count":8,"forks":29,"open_issues":8,"watchers":32,"default_branch":"master"}},"_links":{"self":{"href":"https://api.github.com/repos/torch/cutorch/pulls/120"},"html":{"href":"https://github.com/torch/cutorch/pull/120"},"issue":{"href":"https://api.github.com/repos/torch/cutorch/issues/120"},"comments":{"href":"https://api.github.com/repos/torch/cutorch/issues/120/comments"},"review_comments":{"href":"https://api.github.com/repos/torch/cutorch/pulls/120/comments"},"review_comment":{"href":"https://api.github.com/repos/torch/cutorch/pulls/comments{/number}"},"commits":{"href":"https://api.github.com/repos/torch/cutorch/pulls/120/commits"},"statuses":{"href":"https://api.github.com/repos/torch/cutorch/statuses/00460c067020379da7c55383839fac95d425a00a"}}}}