Commit Graph

39 Commits

Author SHA1 Message Date
Branislav Brzak c2d1776ebd SWDEV-516564 - SWDEV-512817 - Remove mentions of gfx940 and gfx941
Change-Id: Ia069fcb9c6948c3fc9a00961593c9dcc59609375
2025-03-05 04:26:07 -05:00
taosang2 32daa8f384 SWDEV-501963 - Add missing codes for gfx950
Cherry-pick https://gerrit-git.amd.com/c/compute/ec/clr/+/1162997

Change-Id: I6b3c6bf55c61cffd43cd6f17b75998f751b75723
2025-01-31 14:34:49 -05: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
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
Ioannis Assiouras 8f42ad6aa3 SWDEV-464648 - code and comment cleanups
Change-Id: I5ba3f1bff500b3cd5903c2f441017735e688f83f
2024-06-07 22:38:09 +01:00
Ioannis Assiouras 74cae705ae SWDEV-445960 - Fix __umul64hi __mul64hi on windows
Change-Id: I4857d269bc80fcc0e6b77df941f24ff1964a4d42
2024-03-03 12:42:42 -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
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
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
Satyanvesh Dittakavi 270c6d48f1 Revert "SWDEV-431175 - Remove datatypes from HIPRTC"
This reverts commit 629e279f72.

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: I0354ce4e0f5e6c402499a7a8c2aaf43bf5b1bfc7
2023-12-05 02:32:38 -05:00
Satyanvesh Dittakavi f0010f4fdb SWDEV-434362 - Replace __HIP_CLANG_ONLY__ as it is not supported on HIPRTC
__HIP_CLANG_ONLY__ is not recognized in HIPRTC, due to which some
math functions like amd_mixed_dot were not included in hiprtc builtins.

Change-Id: I1fe41e1ddc8911f6a5b5b1405dd4730d0170a4f7
2023-11-29 04:22:19 -05:00
Satyanvesh Dittakavi 629e279f72 SWDEV-431175 - Remove datatypes from HIPRTC
Defining int64_t, uint64_t, int32_t, uint32_t in HIPRTC
seem to result in conflicts with some apps as they use
their own definitions for these types. NVRTC also doesn't
define these. Hence remove them to match the behavior.

Change-Id: I77ef70e846950698cb00375f5d0501b907f01fe3
2023-11-28 11:32:12 -05:00
Joseph Greathouse a0f29b454c SWDEV-431560 - use XCC_ID in __smid() for gfx941 and gfx942
__smid() needs to use both HW_ID and XCC_ID for gfx940, gfx941, and
gfx942. Previously, we only did this for gfx940 and thus XCC_ID
was incorrectly not passed back on the other two architectures.

Change-Id: I9fb13b6cef3280e15463443a180174629d03f8b2
2023-11-15 10:10:47 -05:00
Satyanvesh Dittakavi b5d286a6d3 SWDEV-430963 - Address MIGraphX failures due to unsafe buffer access usage
Compiler seem to be stricter in compiler staging builds related to
safe buffer programming when compared to other component staging builds.
This seem to result in additional errors when -Werror is enabled
in MIGraphX tests.

Removes all the clang pragmas to ignore several type of warnings in all
the headers and adds a single pragma which ignores all warnings using
#pragma clang diagnostic ignored "-Weverything" in hiprtc builtins.

Change-Id: I95f302bb285b2451b19dd5dfdb7df29164b0f750
2023-11-09 04:39:49 -04:00
Jatin Chaudhary 0668f20780 SWDEV-410898 - Add includes in files that use them
Change-Id: I14d9ea233d57c7267da3f9688305b949ced71a53
2023-09-25 08:45:22 -04:00
Satyanvesh Dittakavi d7d0f11318 SWDEV-395996 - Add HIPRTC support for missing hip headers
hip headers such as hip_math_constants.h, library_types.h,
hip_common.h, channel_descriptor.h, device_functions.h, hip_complex.h,
hip_texture_types.h, math_functions.h, surface_types.h
are added in HIPRTC

Change-Id: I4a4c198449ceb609c3ff55e00b43056c1f085431
2023-09-21 03:11:06 -04:00
Jaydeep Patel 9daf8e3177 SWDEV-404118 - Change to uint.
Change-Id: Ib9ced171a4330e868a584bfb86d7ee5418e784ed
2023-06-06 13:17:56 -04:00
Ajay 16c6b365c2 SWDEV-400515 - populate __HIP_NO_IMAGE_SUPPORT flag
MI300 does not support image APIs.
Apps to use __HIP_NO_IMAGE_SUPPORT instead of arch

Change-Id: I91178bfd27ea7b7188e7a958a876c0264f4469aa
2023-06-06 12:11:41 -04:00
Satyanvesh Dittakavi a7041d15a7 SWDEV-399203 - Use clang pragma only with clang
Change-Id: I21d84ccde07e243b6bc06f1980c24b399e74d1c3
2023-05-10 04:19:51 -04:00
Jaydeep Patel a167abefe5 SWDEV-396277 - Offset and bit size correction, Considaration of SA ID for Navi.
Correct sm id formula for GFX10 and GFX11.

Change-Id: I2883c139b8e684e0334ccc81f703a09e3ddec588
2023-05-05 01:04:42 -04:00
Jaydeep Patel 0ad47eb3e5 SWDEV-319526 - SE ID size incresed.
Change-Id: Ia7df94d172102b321f12c8caa821e3a7e4447a21
2023-02-21 22:13:57 -05:00
Matthew Arsenault 8c7391398a SWDEV-1 - Reapply "SWDEV-1 - Directly call __builtin_amdgcn_fence"
This reverts commit 5b39c346ad.

Resubmit after https://github.com/ROCmSoftwarePlatform/rocPRIM/commit/ba8a86166ec9c4cf3a40ecde7017feeb1484e7f6

Change-Id: I194559484a5a0d7f4443381fe77218ea73d2def9
2023-02-03 10:54:27 -05:00
Satyanvesh Dittakavi 0ccd41050b SWDEV-372396 - Address the errors in hiprtc headers with -Werror
Change-Id: Ied553216dfc2ff8b5c72de617cb47b6ca8f52063
2023-01-23 17:25:42 +00:00
Maneesh Gupta 5b39c346ad SWDEV-378321 - Revert "SWDEV-1 - Directly call __builtin_amdgcn_fence"
This reverts commit 3e54bad5b3.

Reason for revert: Root cause for regression SWDEV-378321

Change-Id: If3116c5ad1a23db1840e5fc019d2fe997b37d229
2023-01-19 13:01:27 -05:00
Matt Arsenault 8615a89f24 SWDEV-1 - Use proper builtin for groupstaticsize
Remove the now empty header for asm intrinsic wrappers.

Change-Id: I63495dc6ed899b1ea823062f2be2fb064af8d4f9
2023-01-13 13:14:10 -05:00
Satyanvesh Dittakavi c193503722 SWDEV-376386 - Support __hip_move_dpp with hipRTC
Change-Id: If9a6fe23d6df143a7f522823fca7183416048c6a
2023-01-13 00:05:31 -05:00
Matt Arsenault 3e54bad5b3 SWDEV-1 - Directly call __builtin_amdgcn_fence
The device libraries provide hip.bc, which only contains
__atomic_work_item_fence. This is identical to OpenCL's
atomic_work_item_fence, except prefixed with __. This only existed to
switch over the unknown values the OpenCL function provides. There's
no need to go through the wrapper here since it's using constant
values, and without it we can drop one of the bitcode libraries. I'm
guessing this was a holdover from before the fence builtin existed.

Change-Id: I8b04614f96194eafddeba295215c12c3856e6c42
2023-01-12 12:04:09 -05:00
kjayapra-amd 5ff27faa04 SWDEV-337452 - Adding__ockl_steadyctr_u64 option for wall_clock64.
Change-Id: I4ee6fd4d40d500363ef84975078185093dab5ddd
2022-11-07 08:47:30 -05:00
Tao Sang 9f9fd859db SWDEV-286739 - Support hipDeviceAttributeWallClockRate
Part 3: Add missing declaration of wall_clock64() to fix
compiling issue in device code.
Add querying hipDeviceAttributeWallClockRate.

Change-Id: Ie54771c2f58eeaacdc0248bc116ef193f99eb9b9
2022-07-28 13:02:42 -04:00
Brian Sumner a47887aa0d SWDEV-343278 - fix conversion functions
Change-Id: Ib93f886363314e980462ef562c998981bb958569
2022-07-15 16:00:58 -04:00
Ajay 9eb45dd1e0 SWDEV-345188 - hipamd: warnings treated as errors for catch tests
Change-Id: I8981a530da2e0213c60f7a81b09d9c36e27e9025
2022-07-12 15:34:39 -04:00
Satyanvesh Dittakavi 3cc0a3fcaf SWDEV-321656 - Address incorrect warp size with hiprtc
Change-Id: I97051d4e2cf27cb073c39663c62e43d5c19b1a12
2022-03-15 06:52:21 -04:00
Sourabh Betigeri 7d8c63cc34 SWDEV-294375 - Declares fns() functions static
Change-Id: Ife0150b8d6b9531522e61a82df6b88af3282f437
2022-02-13 23:35:46 -05:00
Sourabh Betigeri 523e5a007c SWDEV-294375 - Adding coalesced_groups to ROCm CG
Change-Id: I4098e1970ae2c388171fa73deb5bc023ed1b7a68
2022-02-07 11:29:46 -05:00
Satyanvesh Dittakavi 673c2deed1 SWDEV-302147 - fix __hip_ds_permutef device function to return float
Change-Id: Ie452e90a7f7410401df4a66de468ad850403953c
2021-09-14 02:15:53 -04:00
agunashe d9d9e81acb SWDEV-293742 - Update copyrights end year for hipamd
Change-Id: I08f620f84563a9214b59f1b943ed091b67229eab
2021-07-09 12:08:39 -04:00
Maneesh Gupta 1e35b26108 SWDEV-234709 - SWDEV-258907 - Reimplement assert() using the new fprintf_stderr facility
This reverts commit 18234e1d24.

Change-Id: Ic1bbf4e4e6f950934d9ba02224ef61a96bdd0e9b
2021-07-07 00:57:36 -07:00
Rahul Garg 19c84bc604 ROCMOPS-1956 - Push restructured code to hipamd
hipamd will have AMD's ROCCLR based HIP backend implementation

Change-Id: Id7de9634519b4ce46fca71a1b61f3d5b1e3fc459
2021-06-07 21:42:44 +00:00