Skip to content

Instantly share code, notes, and snippets.

Embed
What would you like to do?
A bug in AMD OpenCL compiler
#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;
}
}
@ddemidov

This comment has been minimized.

Copy link
Owner Author

@ddemidov ddemidov commented May 14, 2013

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.

@ddemidov

This comment has been minimized.

Copy link
Owner Author

@ddemidov ddemidov commented May 14, 2013

Another possible solution is replacing line 68 with

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

This comment has been minimized.

Copy link
Owner Author

@ddemidov ddemidov commented May 14, 2013

The odd thing is that

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

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

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment