Remove duplicated compiler symbols from hip-rt. hip-rt should use symbols such as EF_AMDGPU_MACH_AMDGCN_XXX and
ELFABIVERSION_AMDGPU_HSA_VX defined in compiler header.
* Analysis report block based filtering is the default now
* Update documentation
* Update CHANGELOG
* Fix tests
* Replace hardware block based filtering tests with report block
based filtering tests
[ROCm/rocprofiler-compute commit: 98bb0f4237]
* Analysis report block based filtering is the default now
* Update documentation
* Update CHANGELOG
* Fix tests
* Replace hardware block based filtering tests with report block
based filtering tests
* Fix null handle
- use .handle=0, not .handle=numeric_limits<>::max()
* Update lib.common.hasher
* Fix ROCPROFILER_CONTEXT_NONE
* Use context operator==
* Update CHANGELOG
* Updated null handle for scratch memory and changed allocation test so that free ops account for null agent
---------
Co-authored-by: Ian Trowbridge <Ian.Trowbridge@amd.com>
[ROCm/rocprofiler-sdk commit: 4d6a61f5e5]
* Fix null handle
- use .handle=0, not .handle=numeric_limits<>::max()
* Update lib.common.hasher
* Fix ROCPROFILER_CONTEXT_NONE
* Use context operator==
* Update CHANGELOG
* Updated null handle for scratch memory and changed allocation test so that free ops account for null agent
---------
Co-authored-by: Ian Trowbridge <Ian.Trowbridge@amd.com>
* make sure symbolStatus after hipGetProcAddress() is the expected in the tests that are failing on Phoenix
* fix that calls to hipMemset2DAsync() and hipMemset2DAsync_spt() to set values on submatrices where not taking into account the pitch, causing test failures on Phoenix
* before testing whether hipMemset2DAsync() works, initialize the whole matrix to a known value. This makes sure the test fails even if the uninitialized memory returned after hipMallocPitch() happens to have the expected value
* Sbefore testing whether hipMemset2DAsync_spt() works, initialize the whole matrix to a known value. This makes sure the test fails even if the uninitialized memory returned after hipMallocPitch() happens to have the expected value
* fix up compiler error: dyn_hipMemset2DAsync_ptr should have been dyn_hipMemset2DAsync_spt_ptr
* when Unit_hipMalloc3D_Basic fails due to a potential memory leak, print the values that hipMemGetInfo() returns before and after the allocation/deallocation pair.
* Also print intermediate free memory
* Make sure Unit_hipMalloc3D_Basic allocates at least PalSettings::subAllocationChunkSize_ on Windows. Otherwise hipMemGetInfo() will not report an increase on available memory after hipFree() is called, as the minimum amount of memory that cause a call to the CoarseMemorySubAllocator::Create() is that chunk size
* Fix up previous commit; allocated too much memory
* skip some extent sizes in Unit_hipFreeMipmappedArrayImplicitSyncArray if allocating them would be require more memory than the actual totalGlobalMem of the device
* Do not expect an exact match when comparing the memory available memory before and after hipMalloc3D() + hipFree()
* Do not allocate more memory than the total GPU memory in Unit_hipFreeMipmappedArrayImplicitSyncArray
* fix expected available memory amount comparison in hipMalloc3D - Basic
* use SUCCEED() macro in Unit_hipFreeMipmappedArrayImplicitSyncArray to log more information when there is not enough memory for the mipmapped array to be allocated
* fix formatting
- Refactored scratch memory handling by introducing fmm_is_scratch_aperture to
replace repeated for-loops.
- Simplified code paths in hsakmt_fmm_release, hsakmt_fmm_map_to_gpu, and
hsakmt_fmm_unmap_from_gpu by using the new helper.
Signed-off-by: Honglei Huang <Honglei1.Huang@amd.com>
* replace azure runners with internal
* change to mi300a for debug
* revert back to mi300
* move some of the load to mi300a
* use mi300a for clang-tidy
---------
Co-authored-by: Kandula, Venkateshwar reddy <Venkateshwarreddy.Kandula@amd.com>
[ROCm/rocprofiler-sdk commit: f4cce0d2d4]
* replace azure runners with internal
* change to mi300a for debug
* revert back to mi300
* move some of the load to mi300a
* use mi300a for clang-tidy
---------
Co-authored-by: Kandula, Venkateshwar reddy <Venkateshwarreddy.Kandula@amd.com>
* Arbitrary host-trap sampling skid (doc)
The host-trap PC sampling might introduce a skid of [0, 2]
instructions. We documented this information and provides
some advice to application developers how to find
hot-spots in the profiles generated by host-trap sampling.
[ROCm/rocprofiler-sdk commit: 650d35bdaa]
* Arbitrary host-trap sampling skid (doc)
The host-trap PC sampling might introduce a skid of [0, 2]
instructions. We documented this information and provides
some advice to application developers how to find
hot-spots in the profiles generated by host-trap sampling.
Baraldi, Giovanni
Kent Russell
Sang, Tao
Pryor, Adam
Sajina PK
Yiannis Papadopoulos
Joseph Macaranas
vedithal-amd
Dittakavi, Satyanvesh
Jiang, Julia
Ma, Bing
Kandula, Venkateshwar reddy
Luca Bruni
Galantsev, Dmitrii
Indic, Vladimir
Madsen, Jonathan
Nagaraj, Sriraksha
Hernandez, Gerardo
Honglei Huang
Bonnell, Jason
systems-assistant[bot]
Joseph Macaranas
Flora Cui
Bhardwaj, Gopesh
Betigeri, Sourabh