Merge pull request #26 from lucidrains/pre-hook-autotuner

use pre hook to fix in-place / autotuner issue

Author: lucidrains

lion_pytorch/lion_pytorch.py

@@ -1,4 +1,3 @@
-from functools import partial
 from typing import Tuple, Optional, Callable
 
 import torch
@@ -34,8 +33,7 @@ def __init__(
         lr: float = 1e-4,
         betas: Tuple[float, float] = (0.9, 0.99),
         weight_decay: float = 0.0,
-        use_triton: bool = False,
-        triton_block_size: int = 1024
+        use_triton: bool = False
     ):
         assert lr > 0.
         assert all([0. <= beta <= 1. for beta in betas])
@@ -49,12 +47,10 @@ def __init__(
         super().__init__(params, defaults)
 
         self.update_fn = update_fn
-        self.use_triton = use_triton
-        self.took_first_step = False
 
         if use_triton:
             from lion_pytorch.triton import update_fn as triton_update_fn
-            self.update_fn = partial(triton_update_fn, BLOCK_SIZE = triton_block_size)
+            self.update_fn = triton_update_fn
 
     @torch.no_grad()
     def step(
@@ -67,8 +63,6 @@ def step(
             with torch.enable_grad():
                 loss = closure()
 
-        # update all parameters
-
         for group in self.param_groups:
             for p in filter(lambda p: exists(p.grad), group['params']):
 

lion_pytorch/triton.py

@@ -1,5 +1,4 @@
 import torch
-from torch import Tensor
 
 try:
     import triton
@@ -8,18 +7,19 @@
     print('triton is not installed, please install by running `pip install triton -U --pre`')
     exit()
 
-# helper functions
+# clone param and exp_avg before autotuning takes place
+# as those are updated in-place
 
-def calc_num_warps(block_size):
-    num_warps = 4
-    if block_size >= 2048:
-        num_warps = 8
-    if block_size >= 4096:
-        num_warps = 16
-    return num_warps
+def clone_inplace_updated_params(nargs):
+    nargs['p_ptr'] = nargs['p_ptr'].clone()
+    nargs['exp_avg_ptr'] = nargs['exp_avg_ptr'].clone()
 
 # triton cuda kernel
 
+@triton.autotune(configs = [
+    triton.Config({'BLOCK_SIZE': 128}, num_warps = 4, pre_hook = clone_inplace_updated_params),
+    triton.Config({'BLOCK_SIZE': 1024}, num_warps = 8, pre_hook = clone_inplace_updated_params),
+], key = ['n_elements'])
 @triton.jit
 def update_fn_kernel(
     p_ptr,
@@ -80,35 +80,26 @@ def update_fn_kernel(
     tl.store(offset_exp_avg_ptr, exp_avg, mask = mask)
 
 def update_fn(
-    p: Tensor,
-    grad: Tensor,
-    exp_avg: Tensor,
+    p: torch.Tensor,
+    grad: torch.Tensor,
+    exp_avg: torch.Tensor,
     lr: float,
     wd: float,
     beta1: float,
-    beta2: float,
-    inplace: bool = True,
-    BLOCK_SIZE: int = 1024
+    beta2: float
 ):
     assert all([t.is_cuda for t in (p, grad, exp_avg)])
-
     n_elements = p.numel()
 
-    block_size = triton.next_power_of_2(BLOCK_SIZE)
-    num_warps = calc_num_warps(block_size)
-    n_rows = triton.cdiv(n_elements, block_size)
-
-    # call triton cuda kernel
+    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)    
 
-    update_fn_kernel[(n_rows,)](
+    update_fn_kernel[grid](
         p,
         grad,
         exp_avg,
         lr,
         wd,
         beta1,
         beta2,
-        n_elements,
-        num_warps = num_warps,
-        BLOCK_SIZE = BLOCK_SIZE
+        n_elements
     )

setup.py

@@ -3,7 +3,7 @@
 setup(
   name = 'lion-pytorch',
   packages = find_packages(exclude=[]),
-  version = '0.1.0',
+  version = '0.1.2',
   license='MIT',
   description = 'Lion Optimizer - Pytorch',
   author = 'Phil Wang',