Commit Graph

1705 Commits

Author SHA1 Message Date
Julia Jiang 3bd32d7d4b SWDEV-497618 - Define GL group of doxygen for HIP API
Change-Id: Ia43f25201df960b905e5fd45481f486e7bffabdb
2024-11-15 11:25:51 -05:00
Sameer Sahasrabuddhe 556390f9c7 SWDEV-490198: _sync() will be enabled by default in 6.4
Change-Id: Id029424a9c0f6b144a7aa0e96fe8acc4a138ec51
2024-10-25 09:54:40 -04:00
Rahul Manocha e729f08704 SWDEV-468039,SWDEV-482579 - Enable FP8 SW Conversions on pre gfx940 archs
1) SW Conversions for ocp and fnuz are enabled on pre mi300 archs
2) for mi300 only fnuz is enabled
3) for gfx1200 only ocp is enabled

Change-Id: I90373752a2d15eff20d5deec874ed396ba4e1788
2024-10-17 11:49:22 -04:00
Jatin Chaudhary b977101893 SWDEV-486137 - match behavior of int variants of hadd/uhadd/rhadd/urhadd
Match cases and handle cases where it can overflow.

Change-Id: I3d6f802686af230a622ef9891a844135ad3d1ae5
2024-10-09 13:47:33 -04:00
Jaydeep Patel d6193a2f23 SWDEV-483436 - User spt stream as def with -fgpu-default-stream=per-thread for hipMemsetAsync.
Change-Id: Ia85c2b4c40fc9250754d3b64fb9fd1c615362572
2024-09-29 01:42:33 -04:00
Jonathan R. Madsen 07c9c7fe56 Fix HIP API trace versioning
Change-Id: I33f2be4668c96e2225d4ca9a253e61ec2dc65102
2024-09-25 10:32:14 -04:00
Julia Jiang bb03ef11a3 SWDEV481762 - Updated definition of 'DEPRECATED' in header file
Change-Id: I88986b8e1815f3d816595f3eb2da8a6c1c1c2993
Jenifer helped make a combined PSDB build, together with the change in hip repos
https://gerrit-git.amd.com/c/compute/ec/hip/+/1114046
Combined PSDB verification passed.
http://rocm-ci.amd.com/job/compute-psdb-staging-hip/17293/
2024-09-05 15:41:04 -04:00
victzhan 8be00b6602 SWDEV-477218 - Implement hipDeviceGetTexture1DLinearMaxWidth
Change-Id: I8103f710abeb869f5f84be61c57a30b24356def6
2024-09-05 15:09:38 -04:00
Ioannis Assiouras 2c84211b58 SWDEV-470372 - Added hipExtHostAlloc API
This change adds a new HIP API `hipExtHostAlloc` which preserves
the functionality of `hipHostMalloc`.

Change-Id: I13504c6fc13465ddd7aed329795bb4f2fef1baff
2024-08-27 08:26:03 -04:00
Jatin Chaudhary bb52d9ed62 SWDEV-480489 - fix unsafeAtomicAdd
Integration into pytorch pointed out some issues, value narrowing, to
fix this we are now using unions. Also removed check for -munsafe*
compiler flag. The check is now just on builtin detection.

Change-Id: I49364503fa429bd862952f9b29879072afa6d553
2024-08-27 06:29:11 -04:00
Sameer Sahasrabuddhe 6df2da65cd SWDEV-480725: missing __ockl_wfall __ockl_wfany in amd_hip_bf16.h
Change-Id: Iff4aeec411bfeaf4cc187c515e2da3d5898f89cb
2024-08-25 22:49:14 -04:00
Satyanvesh Dittakavi d69cc35750 SWDEV-473942 - SWDEV-431367 - Correct atomicMax(_system) and atomicMin(_system)
- Fixes -0.0 and +0.0 comparison. For atomicMax if the value on
address is -0.0 and on val is +0.0, gfx90a's unsafe atomics will swap
them. This behavior should be consistent with cas loop as well.

- _system variants of atomicMax and atomicMin are resulting in
incorrect output. Updated these to use the similar implementation as
atomicMax and atomicMin.

Change-Id: I20df36ee29ae0434a6b564f2ba71193fe41cfa59
2024-08-13 10:38:50 -04:00
Rahul Manocha 39bbc0341d SWDEV-468039 - FP8 host only conversion support on mi200
Change-Id: I0891f42d1b7c0d94d099fe26df5db3eff64ba564
2024-08-07 20:51:00 -04:00
Jaydeep Patel 6344ddb2f3 SWDEV-470886 - Add maybe_undef attribute for shfl device function due to not all lanes of wave define var and compiler needs to know about this.
Change-Id: I3a683887e033305ac55362f356838b491a6d50f2
2024-08-05 00:53:13 -04:00
Sourabh Betigeri 7298b80112 SWDEV-365151 - Fix the fns32 to do 32 bit computation and adds a wrapper to ease porting from CUDA
Change-Id: I0b5a9ca11c98f8c1c40cfba7f4e057bfda2d756e
2024-07-26 11:20:14 -04:00
Sourabh Betigeri 4fbd7abbb2 SWDEV-475394 - Fix for the return type to be in-line with CUDA
Change-Id: I7c833571d47b4e86a86e4a0095b61947d16ecab6
2024-07-25 16:33:33 -04:00
Rahul Manocha 353f15afa6 SWDEV-468039 - FP8 OCP headers
Change-Id: Iecd32c5a0357781da07395d32f894415954b7b22
2024-07-25 12:42:23 -04:00
Jatin Chaudhary e43176bde9 SWDEV-466747 - add shfl functions in bfloat16
Change-Id: Ide7d7e1d449783cced8867abf43ff45f5bce113a
2024-07-23 10:51:02 -04:00
Jatin Chaudhary 6a655a77e7 SWDEV-466747 - optimize conversions for bfloat16 operations
Since we made the members public, we can optimize some operations which
do not require redundant conversions to half_raw types.

Change-Id: I31555ef18e695d8e24b89f0418187fa4e932a38a
2024-07-17 18:37:25 -04:00
Julia Jiang dd30e0e893 SWDEV-472710 - Adding gitattributes and remove trailing spaces
Change-Id: Ic8ad2071745f0ffe6a2e120bfebb6d90bf270f87
2024-07-15 12:39:56 -04:00
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