[Issue]: Simplest offline compiled saxpy kernel fails to load on gfx90c

MathiasMagnus commented 2 months ago

Problem Description

Using all tried releases of Clang (for eg. HIP SDK 5.7.1) a simple saxpy kernel when compiled offline to gfx90c binary and fed to an OpenCL program, it fails to build. The build log upon failure holds:

Error while BRIG Codegen phase: the binary is incomplete

Using cl_amd_offline_devices to manifest a gfx1030 device, it builds as expected.

Operating System

Windows 11 Pro (10.0.22631)

CPU

AMD Ryzen 7 4700U with Radeon Graphics

GPU

AMD Radeon VII

ROCm Version

ROCm 5.7.1

ROCm Component

clr

Steps to Reproduce

Take the following kernel:

kernel void saxpy(
    float a,
    global float* x,
    global float* y
)
{
    const size_t gid = get_global_id(0);

    y[gid] = a * x[gid] + y[gid];
}

and compile using some Clang as such:

& 'C:\Program Files\AMD\ROCm\5.7\bin\clang++.exe' -O2 --target=amdgcn-amd-amdhsa-opencl -mcpu=gfx90c -fshort-wchar .\OpenCL-Cpp-SAXPY.cl -o .\OpenCL-Cpp-SAXPY.bin

Feed the binary to a physically present or offline device (using cl_amd_offline_devices) for eg.:

std::ifstream binary_stream{
    cl::util::executable_folder() + "/OpenCL-Cpp-SAXPY.bin",
    std::ios::binary
};
cl::Program program{
    context,
    { device },
    cl::Program::Binaries{
        cl::vector<unsigned char>(
            std::istreambuf_iterator<char>{ binary_stream },
            std::istreambuf_iterator<char>{}
        )
    }
};
program.build({ device });

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

b-sumner commented 2 months ago

@MathiasMagnus I do not believe what you are attempting is supported. Why can't the runtime compiler be used?

MathiasMagnus commented 2 months ago

@b-sumner Because I want to feed C++ for OpenCL to the runtime, and amdclang (like any other Clang-based compiler) has it.

I'm trying to push the ecosystem further.

b-sumner commented 2 months ago

@MathiasMagnus I appreciate what you're trying to do, but since it's not supported the burden will be on you to properly invoke the LLVM toolchain. The OpenCL runtime use COMgr for online compilation services. If you follow what COMgr does (see https://github.com/ROCm/llvm-project/tree/amd-staging/amd/comgr) you should be able to create correct code objects.

ROCm / clr