Skip to content

[amdgpu][openmp] Treat missing TIMESTAMP_FREQUENCY as non-fatal #70987

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Nov 1, 2023
Merged

[amdgpu][openmp] Treat missing TIMESTAMP_FREQUENCY as non-fatal #70987

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
31 changes: 16 additions & 15 deletions libc/utils/gpu/loader/amdgpu/Loader.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -487,21 +487,22 @@ int load(int argc, char **argv, char **envp, void *image, size_t size,
handle_error(err);
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_clock_freq);

if (hsa_status_t err =
hsa_agent_get_info(dev_agent,
static_cast<hsa_agent_info_t>(
HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
host_clock_freq))
handle_error(err);

void *freq_addr;
if (hsa_status_t err = hsa_executable_symbol_get_info(
freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &freq_addr))
handle_error(err);

if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,
host_agent, sizeof(uint64_t)))
handle_error(err);
if (HSA_STATUS_SUCCESS ==
hsa_agent_get_info(dev_agent,
static_cast<hsa_agent_info_t>(
HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
host_clock_freq)) {

void *freq_addr;
if (hsa_status_t err = hsa_executable_symbol_get_info(
freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
&freq_addr))
handle_error(err);

if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,
host_agent, sizeof(uint64_t)))
handle_error(err);
}
}

// Obtain a queue with the minimum (power of two) size, used to send commands
Expand Down
10 changes: 6 additions & 4 deletions openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1810,10 +1810,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
return Err;
GridValues.GV_Warp_Size = WavefrontSize;

// Get the frequency of the steady clock.
if (auto Err = getDeviceAttr(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY,
ClockFrequency))
return Err;
// Get the frequency of the steady clock. If the attribute is missing
// assume running on an older libhsa and default to 0, omp_get_wtime
// will be inaccurate but otherwise programs can still run.
if (auto Err = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

So Raw does not emit an error. I thought the issue would be that HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY may not be defined at all so it would fail to compile?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If the enum field is missing, no compilation, no progress. I don't think there's a reasonable way to reflect on whether an enum field is present or not (i.e. no #ifdef equivalent).

However the enum field is present in the dynamic_hsa header, so openmp compiles just fine against that.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Was assuming there'd be some version field in HSA that indicates it, but I don't know. This should be better than nothing for now.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is no version in the HSA library, it's on the wishlist. There is one in the header, but as the header may bear little relation to the shared library, that is a partial help.

ClockFrequency))
ClockFrequency = 0;

// Load the grid values dependending on the wavefront.
if (WavefrontSize == 32)
Expand Down