cocl-issue3-a.cpp

#include "CL/cl.h"

#include <iostream>
#include <stdexcept>

using namespace std;

string kernelSource = R"(
kernel void _z8setValuePfif(global float* data, long data_offset, int idx, float value) {
   data = (global float*)((global char *)data + data_offset);

   label0:;
   int v1 = get_local_id(0);
    bool v2 = v1 == 0;
    if(v2) {
        goto v4;
    } else {
        goto v5;
    }
    v4:;
    long v6 = idx;
    global float* v7 = (&data[v6]);
    v7[0] = value;
        goto v5;
    v5:;
    return;
}
)";

int main(int argc, char *argv[]) {
    cl_int err;
    cl_platform_id platform = 0;
    cl_device_id device = 0;
    cl_context_properties props[3] = { CL_CONTEXT_PLATFORM, 0, 0 };
    cl_context context = 0;
    cl_command_queue queue = 0;
    cl_mem bufA, bufB, bufC;
    cl_event event = NULL;
    int ret = 0;

    /* Setup OpenCL environment. */
    err = clGetPlatformIDs(1, &platform, NULL);
    if (err != CL_SUCCESS) {
        printf( "clGetPlatformIDs() failed with %d\n", err );
        return 1;
    }

    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    if (err != CL_SUCCESS) {
        printf( "clGetDeviceIDs() failed with %d\n", err );
        return 1;
    }

    props[1] = (cl_context_properties)platform;
    context = clCreateContext(props, 1, &device, NULL, NULL, &err);
    if (err != CL_SUCCESS) {
        printf( "clCreateContext() failed with %d\n", err );
        return 1;
    }

    queue = clCreateCommandQueue(context, device, 0, &err);
    if (err != CL_SUCCESS) {
        printf( "clCreateCommandQueue() failed with %d\n", err );
        clReleaseContext(context);
        return 1;
    }

    size_t src_size = kernelSource.size();
    const char *source_pchar = kernelSource.c_str();
    cl_program program = clCreateProgramWithSource(context, 1, &source_pchar, &src_size, &err);
    if (err != CL_SUCCESS) {
        printf( "clCreateProgramWithSource() failed with %d\n", err );
        clReleaseContext(context);
        return 1;
    }

    string options = "";
    err = clBuildProgram(program, 1, &device, options.c_str(), NULL, NULL);

    char* build_log;
    size_t log_size;
    err = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
    if (err != CL_SUCCESS) {
        printf( "clGetProgramBuildInfo() failed with %d\n", err );
        clReleaseContext(context);
        return 1;
    }
    build_log = new char[log_size+1];
    err = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL);
    if (err != CL_SUCCESS) {
        printf( "clGetProgramBuildInfo() failed with %d\n", err );
        clReleaseContext(context);
        return 1;
    }
    build_log[log_size] = '\0';
    string buildLogMessage = "";
    string sourceFilename = "filename";
    if(log_size > 2) {
        buildLogMessage = sourceFilename + " build log: "  + "\n" + build_log;
        cout << buildLogMessage << endl;
    }
    delete[] build_log;

    string kernelName = "_z8setValuePfif";
    cl_kernel kernel = clCreateKernel(program, kernelName.c_str(), &err);
    if(err != CL_SUCCESS) {
        cout << "kernel build err" << endl;
        return 1;
    }


    int N = 32;
    cl_float *a = new cl_float[N];
    cl_long offset = 0;
    cl_int pos = 2;
    cl_float value = 123;

    for(int i = 0; i < N; i++) {
        a[i] = 555;
    }
    cl_mem a_gpu = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * N, 0, &err);
    if(err != CL_SUCCESS) {
        throw runtime_error("a_gpu create failed");
    }

    err = clEnqueueWriteBuffer(queue, a_gpu, CL_TRUE, 0,
                                          N * sizeof(cl_float), a, 0, NULL, NULL);
    if(err != CL_SUCCESS) {
        throw runtime_error("a_gpu copy from host failed");
    }

    err = clSetKernelArg(kernel, 0, sizeof(a_gpu), &a_gpu);
    if(err != CL_SUCCESS) {
        throw runtime_error("arg0 failed");
    }
    err = clSetKernelArg(kernel, 1, sizeof(cl_long), &offset);
    if(err != CL_SUCCESS) {
        throw runtime_error("arg1 failed");
    }
    err = clSetKernelArg(kernel, 2, sizeof(cl_int), &pos);
    if(err != CL_SUCCESS) {
        throw runtime_error("arg2 failed");
    }
    err = clSetKernelArg(kernel, 3, sizeof(cl_float), &value);
    if(err != CL_SUCCESS) {
        throw runtime_error("arg3 failed");
    }

    size_t global_ws[3];
    size_t local_ws[3];
    for(int i = 0; i < 3; i++) {
        global_ws[i] = 0;
        local_ws[i] = 0;
    }
    int workgroupSize = 32;
    int numWorkgroups = (N + workgroupSize - 1) / workgroupSize;
    int globalSize = numWorkgroups * workgroupSize;
    global_ws[0] = globalSize;
    local_ws[0] = workgroupSize;
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_ws, local_ws, 0, NULL, 0);
    if(err != CL_SUCCESS) {
        cout << "err" << err << endl;
        throw runtime_error("clenqueuendrangekernel failed");
    }
    cout << "queued kernel ok" << endl;

   err = clEnqueueReadBuffer(queue, a_gpu, CL_TRUE, 0,
                                         sizeof(cl_float) * N, a, 0, NULL, NULL);

    err = clFinish(queue);
    if(err != CL_SUCCESS) {
        throw runtime_error("finish failed");
    }
    cout << "clfinish finished ok" << endl;

    for(int i = 0; i < 5; i++) {
        cout << "a[" << i << "]=" << a[i] << endl;
    }

    return 0;
}

runcoclissue3a.sh

#!/bin/bash

set -x
set -e

mkdir -p build
clang++-3.8 -c -o build/cocl-issue3-a.o cocl-issue3-a.cpp -std=c++11
g++ -o build/cocl-issue3-a build/cocl-issue3-a.o -lOpenCL
build/cocl-issue3-a