[Beignet] Bug in fmod
Zhigang Gong
zhigang.gong at linux.intel.com
Thu Feb 12 20:44:42 PST 2015
On Tue, Feb 10, 2015 at 01:43:47PM -0500, Pavan Yalamanchili wrote:
> Hi Zhigang,
>
> Running with OCL_STIRCT_CONFORMANCE=1 fixed this particular issues.
> The rest of the tests that are still failing. It will require a little
> bit more work on my end to narrow down the other problems.
One reminder here, for those errors, you may check whether there is a GPU
hang. For more detail of how to check the GPU hang and how to disable the
hang check if there is, you can check the "Known Issues" section in the
following wiki page or just the README.md.
http://www.freedesktop.org/wiki/Software/Beignet/
Thanks,
Zhigang Gong.
>
> --
> Pavan
>
> On Mon, Feb 9, 2015 at 7:42 PM, Zhigang Gong
> <zhigang.gong at linux.intel.com> wrote:
> > Hi Pavan,
> >
> > Could you run all of your test cases with the strict conformance mode as below?
> >
> > OCL_STRICT_CONFORMANCE=1 ./test_cases
> >
> > Thanks,
> > Zhigang Gong.
> >
> > On Mon, Feb 09, 2015 at 06:38:32PM -0500, Pavan Yalamanchili wrote:
> >> Hi Everyone,
> >>
> >> We have recently began testing our library (
> >> https://github.com/arrayfire/arrayfire) with Beignet. About 10 of our 44
> >> tests are failing for this architecture. I am trying to narrow down the
> >> cause of all the issues.
> >>
> >> Today I was able to narrow down the first, simplest issue. I am attaching a
> >> stand alone file that should reproduce the problem. The output of "fmod"
> >> seems to be incorrect for atleast some values.
> >>
> >> System information:
> >>
> >> Device: Intel(R) HD Graphics IvyBridge GT2
> >> Beignet version: Release_v1.0.0
> >> Linux Kernel: 3.18.4-1-ARCH
> >>
> >>
> >> The source code to reproduce the problem is attached below. Let me know if
> >> further information is needed from my end.
> >>
> >> --
> >> Pavan
> >
> >> #include <stdio.h>
> >> #include <CL/cl.h>
> >> #include <math.h>
> >>
> >> #define CL_CHECK(fn) do { \
> >> cl_int _err = fn; \
> >> if (_err == CL_SUCCESS) break; \
> >> printf("%s(%d) OpenCL Error: %d\n", \
> >> __FILE__, __LINE__, _err); \
> >> return _err; \
> >> } while (0) \
> >>
> >> const char fmod_ker_str[] = ""
> >> "__kernel\n"
> >> "void fmod_ker(__global const float * const lhs,\n"
> >> "float rhs,\n"
> >> "__global float * const out) {\n"
> >> "unsigned idx = get_global_id(0);\n"
> >> "out[idx] = fmod(lhs[idx], rhs);\n"
> >> "}\n";
> >>
> >>
> >> int main(int argc, char **args)
> >> {
> >> int pid = argc < 2 ? 0 : atoi(args[1]);
> >> int did = argc < 3 ? 0 : atoi(args[2]);
> >>
> >> unsigned num_platforms = 0;
> >> CL_CHECK(clGetPlatformIDs(0, NULL, &num_platforms));
> >> if (num_platforms <= pid) {
> >> printf("Invalid Platform ID\n");
> >> return -pid;
> >> }
> >>
> >> cl_platform_id *platforms = (cl_platform_id *)malloc(num_platforms * sizeof(cl_platform_id));
> >> CL_CHECK(clGetPlatformIDs(num_platforms, platforms, NULL));
> >> cl_platform_id platform = platforms[pid];
> >>
> >> unsigned num_devices = 0;
> >> CL_CHECK(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices));
> >> if (num_devices <= did) {
> >> printf("Invalid Device ID\n");
> >> return -did;
> >> }
> >>
> >> cl_device_id *devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id));
> >> CL_CHECK(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, devices, NULL));
> >> cl_device_id device = devices[did];
> >>
> >> char platform_name[64];
> >> char device_name[256];
> >>
> >> CL_CHECK(clGetPlatformInfo(platform,
> >> CL_PLATFORM_NAME,
> >> sizeof(platform_name), &platform_name,
> >> NULL));
> >>
> >>
> >> CL_CHECK(clGetDeviceInfo(device,
> >> CL_DEVICE_NAME,
> >> sizeof(device_name), &device_name,
> >> NULL));
> >>
> >> printf("Platform: %s\n", platform_name);
> >> printf("Device: %s\n", device_name);
> >> printf("Kernel String:\n%s\n", fmod_ker_str);
> >>
> >> cl_context_properties cps[3] = {CL_CONTEXT_PLATFORM,
> >> (cl_context_properties)(platform),
> >> 0};
> >>
> >> cl_int err;
> >> cl_context ctx = clCreateContext(cps, 1, &device, NULL, NULL, &err);
> >> CL_CHECK(err);
> >>
> >> cl_command_queue queue = clCreateCommandQueue(ctx, device, 0, &err);
> >> CL_CHECK(err);
> >>
> >> const int elements = 256;
> >> int bytes = elements * sizeof(float);
> >> float *h_in = (float *)malloc(bytes);
> >> float *h_out = (float *)malloc(bytes);
> >>
> >> for (int i = 0; i < elements; i++) {
> >> h_in[i] = 3.0;
> >> h_out[i] = 0.3; //10101;
> >> }
> >>
> >> cl_mem b_in = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
> >> CL_CHECK(err);
> >> CL_CHECK(clEnqueueWriteBuffer(queue, b_in, true, 0, bytes, h_in, 0, NULL, NULL));
> >>
> >> cl_mem b_out = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
> >> CL_CHECK(err);
> >>
> >> float val = 0.3;
> >>
> >> size_t local = elements;
> >> size_t global = elements;
> >>
> >> const char *ker_sources[] = {fmod_ker_str, NULL};
> >> const size_t ker_lens[] = { sizeof(fmod_ker_str), 0 };
> >> cl_program program = clCreateProgramWithSource(ctx, 1, ker_sources,
> >> ker_lens, &err);
> >> CL_CHECK(err);
> >> CL_CHECK(clBuildProgram(program, 0, NULL, "", NULL, NULL));
> >>
> >> cl_kernel kernel = clCreateKernel(program, "fmod_ker", &err);
> >> CL_CHECK(err);
> >>
> >> CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &b_in ));
> >> CL_CHECK(clSetKernelArg(kernel, 1, sizeof(float ), &val ));
> >> CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &b_out));
> >>
> >> CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, NULL));
> >> CL_CHECK(clFlush(queue));
> >> CL_CHECK(clFinish(queue));
> >>
> >> CL_CHECK(clEnqueueReadBuffer(queue, b_out, true, 0, bytes, h_out, 0, NULL, NULL));
> >>
> >> for (int i = 0; i < 10; i++) {
> >> printf("Expected: %f, Calculated: %f\n", fmodf(h_in[i], val), h_out[i]);
> >> }
> >>
> >> free(platforms);
> >> free(devices);
> >> free(h_in);
> >> free(h_out);
> >> }
> >
> >> _______________________________________________
> >> Beignet mailing list
> >> Beignet at lists.freedesktop.org
> >> http://lists.freedesktop.org/mailman/listinfo/beignet
> >
More information about the Beignet
mailing list