ROCm
How to get isa code with HIP RTC?
I am doing a program with HIPRTC. I want to have a look at the isa code to consider I can enhance its performance. I have known that in cuda we can use nvrtcGetPTXSize and nvrtcGetPTX to get ptx code with nvrtc. Is there any familiar method to get isa code with hiprtc?
The hipRTC equivalent APIs for nvrtcGetPTXSize and nvrtcGetPTX are hiprtcGetCodeSize and hiprtcGetCode respectively. Please use these and check if it helps.
The hipRTC equivalent APIs for nvrtcGetPTXSize and nvrtcGetPTX are hiprtcGetCodeSize and hiprtcGetCode respectively. Please use these and check if it helps.
Thank you very much. But I've found that 'hiprtcGetCode' is going to generate binary file, and it's not readable. Is there any way to read it?
@yxsamliu Any suggestions regarding this?
I've added '--save-temps' to the function 'hiprtcCompileProgram', and I can get the readable isa now, thank you very much for help.
@yxsamliu Any suggestions regarding this?
Sorry, I have a other question. When I am using '--save-temps' to get the isa, it returns HIPRTC_ERROR_COMPILATION for the function 'hiprtcCompileProgram'. Do you know the reason? Below is how I use it, and my compile command is 'hipcc -I /opt/rocm/hip/include/hip saxpy.hip.cpp'
.
It seem to work fine when verified with latest hip using --save-temps . Please share the sample and the release with which you are trying in case you still see the issue.
It seem to work fine when verified with latest hip using --save-temps . Please share the sample and the release with which you are trying in case you still see the issue.
My Code is this:
#include <hip/hip_runtime.h>
#include <hiprtc.h>
#include <cassert>
#include <iostream>
#include <vector>
std::string source = R"(
extern "C"
__global__
__launch_bounds__(256)
void saxpy(int n, const float a, float const* x, int incx, float* y, int incy)
{
int i = blockDim.x*blockIdx.x + threadIdx.x;
y[i] += a*x[i];
}
)";
char const* src = source.c_str();
#define HIPRTC_CALL(call) assert(call == HIPRTC_SUCCESS);
#define HIP_CALL(call) assert(call == hipSuccess);
int main()
{
hiprtcProgram prog;
HIPRTC_CALL(hiprtcCreateProgram(&prog, src, nullptr, 0, nullptr, nullptr));
hiprtcResult result;
const char *opts[] = {"--save-temps"};
result = hiprtcCompileProgram(prog, 1, opts);
if (result != HIPRTC_SUCCESS) {
std::string s_error = hiprtcGetErrorString(result);
std::cout << s_error << std::endl;
std::size_t log_size;
hiprtcGetProgramLogSize(prog, &log_size);
std::string log;
log.reserve(log_size);
hiprtcGetProgramLog(prog, &log[0]);
std::cout << log << std::endl;
exit(EXIT_FAILURE);
}
return 0;
}
I am using rocm-5.2.0, and my command to build it is $hipcc -I /opt/rocm/hip/include/hip saxpy.hip.cpp Thank you very much!
@guangzlu Thanks for sharing the sample. The issue was due to hiprtcCreateProgram expecting a program name instead of nullptr when "--save-temps" option is passed. Fix has been added internally. However, it might take few days to be seen in the github. Meanwhile, as a workaround, please use some valid program name instead of nullptr to avoid this error. eg: HIPRTC_CALL(hiprtcCreateProgram(&prog, src, "saxpy", 0, nullptr, nullptr));
@guangzlu Hi, is this resolved on the latest HIP? If so can we please close this ticket?
