A bug in AMD OpenCL compiler

compiler_bug.cpp

#include <iostream>
#include <vector>
#include <string>
#include <stdexcept>

#define __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>

//---------------------------------------------------------------------------
void precondition(bool cond, const std::string &msg) {
    if (!cond) throw std::runtime_error(msg);
}

//---------------------------------------------------------------------------
cl::Device get_device() {
    // Get list of OpenCL platforms.
    std::vector<cl::Platform> platform;
    cl::Platform::get(&platform);

    precondition(!platform.empty(), "No OpenCL platforms.");

    // Get first available device.
    for(auto p = platform.begin(); p != platform.end(); p++) {
        if (p->getInfo<CL_PLATFORM_NAME>().find("AMD") == std::string::npos) continue;

        std::vector<cl::Device> device;
        p->getDevices(CL_DEVICE_TYPE_ALL, &device);

        if (!device.empty()) return device[0];
    }

    precondition(false, "No compute devices.");
}

//---------------------------------------------------------------------------
cl::Program build_program(
        const cl::Context &context,
        const std::vector<cl::Device> &device,
        const std::string &source
        )
{
    cl::Program program(context,
            cl::Program::Sources(1, std::make_pair(source.c_str(), source.size()))
            );
    program.build(device);
    return program;
}

//---------------------------------------------------------------------------
int main() {
    try {
        std::vector<cl::Device> device;
        device.push_back(get_device());
        std::cout << device[0].getInfo<CL_DEVICE_NAME>() << std::endl;

        cl::Context context(device);

        // Compile OpenCL program for the device.
        cl::Program program = build_program(context, device,
            "typedef struct {\n"
            "    int first;\n"
            "    float second;\n"
            "} _pair_int_float_t;\n"
            "\n"
            "__kernel void copy(__global _pair_int_float_t* _buf0)\n"
            "{\n"
            "const uint i = get_global_id(0);\n"
            "_buf0[i]=(_pair_int_float_t){4, 2};\n"
            "\n"
            "}\n"
            );
    } catch (const std::exception &err) {
        std::cerr << err.what() << std::endl;
        return 1;
    }
}

When this is compiled with g++ -o compiler_bug compiler_bug.cpp -std=c++0x -lOpenCL, it outputs

$ ./compiler_bug
Capeverde
UNREACHABLE executed!
Aborted

If line 68 is replaced with

            "_buf0[i].first = 4;\n"
            "_buf0[i].second = 2;\n"

it completes successfully.

The program works unchanged with NVIDIA or Intel OpenCL implementations.

Another possible solution is replacing line 68 with

            "_pair_int_float_t s = {4, 2};\n"
            "_buf0[i]=s;\n"

The odd thing is that

            "_buf0[i]=(_pair_int_float_t){0, 0};\n"

(I've changed constants in right hand side) actually works.