Commit grafiek

1685 Commits

Auteur SHA1 Bericht Datum
Marko Arandjelovic 152f343124 SWDEV-472345 - Fix coalesced group size
In case when the tile size is greater than the number of active threads,
the coalesced group size should be equal to the number of active threads.

Change-Id: I1d41322f2428a07862a590cb5d34b01243383b7c
2024-07-12 04:29:53 -04:00
sdashmiz 57e79802cd SWDEV-421021 - Add APIs cuMemcpyNodeGet/Set params
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I24bc0da56aad34c9d5876a3d83b59515f11dc3ea
2024-07-10 12:53:01 -04:00
Jatin Chaudhary ecd812b2d8 SWDEV-460834 - add unsafe atomic add for fp16 and bf16
Change-Id: I6de5c2c425c9f8ac7f6c4e5c83c8b8b7ac8fe4cb
2024-07-02 16:45:17 -04:00
Jatin Chaudhary 36d58c6518 SWDEV-450109 - add missing coop group operations
any, all, ballot, match_any, match_all for coalesced_group and
tiled_group. add binary_partition.

Change-Id: I8ed0e654f36d788653352a0617beb5b8d44749bc
2024-06-20 10:42:06 -04:00
Ioannis Assiouras 8f42ad6aa3 SWDEV-464648 - code and comment cleanups
Change-Id: I5ba3f1bff500b3cd5903c2f441017735e688f83f
2024-06-07 22:38:09 +01:00
Fabian Ritter 819e537dc5 Add __device__ specifier to HIP's warpSize constexpr variable
The warpSize variable is set to the value of the __AMDGCN_WAVEFRONT_SIZE macro,
which is a meaningless default in host code.
The resolution for SWDEV-449015 will introduce diagnostics for uses of this
macro in host code, which includes the current definition of the warpSize
variable. With the __device__ specifier, the definition of the warpSize
variable will not cause these diagnostics.

This change does not stop the variable from being used in host code since clang
intentionally does not diagnose uses of __device__ constexpr variables in host
code.

Change-Id: I0317217affe94fdf2dfd9ad0f134e68f5173245f
2024-06-04 09:32:52 -04:00
Saleel Kudchadker f18a3405d7 SWDEV-301667 - Add enum for TASK command
Change-Id: I5d2712d913bdbc42bc6eaccf9eb2f8cb9976b22e
2024-05-28 14:18:46 -04:00
Jatin Chaudhary 2b290953b0 SWDEV-462027 - bfloat162 struct members are public
bfloat16 member is protected, bfloat162 members are public

Change-Id: Ice090401553b9526cb34d003b3996b85d538bc4e
2024-05-28 06:31:10 +00:00
Ioannis Assiouras 7b46701142 SWDEV-449620 - Added hipExternalMemoryGetMappedMipmappedArray in the dispatch table
Change-Id: I79952f131dea1949296685ee92cae05a4f9aec30
2024-05-28 06:31:10 +00:00
shadi e705e5e0d9 SWDEV-421027 - Add more Graph APIs
Signed-off-by: shadi <shadi.dashmiz@amd.com>
Change-Id: I0a1fc284e48317a49ca88d4ed4e3a10e752efd58
2024-05-28 06:28:17 +00:00
Anusha GodavarthySurya 45dccd6eab SWDEV-453739 - Added definition for undeclared methods
Change-Id: I3f10e0a3fb4bbaf9d873a6a988847f05ac43a312
2024-05-28 06:28:17 +00:00
shadi f2b01782ac SWDEV-420016 - Add more driver side graph APIs
Signed-off-by: shadi <shadi.dashmiz@amd.com>
Change-Id: Iff3ee7dcbcd24836f227fdc9bd5ff4b554ac914f
2024-04-25 12:50:43 -04:00
Ioannis Assiouras 2841aab017 SWDEV-451099 - Added include for __half type definitions for non-HIP code on windows
Change-Id: Id80cef5a36db8707276de052cbaf73b6826d222f
2024-04-24 15:31:31 -04:00
Jatin Chaudhary d7b0d78fad SWDEV-379007 - fix bool check for fp8_fnuz
for fnuz nums zero val is 0x00, -0 i.e. 0x80 would be a NaN

Change-Id: Ibdc4fb4b9fb307b5952434f08d45a8ddd6262db8
2024-04-17 05:31:21 -04:00
Sameer Sahasrabuddhe 03562a2547 SWDEV-454959 - ignore the upper half of the mask in wave32 mode
Change-Id: If027dd8cbe5cbe142fff353cb72c16f08e9aea8e
2024-04-17 10:12:57 +05:30
Jatin Chaudhary 49349f168c SWDEV-379007 - use avx instruction for bf16 cvt
AMD CPUs have had avx512_bf16 support for quite some time now (from
consumer Ryzen 7000 series to enterprise grade CPUs). This
patch should allow users to use the hardware bf16 unit when running the
__host__ variants of the function. This can be enabled via `hipcc ...
-mavx512vl -mavx512bf16`.

Change-Id: I67c377afc95ddfe8d45a048dce078a247d4a1878
2024-04-16 18:35:08 -04:00
Jatin Chaudhary ca07f59fb1 SWDEV-379007 initial implementation of fp8 header
Change-Id: Id9a5a85641882961e4d860a815217c641e6f3387
2024-04-16 05:37:59 -04:00
Jatin Chaudhary 481912a1fd SWDEV-379007 - add __hip_bfloat16_raw types
This also brings bfloat16 implementation closer to CUDA's.

Change-Id: I23f381141faacd6537923ae9b88ada4d661db496
2024-04-09 05:32:13 -04:00
kjayapra-amd 5b28e386f6 SWDEV-448288 - Remove references to deprecated llvm references.
Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72
2024-03-26 10:23:18 -04:00
Satyanvesh Dittakavi 684fd60c8f SWDEV-444021 - Implement hipGetFuncBySymbol
Change-Id: I7ef13d02c5b5c6ce2386ccb92b5602d005b35988
2024-03-14 02:56:30 -04:00
Sameer Sahasrabuddhe 0421a35f71 SWDEV-295298: hide warp sync builtins in ROCm 6.2 (absent in 6.1)
Change-Id: I6d59e4f3a133dce5ea0521fc9d09f93f9fd0f376
2024-03-11 02:44:20 -04:00
Anusha GodavarthySurya 4feb1f9337 SWDEV-448586 - Added implementation for new API hipStreamBeginCaptureToGraph
Change-Id: I1ce802102cef2b66c92d3375f769983841de793f
2024-03-07 05:24:49 +00:00
Rahul Manocha e5e7687fd8 SWDEV-448281 - Fix for CG Tile Partition Dynamic Size
Change-Id: Ibdccec45470caf55a8c00194e6da0dde73e4b180
2024-03-04 13:00:24 -05:00
Ioannis Assiouras 74cae705ae SWDEV-445960 - Fix __umul64hi __mul64hi on windows
Change-Id: I4857d269bc80fcc0e6b77df941f24ff1964a4d42
2024-03-03 12:42:42 -05:00
Jonathan R. Madsen d84c5ae3af SWDEV-443227 - Update hip_api_trace.cpp
- enforce incrementing the table versioning number when a table size changes outside of ifdef for ROCPROFILER_REGISTER
- add new HIP_ENFORCE_ABI entries
- update the HipDispatchTable size and bump HIP_RUNTIME_API_TABLE_STEP_VERSION to 1
- re-enable rocprofiler-register

Change-Id: Ie0cc1d8491c5640056e5dd393ea243e4dce4e8a9
2024-02-27 13:29:33 -05:00
Jaydeep Patel d1b75659f2 SWDEV-445539, SWDEV-445549 - Restore kernel launch prof data.
Issue is due to INIT_hipLaunchKernel_CB_ARGS_DATA turns out to be empty when we generate new hip_prof_str.h with since
extra space introduced by https://gerrit-git.amd.com/c/compute/ec/clr/+/846275/77/hipamd/src/hip_module.cpp#667

Change-Id: Id251956c1663483cc7ab11a3722769cb5b0d3664
2024-02-14 00:05:41 -05:00
jiabaxie 0479cdb3dd SWDEV-437702 - implement hipGetProcAddress
This should be used in place of dlsym or GetProcAddress (linux and windows respectively)

Change-Id: I5501b538e03892e8e5a2282678d848fcaf21d911
2024-02-08 11:26:02 -05:00
Anusha GodavarthySurya 5303cd2740 SWDEV-439637 - Fix undefined symbol error hipTexRefGetArray & hipTexRefGetBorderColor (github id- #3382)
Change-Id: Iae1956779e84759f23011151b60a01cef0d0ef32
2024-02-06 03:04:07 -05:00
Rahul Manocha 1a3901fa49 SWDEV-421025 - hipGraphInstantiateWithParams API changes
Change-Id: Ib07d4dd1698220b68ed27f91d58d3bd315a8804c
2024-02-05 05:08:11 +00:00
Rahul Manocha ceffb24416 SWDEV-421025 - Hip Launch Attribute ID/Value datatype for rocprof
Change-Id: Icfcd1ecf2f2956ba015820ba250f06b438169a6a
2024-02-05 05:08:11 +00:00
Sameer Sahasrabuddhe d80168eb87 SWDEV-295298: hide warp sync builtins with a macro in ROCm 6.1
Change-Id: Ie2efd233c0bcf8ad0e06223ee854fe4bd1060443
2024-02-05 05:08:11 +00:00
Anusha GodavarthySurya 242132caa7 SWDEV-439637 - Fix undefined symbol error(github id- #3382)
Change-Id: Ic7812e21618f48c737ce2c51b22a153099cc5d66
2024-02-05 05:08:11 +00:00
Jatin Chaudhary a01257f000 SWDEV-431367 - fix float compare for atomicMax/Min where -0.0 < +0.0
Change-Id: Ief99aac3775b0a0c6372215e5a0a9ecbbdca4ca5
2024-02-05 05:08:11 +00:00
Satyanvesh Dittakavi 0f9f8b52da SWDEV-435750 - Make threadIdx, blockIdx compatible with CUDA
Change-Id: If0bd2c2c72e8418e30b5a202d9c10da334eff75e
2024-01-18 02:16:15 -05:00
Jaydeep Patel fb2c100d8c SWDEV-398109 - Enable smid support for CU mode for gfx10+.
Change-Id: Ief1e7c952c3de7e4a547bbdcc9817f51aa4d5d36
2024-01-16 20:13:23 -05:00
Jonathan R. Madsen 0f1fdc7794 Add rocprofiler-register support
- Add API table versioning defines to hip_api_trace.hpp
- Add rocprofiler-register implementation in hip_api_trace.cpp
  - created ToolInit implementation
  - moved some local functions into anonymous namespace
  - renamed some function to be overloads
  - made most of the initialization code templated (reduced code duplication)
  - enforce ABI
- Updated hipamd/src/CMakeLists.txt
  - find rocprofiler-register package (enabled by default but not required)
- Address review comments
  - simplify size calculation for dispatch table
  - remove setting .size field in GetDispatchTableImpl
- Fix windows build
- Do not enforce ABI on Windows

Change-Id: I08766e8a083528a1236996274bf4522e0e8a1b9f
2024-01-13 01:02:59 -06:00
sdashmiz d23835ffbe SWDEV-421027 - Add hipGraphAddNode
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: Ic8cf293ff483ee2547b52d2975062bcb9a6f5d17
2024-01-12 11:36:30 -05:00
Sameer Sahasrabuddhe c5ab5680b4 SWDEV-295298: introduce warp sync functions
The following builtins from the CUDA spec are implemented:

  - __all_sync, __any_sync, __ballot_sync and __activemask
  - __match_any_sync and __match_all_sync
  - __shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync

The following builtins are NOT implemented, pending support in the compiler:

  - __reduce_add_sync, __reduce_min_sync, __reduce_max_sync
  - __reduce_and_sync, __reduce_or_sync, __reduce_xor_sync

Change-Id: I07dedbbfe5449f4b5c9b040bed59f5603ccec8c3
2024-01-11 22:31:27 -05:00
jujiang 2a2fea6c07 SWDEV-410522 - Fix a doxygen warning in a header in AMD platform
Change-Id: I29cd222af453c42a5d472081747bacf0c3f8d17c
2024-01-11 13:40:25 -05:00
sdashmiz 087767d9ac SWDEV-403382 - Add missing graph APIs
- add external semaphor signal and wait node

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: If205d48bc4fb60ce2966ca6ed0419e7700755bdd
2024-01-09 14:19:50 -05:00
Jaydeep Patel c3fc607e7d SWDEV-437440 - Specify which defination to use at compile time.
Change-Id: Ibeee914f293d26fa68ab793cac8b662aa34e8df3
2023-12-21 19:59:39 -05:00
Rahul Garg 92aa9d0fba SWDEV-388256 - Add support of hipExtGetLastError
Change-Id: Ie119d5ea611019996b8a80f7b820ca6a160496a6
2023-12-18 22:03:22 -05:00
Anusha GodavarthySurya 3e72b8d1e1 SWDEV-436405 - Add hipGetDevicePropertiesR0000 and hipChooseDeviceR0000 to hip dispatch table
Change-Id: I5b373ac030502eb88477d20a1d216bc48369b51d
2023-12-15 06:26:52 -05:00
sdashmiz 3a6e1b8fc5 SWDEV-421021 - move new fn pointers to end of the file
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I98d2379b3ef90282de64b22138b1dedcc3586a1b
2023-12-14 09:15:07 -05:00
Sameer Sahasrabuddhe 7137a296dd SWDEV-437090: move the __hip_assert macro to a common place
It cannot be moved to amd_device_functions.h because that causes circular
dependences when trying to use the macro in other files. So we create a new
header and move all assert/abort macros to that common header.

As a side-effect, also fix the macro to correctly expand the entire condition
argument, and also consume the trailing semicolon.

Change-Id: I43688c8e61183503a3a1a039b91321a3779152af
2023-12-14 09:03:14 -05:00
sdashmiz b8e820f835 SWDEV-421021 - Add hipDrvmemsetnode for graph
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I0240a162a16e61549d46b5d086c831404550e833
2023-12-13 09:30:36 -05:00
Ioannis Assiouras a562194562 SWDEV-433821 - Make hipStreamWaitEvent flags a default argument
Change-Id: I0192c1c5f69a1bbcc3f9919f76c947dc40737cb8
2023-12-06 12:37:35 -05:00
Jatin Chaudhary 77c581a3eb SWDEV-435702 - the functions in bf16 header need to be static
If the compiler decides not to inline these functions, we might break ODR (one definition rule) due to this file being included in multiple files and being linked together

Change-Id: Iacbfdabb53f5b4e5db8c690b23f3730ec9af16c0
2023-12-06 03:59:54 -05:00
Satyanvesh Dittakavi 7febad13be Revert "SWDEV-419480 - HIPRTC implementation of standard headers shouldn't be in std namespace"
This reverts commit 353dbe6e3b.

Reason for revert: This is considered a breaking change and requires
multiple apps to change their behavior. This will be reintroduced in later releases.

Change-Id: I3481627115af1872785585a155cc6a0ecfbe1372
2023-12-06 03:11:59 -05:00
Anusha GodavarthySurya d347f1d67d SWDEV-351966 - Remove hipRegisterTracerCallback from dispatch table
Change-Id: Ib343e2c8414a98c0cee6c60fb74a47bb859d8f45
2023-12-06 01:23:33 -05:00