Problem Description
The ROCm 6.4 release notes document that, as part of deprecating roc-obj, llvm-objdump --offloading supports --arch-name to extract only the code objects matching a given target architecture.
However, it seems that in ROCm 7.14, --arch-name is broken. No objects are returned if this flag is passed. Moreover, the release notes do not appear to mention changes to the tooling.
Reproducer
main.cpp
#include <concepts>
#include <iostream>
#include "hip/hip_runtime.h"
#define CHECK_HIP_CALL(__statement__) \
{ \
const auto __status__ = __statement__; \
if(__status__!= hipSuccess) \
{ \
std::cout << #__statement__ " status: " \
<< __status__ \
<< " " << hipGetErrorString(__status__) \
<< std::endl; \
throw std::runtime_error( \
std::string("Statement ") + #__statement__ " failed at " + __FILE__ + ":" + std::to_string(__LINE__)); \
} \
}
__global__ void my_kernel(int* const data)
{
const int idx = threadIdx.y;
data[idx] = idx;
}
int main()
{
//! Retrieve kernel attributes.
hipFuncAttributes attrs{};
CHECK_HIP_CALL(hipFuncGetAttributes(&attrs, (void*)my_kernel));
std::cout << "> Kernel attributes:" << std::endl;
std::cout << "\t- number of registers: " << attrs.numRegs << std::endl;
//! Create a stream.
hipStream_t stream = nullptr;
CHECK_HIP_CALL(hipStreamCreate(&stream));
//! Allocate data.
constexpr size_t size = 128;
int* data = nullptr;
CHECK_HIP_CALL(hipMalloc((void**)&data, size * sizeof(int)));
//! Launch kernel.
constexpr dim3 grid {1, 1, 1};
constexpr dim3 block {1, size, 1};
constexpr unsigned int shared = 0;
my_kernel <<<grid, block, shared, stream>>>(data);
//! Synchronize and cleanup.
CHECK_HIP_CALL(hipStreamSynchronize(stream));
CHECK_HIP_CALL(hipFree(data));
CHECK_HIP_CALL(hipStreamDestroy(stream));
}
hipcc --offload-arch=gfx942 -std=c++20 main.cpp
On ROCm 6.4.3:
# extracts `a.out:0.hipv4-amdgcn-amd-amdhsa--gfx942`
llvm-objdump --arch-name=gfx942 --offloading a.out
On ROCm 7.14:
# extracts no gpu code object
llvm-objdump --arch-name=gfx942 --offloading a.out
# extracts `a.out.0.hipv4-amdgcn-amd-amdhsa--gfx942`
llvm-objdump --arch=gfx942 --offloading a.out
In addition, in both releases, the documentation of --arch-name appears erroneous. In both releases, the docstring refers to --version to know the valid targets that may be passed. But the list of valid targets includes only targets such as amdgcn; in both releases, specific architectures such as amd_gfx1201 are not in the list.
This issue is thus a question to know what is the recommended way to extract GPU code objects on ROCm as of 7.14, as well as a suggestion to reflect changes in the documentation and in the release notes.
Operating System
Ubuntu
CPU
independent of cpu
GPU
independent of gpu
ROCm Version
6.4.3 vs 7.14
ROCm Component
No response
Steps to Reproduce
No response
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
rocminfo --support output
Additional Information
No response
Problem Description
The ROCm 6.4 release notes document that, as part of deprecating
roc-obj,llvm-objdump --offloadingsupports--arch-nameto extract only the code objects matching a given target architecture.However, it seems that in ROCm 7.14,
--arch-nameis broken. No objects are returned if this flag is passed. Moreover, the release notes do not appear to mention changes to the tooling.Reproducer
main.cpp
On ROCm 6.4.3:
On ROCm 7.14:
In addition, in both releases, the documentation of
--arch-nameappears erroneous. In both releases, the docstring refers to--versionto know the valid targets that may be passed. But the list of valid targets includes only targets such asamdgcn; in both releases, specific architectures such asamd_gfx1201are not in the list.This issue is thus a question to know what is the recommended way to extract GPU code objects on ROCm as of 7.14, as well as a suggestion to reflect changes in the documentation and in the release notes.
Operating System
Ubuntu
CPU
independent of cpu
GPU
independent of gpu
ROCm Version
6.4.3 vs 7.14
ROCm Component
No response
Steps to Reproduce
No response
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
rocminfo --support output
Additional Information
No response