1. 28 Oct, 2022 1 commit
  2. 27 Oct, 2022 2 commits
  3. 26 Oct, 2022 1 commit
  4. 24 Oct, 2022 1 commit
  5. 19 Oct, 2022 2 commits
  6. 18 Oct, 2022 1 commit
  7. 13 Oct, 2022 2 commits
  8. 04 Oct, 2022 2 commits
  9. 03 Oct, 2022 1 commit
    • Umang Yadav's avatar
      Add output_alias and runs_on_offload_target flags for the custom ops (#1309) · c9ffb38d
      Umang Yadav authored
      Adds two methods for the custom_ops virtual class.
      
      bool runs_on_offload_target(), if the custom op runs directly on the gpu then it should be set to true. in this case, custom op expects its parameters to reside in GPU memory and writes output to the GPU memory. If it is set to false then, custom op expects it's parameter to reside on the host and puts back the result into the host memory.
      
      output_alias, if output of the custom op is aliasing the input buffer. i.e. interpreting the same input buffer with differnet shape and strides.
      
      Update as_vector() in C++ API to handle non-standard shapes. It required exposing element_index to space_index conversion method for the shape class.
      c9ffb38d
  10. 29 Sep, 2022 1 commit
  11. 28 Sep, 2022 1 commit
    • Umang Yadav's avatar
      Add compute_fp32 flag for quant_gemm tests (#1360) · 70e63960
      Umang Yadav authored
      test_gpu_pack_int8_args fails on gfx908 machine, because it doesn't set compute_fp32 flag correctly. This PR fixes the test such that it checks for the device-name, and rocblas-versions and sets this flag accordingly.
      70e63960
  12. 27 Sep, 2022 1 commit
  13. 26 Sep, 2022 1 commit
  14. 23 Sep, 2022 1 commit
  15. 21 Sep, 2022 1 commit
  16. 19 Sep, 2022 1 commit
    • Paul Fultz II's avatar
      Improve layernorm and reductions performance (#1348) · 97a1ed2d
      Paul Fultz II authored
      Compute mean and variance in same reduction
      Set block size to numbers divisible by 32 instead powers of 2
      Global is also set exactly instead of being divisible by block size
      More exact matching of global/local can help get rid of branching/loops
      Reduce vectors first before doing dpp_reduce
      Explicitly vectorize array operators since the compiler doesnt always vectorize them
      Still uses old for loop when its computing at compile-time since the reinterpret_cast nor the all the vector types is supported
      97a1ed2d
  17. 16 Sep, 2022 1 commit
  18. 15 Sep, 2022 1 commit
  19. 14 Sep, 2022 1 commit
  20. 13 Sep, 2022 1 commit
    • turneram's avatar
      Use rocblas_gemm_ex for batched gemms with broadcasted B (#1354) · a10a8ef1
      turneram authored
      Improves performance for 4/6 GEMMs used by huggingface BERT models with batch_size>1 by using a non-batched rocBLAS call for GEMMs where the B input has a broadcasted batch dimension.
      The four verify tests added reflect the actual configurations used by bert-base-cased, with varied batch sizes.
      
      Also adds a matcher to simplify_reshapes to move multibroadcasts after concats.
      a10a8ef1
  21. 08 Sep, 2022 1 commit
  22. 07 Sep, 2022 1 commit
  23. 06 Sep, 2022 1 commit
  24. 31 Aug, 2022 1 commit
  25. 27 Aug, 2022 2 commits
  26. 21 Aug, 2022 1 commit
    • varunsh's avatar
      Update is_supported (#1334) · 79e15ca9
      varunsh authored
      * Update is_supported
      * Return object from is_supported
      * Return by reference in interator
      79e15ca9
  27. 19 Aug, 2022 1 commit
  28. 17 Aug, 2022 1 commit
  29. 16 Aug, 2022 1 commit
  30. 12 Aug, 2022 1 commit
  31. 04 Aug, 2022 1 commit
    • Charlie Lin's avatar
      Dynamic ref convolution op (#1224) · 67f77ac1
      Charlie Lin authored
      
      
      * Dynamic shape handling in shape object
      
      * rewrite empty lens multibroadcast test
      
      * Shape class changes to handle dynamic
      * More throw errors for functions that don't make sense for dynamic shape
      * Print output changes
      * Serialization changes
      
      * Fixing serialization errors
      
      * Remove const on dyn_dim copy getters
      
      * Dynamic shape tests
      
      * Fix serialize errors
      
      * Add dyn_data struct to avoid ambiguous constructor
      
      * Tidy fix: emplace_back() over for loop
      
      * Tidy fix: use move
      
      * Use std::initializer_list in constructor
      Reverts the dyn_data struct change
      Should get around the ambiguous braced initialization list error
      
      * avoid typedef
      
      * element_space, min,max,opt _lens change
      
      * formatting
      
      * Comments fix
      
      * dynamic bytes() test
      
      * Seralize and reflect changes
      
      * formatting
      
      * Test the dynamic lens functions
      
      * progress
      
      * Formatting
      
      * Dynamic conv draft progress
      
      * Add operator<< tests for coverage
      
      * Coverage update
      
      * Add to conv dynamic batch test
      
      * Dynamic image size test
      
      * Dynamic weight handling
      
      * Dyn image shape test change, fix dyn weight cond
      
      * Comment update
      
      * Dynamic weights shape test and fix
      
      * Use ternary operator
      
      * Tidy fixes
      
      * Handle dynamic graph input shapes in ONNX parser
      
      * Formatting
      
      * Handle dynamic shape for convolution
      
      * formatting
      
      * cppcheck fixes
      
      * Add onnx test files
      
      * Fix typo
      
      * Disable auto_pad for dynamic input shape
      
      * check_shapes object checks for allowing dynamic shapes
      
      * Fix any_of
      
      * Change to maintain const objectness
      
      * Formatting
      
      * Check shapes allow dynamic
      
      * Refactor compute_shape() call into op.compute()
      Allows for per operator differences with handling dynamic shape
      Fix operation.hpp change to use the generator
      
      * Comment fix
      
      * Refactor normalize_attributes() calls to use max_lens()
      
      * Comment addition
      
      * Update other normalize_attributes() calls
      
      * Change to using constructor and add tests
      
      * Use const member function
      
      * Add more dynamic shape support
      
      * Add tests for error code coverage
      
      * Fix opt shape bug and add shape tests
      
      * capture all by ref
      
      * Fix typo with img shape calculation
      
      * Add more tests
      
      * dynamic auto pad attempt
      Linker error with pad_calc.cpp
      
      * Fix parse dyn auto_pad
      Should only need to use dynamic auto pad when the image shape or kernel
      shape are dynamic. For a dynamic batch size, the auto pad calculation is
      the same.
      
      * Fix linking error
      
      * Fix auto_pad bug
      Fixed input tensor with auto_pad setting on
      
      * auto_pad onnx tests
      
      * Fix auto_pad calculation, evaluate in ref_conv
      add ref_ops tests
      
      * Add shape tests, fix bugs
      
      * Refactor first two output dynamic len calculation
      
      * Conv MLIR test update
      
      * i64 MLIR test fix
      
      * Fix MLIR test typo
      Co-authored-by: default avatarChris Austen <causten@users.noreply.github.com>
      67f77ac1
  32. 02 Aug, 2022 1 commit
  33. 29 Jul, 2022 1 commit
    • Umang Yadav's avatar
      Avoid registering host buffer ptr multiple times during hip copies (#1245) · 7596f3f1
      Umang Yadav authored
      Currently, while copying a host buffer to the device, it first registers/maps the host buffer pointer to address space of the device.
      
      If the host buffer has been allocated by the hipHostMalloc then, it is implicitly registered to the device's address space, and no need to register again. This PR adds a check for the same.
      7596f3f1
  34. 25 Jul, 2022 2 commits
    • Ted Themistokleous's avatar
      Add onnx mod operator (#1302) · 77e80b8e
      Ted Themistokleous authored
      * Add in changes for onnx Mod operator
      
      Initial operator for mod implementation and test cases for integer and floating based types.
      
      Need to use fmod from stdlib for floating point types. half_float::half thankfully is specced to the use the existing std::fmod() call when looking at the half.hpp implementation.
      
      fmod_flag should mirror the onnx fmod attribute. Right now using a floating point type without setting that on the user side to true will result in an exception.
      
      Ref ticket #1283 
      77e80b8e
    • varunsh's avatar
      Add fpga target (#1304) · 8a30d698
      varunsh authored
      * Add is_supported to the target
      * Add get_target_assignments
      * Rename assignment to target_assignments
      * Add ref target header to test
      * Add fpga target
      * Make context const in compute
      8a30d698