Пример #1
0
        private static long PoissonRBSOR(Cl.Device device, Cl.Context context, Cl.Program program, bool lmem,
                                         float x0, float y0, float x1, float y1, int dimX, int dimY, int N, float omega,
                                         string fileName = null, string options = "")
        {
            Cl.ErrorCode error;
            Cl.Event     clevent;

            // build program
            clSafeCall(Cl.BuildProgram(program, 1, new[] { device }, options, null, IntPtr.Zero));
            Cl.BuildStatus status = Cl.GetProgramBuildInfo(program, device, Cl.ProgramBuildInfo.Status, out error).CastTo <Cl.BuildStatus>();
            if (status != Cl.BuildStatus.Success)
            {
                throw new Exception(status.ToString());
            }

            // save binary
            if (fileName != null)
            {
                Cl.InfoBuffer binarySizes = Cl.GetProgramInfo(program, Cl.ProgramInfo.BinarySizes, out error);
                clSafeCall(error);
                Cl.InfoBufferArray binaries = new Cl.InfoBufferArray(
                    binarySizes.CastToEnumerable <IntPtr>(Enumerable.Range(0, 1)).Select(sz => new Cl.InfoBuffer(sz)).ToArray());
                IntPtr szRet;
                clSafeCall(Cl.GetProgramInfo(program, Cl.ProgramInfo.Binaries, binaries.Size, binaries, out szRet));
                byte[] binary = binaries[0].CastToArray <byte>(binarySizes.CastTo <IntPtr>(0).ToInt32());
                File.WriteAllBytes(fileName, binary);
            }

            // create kernel
            Cl.Kernel kernel = Cl.CreateKernel(program, "PoissonRBSOR" + (lmem ? "_LMem" : ""), out error);
            clSafeCall(error);

            // create command queue
            Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(context, device, Cl.CommandQueueProperties.None, out error);
            clSafeCall(error);

            float hx = (x1 - x0) / dimX;
            float hy = (y1 - y0) / dimY;

            // boundary values

            float[] hgrid = new float[dimX * dimY];

            int gstride = dimX;

            for (int i = 1; i < dimY - 1; i++)
            {
                int   y_idx = i * gstride;
                float y_val = y0 + i * hy;
                hgrid[y_idx]            = u(x0, y_val);
                hgrid[y_idx + dimX - 1] = u(x0 + (dimX - 1) * hx, y_val);
            }

            for (int j = 1; j < dimX - 1; j++)
            {
                float x_val = x0 + j * hx;
                hgrid[j] = u(x_val, y0);
                hgrid[j + (dimY - 1) * gstride] = u(x_val, y0 + (dimY - 1) * hy);
            }

            // laplacian values

            float[] hlaplacian = new float[(dimX - 2) * (dimY - 2)];

            int lstride = dimX - 2;

            for (int i = 1; i < dimY - 1; i++)
            {
                for (int j = 1; j < dimX - 1; j++)
                {
                    hlaplacian[j - 1 + (i - 1) * lstride] = J(x0 + j * hx, y0 + i * hy);
                }
            }

            // allocate device vectors
            Cl.Mem dgrid = Cl.CreateBuffer(context, Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadWrite,
                                           (IntPtr)(sizeof(float) * hgrid.Length), hgrid, out error);
            clSafeCall(error);
            Cl.Mem dlaplacian = Cl.CreateBuffer(context, Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadOnly,
                                                (IntPtr)(sizeof(float) * hlaplacian.Length), hlaplacian, out error);
            clSafeCall(error);

            // setup kernel arguments
            clSafeCall(Cl.SetKernelArg(kernel, 0, dgrid));
            clSafeCall(Cl.SetKernelArg(kernel, 1, dlaplacian));
            clSafeCall(Cl.SetKernelArg(kernel, 2, dimX));
            clSafeCall(Cl.SetKernelArg(kernel, 3, dimY));
            clSafeCall(Cl.SetKernelArg(kernel, 4, gstride));
            clSafeCall(Cl.SetKernelArg(kernel, 5, lstride));
            clSafeCall(Cl.SetKernelArg(kernel, 6, hx));
            clSafeCall(Cl.SetKernelArg(kernel, 7, hy));
            clSafeCall(Cl.SetKernelArg(kernel, 8, omega));
            if (lmem)
            {
                clSafeCall(Cl.SetKernelArg(kernel, 10, (AREA_SIZE_Y + 2) * (AREA_SIZE_X + 2) * sizeof(float), null));
            }

            IntPtr[] lo = { (IntPtr)TILE_SIZE_X, (IntPtr)TILE_SIZE_Y };
            IntPtr[] gl =
            {
                (IntPtr)((dimX - 2 + (lmem ? AREA_SIZE_X : TILE_SIZE_X) - 1) /
                         (lmem ? AREA_SIZE_X : TILE_SIZE_X) * TILE_SIZE_X),
                (IntPtr)((dimY - 2 + (lmem ? AREA_SIZE_Y : TILE_SIZE_Y) - 1) /
                         (lmem ? AREA_SIZE_Y : TILE_SIZE_Y) * TILE_SIZE_Y)
            };

            // execute RED kernel
            clSafeCall(Cl.SetKernelArg(kernel, 9, 1));
            clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));

            // execute BLACK kernel
            clSafeCall(Cl.SetKernelArg(kernel, 9, 0));
            clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));

            clSafeCall(Cl.Finish(cmdQueue));

            Stopwatch stopwatch = new Stopwatch();

            stopwatch.Start();

            for (int idx = 1; idx < N; idx++)
            {
                // execute RED kernel
                clSafeCall(Cl.SetKernelArg(kernel, 9, 1));
                clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));

                // execute BLACK kernel
                clSafeCall(Cl.SetKernelArg(kernel, 9, 0));
                clSafeCall(Cl.EnqueueNDRangeKernel(cmdQueue, kernel, 2, null, gl, lo, 0, null, out clevent));
            }

            clSafeCall(Cl.Finish(cmdQueue));

            stopwatch.Stop();

            // copy results from device back to host
            clSafeCall(Cl.EnqueueReadBuffer(cmdQueue, dgrid, Cl.Bool.True, IntPtr.Zero,
                                            (IntPtr)(sizeof(float) * hgrid.Length), hgrid, 0, null, out clevent));

            clSafeCall(Cl.Finish(cmdQueue));

            cmdQueue.Dispose();
            kernel.Dispose();
            dgrid.Dispose();

            float avgerr = 0, maxerr = 0;

            for (int i = 1; i < dimY - 1; i++)
            {
                for (int j = 1; j < dimX - 1; j++)
                {
                    float theory = u(x0 + j * hx, y0 + i * hy);
                    float err    = Math.Abs(theory - hgrid[j + i * gstride]) / Math.Abs(theory);
                    avgerr += err;
                    maxerr  = Math.Max(maxerr, err);
                }
            }
            avgerr /= dimX * dimY;

            long elapsedTime = stopwatch.ElapsedMilliseconds;

            Console.WriteLine("average error = {0}%\nmaximal error = {1}%\nelapsed time: {2}ms\niterations per second: {3}",
                              avgerr * 100, maxerr * 100, elapsedTime, (double)N / (double)elapsedTime * 1000.0d);

            return(elapsedTime);
        }
Пример #2
0
        private static byte[] ToAMDClBinary(this Program irprog, Cl.Device device)
        {
            // create OpenCL stub before registers naming

            string clstub = string.Join("\n", irprog.Kernels.Select(kernel => string.Format(
                                                                        "__kernel void {0}({1}) {{ {2} }}",
                                                                        kernel.Name,
                                                                        string.Join(", ", kernel.FormalParameters.Select(
                                                                                        (fp, idx) => {
                switch (fp.StateSpace)
                {
                case StateSpaces.GLOBAL:
                    return(string.Format("__global {0}* p{1}", fp.UnderlyingType.FormatC(), idx));

                case StateSpaces.SHARED:
                    return(string.Format("__local {0}* p{1}", fp.UnderlyingType.FormatC(), idx));

                case StateSpaces.CONSTANT:
                    return(string.Format("__constant {0}* p{1}", fp.UnderlyingType.FormatC(), idx));

                case StateSpaces.REG:
                default:
                    return(string.Format("{0} p{1}", fp.UnderlyingType.FormatC(), idx));
                }
            })),
                                                                        string.Join(" ", kernel.FormalParameters.Where(
                                                                                        fp => fp.StateSpace != StateSpaces.REG && fp.StateSpace != StateSpaces.CONSTANT).Select(
                                                                                        fp => string.Format("*p{0} = 0;", kernel.FormalParameters.IndexOf(fp)))))));

            // create template binary from OpenCL stub

            Cl.ErrorCode error;

            Cl.Context context = Cl.CreateContext(null, 1, new[] { device }, null, IntPtr.Zero, out error);
            clSafeCall(error);

            Cl.Program program = Cl.CreateProgramWithSource(context, 1, new[] { clstub }, null, out error);
            clSafeCall(error);

            clSafeCall(Cl.BuildProgram(program, 1, new[] { device },
                                       "-fno-bin-source -fno-bin-llvmir -fno-bin-exe -fbin-amdil", null, IntPtr.Zero));
            Cl.BuildStatus status = Cl.GetProgramBuildInfo(program, device, Cl.ProgramBuildInfo.Status, out error).CastTo <Cl.BuildStatus>();
            if (status != Cl.BuildStatus.Success)
            {
                throw new Exception(status.ToString());
            }

            Cl.InfoBuffer binarySizes = Cl.GetProgramInfo(program, Cl.ProgramInfo.BinarySizes, out error);
            clSafeCall(error);
            Cl.InfoBufferArray binaries = new Cl.InfoBufferArray(
                binarySizes.CastToEnumerable <IntPtr>(Enumerable.Range(0, 1)).Select(sz => new Cl.InfoBuffer(sz)).ToArray());
            IntPtr szRet;

            clSafeCall(Cl.GetProgramInfo(program, Cl.ProgramInfo.Binaries, binaries.Size, binaries, out szRet));

            program.Dispose();
            context.Dispose();

            // inject generated code into the elf binary

            LinkingView   elf    = new LinkingView(binaries[0].CastToArray <byte>(binarySizes.CastTo <IntPtr>(0).ToInt32()));
            SymTabSection symtab = (SymTabSection)elf[".symtab"];
            Section       amdil  = elf[".amdil"];
            Section       rodata = elf[".rodata"];

            MemoryStream amdilcode = new MemoryStream();

            foreach (Kernel kernel in irprog.Kernels)
            {
                SymbolWrapper _metadata = symtab["__OpenCL_" + kernel.Name + "_metadata"];

                string[] str_metadata = Marshal.PtrToStringAnsi(Marshal.UnsafeAddrOfPinnedArrayElement(
                                                                    rodata.Data, (int)_metadata.st_value), (int)_metadata.st_size).Split('\n');

                int setup_id = (from line in str_metadata let prms = line.Split(':')
                                                                     where prms[0] == ";uniqueid" select int.Parse(prms[1])).Single();

                int raw_uav_id = (from line in str_metadata let prms = line.Split(':')
                                                                       where prms[0] == ";uavid" select int.Parse(prms[1])).Single();

                SymbolWrapper _fmetadata = symtab["__OpenCL_" + kernel.Name + "_fmetadata"];

                string[] str_fmetadata = Marshal.PtrToStringAnsi(Marshal.UnsafeAddrOfPinnedArrayElement(
                                                                     rodata.Data, (int)_fmetadata.st_value), (int)_fmetadata.st_size).Split('\n');

                int func_id = (from line in str_fmetadata let prms = line.Split(':')
                                                                     where prms[0] == ";uniqueid" select int.Parse(prms[1])).Single();

                // ugly, i know!!!
                raw_uav_id = Math.Max(raw_uav_id, 11);
                int arena_uav_id = raw_uav_id;

                byte[] code = Encoding.Convert(Encoding.Unicode, Encoding.ASCII, Encoding.Unicode.GetBytes(
                                                   irprog.ToAMDIL(kernel.Name, setup_id, func_id, raw_uav_id, arena_uav_id)));

                SymbolWrapper _amdil = symtab["__OpenCL_" + kernel.Name + "_amdil"];

                _amdil.st_value = (uint)amdilcode.Position;
                _amdil.st_size  = (uint)code.Length;

                foreach (byte b in code)
                {
                    amdilcode.WriteByte(b);
                }
            }

            amdil.Data = amdilcode.ToArray();

            return(elf.BuildBinary());
        }