MacOSX OpenCL problems

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++) {
        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()))
            );
    try {
        program.build(device);
    } catch(const cl::Error&) {
        std::cerr
            << "OpenCL compilation error" << std::endl
            << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device[0])
            << std::endl;
        throw;
    }
    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,
                "#if defined(cl_khr_fp64)\n"
                "#  pragma OPENCL EXTENSION cl_khr_fp64: enable\n"
                "#elif defined(cl_amd_fp64)\n"
                "#  pragma OPENCL EXTENSION cl_amd_fp64: enable\n"
                "#endif\n"
                " \n"
                "kernel void vexcl_vector_kernel\n"
                "(\n"
                "  ulong n,\n"
                "  global double * prm_1\n"
                ")\n"
                "{\n"
                "  ulong chunk_size  = (n + get_global_size(0) - 1) / get_global_size(0);\n"
                "  ulong chunk_start = get_global_id(0) * chunk_size;\n"
                "  ulong chunk_end   = chunk_start + chunk_size;\n"
                "  if (n < chunk_end) chunk_end = n;\n"
                "  for(ulong idx = chunk_start; idx < chunk_end; ++idx)\n"
                "  {\n"
                "    prm_1[idx] = 42.0;\n"
                "  }\n"
                "}\n"
                );
    } catch (const std::exception &err) {
        std::cerr << err.what() << std::endl;
        return 1;
    }
}

Compile with something like g++ -o compiler_bug compiler_bug.cpp -std=c++0x -lOpenCL

@ddemidov This crashes the compiler/linker as well. If I replace the min() with 0 it works so there's something happening in the compiler there.