{"action":"created","issue":{"url":"https://api.github.com/repos/pytorch/pytorch/issues/9628","repository_url":"https://api.github.com/repos/pytorch/pytorch","labels_url":"https://api.github.com/repos/pytorch/pytorch/issues/9628/labels{/name}","comments_url":"https://api.github.com/repos/pytorch/pytorch/issues/9628/comments","events_url":"https://api.github.com/repos/pytorch/pytorch/issues/9628/events","html_url":"https://github.com/pytorch/pytorch/pull/9628","id":343136202,"node_id":"MDExOlB1bGxSZXF1ZXN0MjAyODcxNDk2","number":9628,"title":"Add CTC loss","user":{"login":"t-vi","id":20787943,"node_id":"MDQ6VXNlcjIwNzg3OTQz","avatar_url":"https://avatars2.githubusercontent.com/u/20787943?v=4","gravatar_id":"","url":"https://api.github.com/users/t-vi","html_url":"https://github.com/t-vi","followers_url":"https://api.github.com/users/t-vi/followers","following_url":"https://api.github.com/users/t-vi/following{/other_user}","gists_url":"https://api.github.com/users/t-vi/gists{/gist_id}","starred_url":"https://api.github.com/users/t-vi/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/t-vi/subscriptions","organizations_url":"https://api.github.com/users/t-vi/orgs","repos_url":"https://api.github.com/users/t-vi/repos","events_url":"https://api.github.com/users/t-vi/events{/privacy}","received_events_url":"https://api.github.com/users/t-vi/received_events","type":"User","site_admin":false},"labels":[],"state":"closed","locked":false,"assignee":{"login":"SsnL","id":5674597,"node_id":"MDQ6VXNlcjU2NzQ1OTc=","avatar_url":"https://avatars2.githubusercontent.com/u/5674597?v=4","gravatar_id":"","url":"https://api.github.com/users/SsnL","html_url":"https://github.com/SsnL","followers_url":"https://api.github.com/users/SsnL/followers","following_url":"https://api.github.com/users/SsnL/following{/other_user}","gists_url":"https://api.github.com/users/SsnL/gists{/gist_id}","starred_url":"https://api.github.com/users/SsnL/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/SsnL/subscriptions","organizations_url":"https://api.github.com/users/SsnL/orgs","repos_url":"https://api.github.com/users/SsnL/repos","events_url":"https://api.github.com/users/SsnL/events{/privacy}","received_events_url":"https://api.github.com/users/SsnL/received_events","type":"User","site_admin":false},"assignees":[{"login":"SsnL","id":5674597,"node_id":"MDQ6VXNlcjU2NzQ1OTc=","avatar_url":"https://avatars2.githubusercontent.com/u/5674597?v=4","gravatar_id":"","url":"https://api.github.com/users/SsnL","html_url":"https://github.com/SsnL","followers_url":"https://api.github.com/users/SsnL/followers","following_url":"https://api.github.com/users/SsnL/following{/other_user}","gists_url":"https://api.github.com/users/SsnL/gists{/gist_id}","starred_url":"https://api.github.com/users/SsnL/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/SsnL/subscriptions","organizations_url":"https://api.github.com/users/SsnL/orgs","repos_url":"https://api.github.com/users/SsnL/repos","events_url":"https://api.github.com/users/SsnL/events{/privacy}","received_events_url":"https://api.github.com/users/SsnL/received_events","type":"User","site_admin":false}],"milestone":null,"comments":12,"created_at":"2018-07-20T14:47:38Z","updated_at":"2018-07-31T23:11:10Z","closed_at":"2018-07-31T18:10:56Z","author_association":"CONTRIBUTOR","pull_request":{"url":"https://api.github.com/repos/pytorch/pytorch/pulls/9628","html_url":"https://github.com/pytorch/pytorch/pull/9628","diff_url":"https://github.com/pytorch/pytorch/pull/9628.diff","patch_url":"https://github.com/pytorch/pytorch/pull/9628.patch"},"body":"The CPU and CUDA variants are a direct transposition of Graves et al.'s description of the algorithm with the\r\nmodification that is is in log space.\r\nThe there also is a binding for the (much faster) CuDNN implementation.\r\n\r\nThis could eventually fix #3420 \r\n\r\nI still need to add tests (TestNN seems much more elaborate than the other testing) and fix the bugs than invariably turn up during the testing. Also, I want to add some more code comments.\r\n\r\nI could use feedback on all sorts of things, including:\r\n- Type handling (cuda vs. cpu for the int tensors, dtype for the int tensors)\r\n- Input convention. I use log probs because that is what the gradients are for.\r\n- Launch parameters for the kernels\r\n- Errors and obmissions and anything else I'm not even aware of.\r\n\r\nThank you for looking!\r\n\r\nIn terms of performance it looks like it is superficially comparable to WarpCTC (and thus, but I have not systematically investigated this).\r\nI have read CuDNN is much faster than implementations because it does *not* use log-space, but also the gathering step is much much faster (but I avoided trying tricky things, it seems to contribute to warpctc's fragility). I might think some more which existing torch function (scatter or index..) I could learn from for that step.\r\nAverage timings for the kernels from nvprof for some size:\r\n\r\n```\r\nCuDNN:\r\n60.464us compute_alphas_and_betas\r\n16.755us compute_grads_deterministic\r\nCuda:\r\n121.06us ctc_loss_backward_collect_gpu_kernel (= grads)\r\n109.88us ctc_loss_gpu_kernel (= alphas)\r\n98.517us ctc_loss_backward_betas_gpu_kernel (= betas)\r\nWarpCTC:\r\n299.74us compute_betas_and_grad_kernel\r\n66.977us compute_alpha_kernel\r\n```\r\n\r\nOf course, I still have the (silly) outer blocks loop rather than computing consecutive `s` in each thread which I might change, and there are a few other things where one could look for better implementations.\r\n\r\nFinally, it might not be unreasonable to start with these implementations, as the performance of the loss has to be seen in the context of the entire training computation, so this would likely dilute the relative speedup considerably.\r\n\r\nMy performance measuring testing script:\r\n```\r\nimport timeit\r\nimport sys\r\nimport torch\r\nnum_labels = 10\r\ntarget_length  = 30\r\ninput_length = 50\r\neps = 1e-5\r\nBLANK = 0#num_labels\r\nbatch_size = 16\r\n\r\n\r\ntorch.manual_seed(5)\r\nactivations = torch.randn(input_length, batch_size, num_labels + 1)\r\nlog_probs = torch.log_softmax(activations, 2)\r\nprobs = torch.exp(log_probs)\r\ntargets = torch.randint(1, num_labels+1, (batch_size * target_length,), dtype=torch.long)\r\ntargets_2d = targets.view(batch_size, target_length)\r\ntarget_lengths = torch.tensor(batch_size*[target_length])\r\ninput_lengths = torch.tensor(batch_size*[input_length])\r\nactivations = log_probs.detach()\r\n\r\n\r\ndef time_cuda_ctc_loss(grout, *args):\r\n    torch.cuda.synchronize()\r\n    culo, culog_alpha = torch._ctc_loss(*args)\r\n    g, = torch.autograd.grad(culo, args[0], grout)\r\n    torch.cuda.synchronize()\r\n\r\ndef time_cudnn_ctc_loss(groupt, *args):\r\n    torch.cuda.synchronize()\r\n    culo, cugra= torch._cudnn_ctc_loss(*args)\r\n    g, = torch.autograd.grad(culo, args[0], grout)\r\n    torch.cuda.synchronize()\r\n\r\ndef time_warp_ctc_loss(grout, *args):\r\n    torch.cuda.synchronize()\r\n    culo = warpctc.ctc_loss(*args, blank_label=BLANK, size_average=False, length_average=False, reduce=False)\r\n    g, = torch.autograd.grad(culo, args[0], grout)\r\n    torch.cuda.synchronize()\r\n\r\nif sys.argv[1] == 'cuda':\r\n    lpcu = log_probs.float().cuda().detach().requires_grad_()\r\n    args = [lpcu, targets_2d.cuda(), input_lengths.cuda(), target_lengths.cuda(), BLANK]\r\n    grout = lpcu.new_ones((batch_size,))\r\n    torch.cuda.synchronize()\r\n    print(timeit.repeat(\"time_cuda_ctc_loss(grout, *args)\", number=1000, globals=globals()))\r\nelif sys.argv[1] == 'cudnn':\r\n    lpcu = log_probs.float().cuda().detach().requires_grad_()\r\n    args = [lpcu, targets.int(), input_lengths.int(), target_lengths.int(), BLANK, True]\r\n    grout = lpcu.new_ones((batch_size,))\r\n    torch.cuda.synchronize()\r\n    print(timeit.repeat(\"time_cudnn_ctc_loss(grout, *args)\", number=1000, globals=globals()))\r\nelif sys.argv[1] == 'warpctc':\r\n    import warpctc\r\n    activations = activations.cuda().detach().requires_grad_()\r\n    args = [activations, input_lengths.int(), targets.int(), target_lengths.int()]\r\n    grout = activations.new_ones((batch_size,), device='cpu')\r\n    torch.cuda.synchronize()\r\n\r\n    print(timeit.repeat(\"time_warp_ctc_loss(grout, *args)\", number=1000, globals=globals()))\r\n```\r\nI'll also link to a notebook that I used for writing up the algorithm in simple form and then test the against implementations against it.\r\n"},"comment":{"url":"https://api.github.com/repos/pytorch/pytorch/issues/comments/409397223","html_url":"https://github.com/pytorch/pytorch/pull/9628#issuecomment-409397223","issue_url":"https://api.github.com/repos/pytorch/pytorch/issues/9628","id":409397223,"node_id":"MDEyOklzc3VlQ29tbWVudDQwOTM5NzIyMw==","user":{"login":"soumith","id":1310570,"node_id":"MDQ6VXNlcjEzMTA1NzA=","avatar_url":"https://avatars0.githubusercontent.com/u/1310570?v=4","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},"created_at":"2018-07-31T23:11:09Z","updated_at":"2018-07-31T23:11:09Z","author_association":"MEMBER","body":"as reported by someone internally, the example is broken.\r\n\r\n```\r\nTypeError: ctc_loss(): argument 'input_lengths' (position 3) must be tuple of ints, not Tensor\r\n```\r\n\r\nCan we allow input_lengths and target_lengths to optionally also be a Tensor, not just tuple of ints."}}
{"actor":{"display_login":"soumith"}}