13 Replies Latest reply: Jul 8, 2014 12:11 PM by endoerner RSS

LLVM error when compiling kernel

s23ba32 Newbie
Currently Being Moderated

I am getting this error message:

LLVM ERROR: Cannot select: 0xb1e670: i8 = setcc 0xa94230, 0xf72f00, 0xae22a0 [ID=25]

  0xa94230: i32 = AMDILISD::ADD 0xf72f00, 0xf6f7d0 [ID=22]

    0xf72f00: i32 = mul 0xb30210, 0xf99250 [ORD=831] [ID=19]

      0xb30210: i32,ch = CopyFromReg 0xbbea30, 0xb71800 [ORD=830] [ID=14]

        0xb71800: i32 = Register %vreg41 [ORD=830] [ID=1]

      0xf99250: i32 = Constant<-83941> [ORD=831] [ID=3]

    0xf6f7d0: i32,ch = CopyFromReg 0xbbea30, 0xae37b0 [ORD=830] [ID=15]

      0xae37b0: i32 = Register %vreg42 [ORD=830] [ID=2]

  0xf72f00: i32 = mul 0xb30210, 0xf99250 [ORD=831] [ID=19]

    0xb30210: i32,ch = CopyFromReg 0xbbea30, 0xb71800 [ORD=830] [ID=14]

      0xb71800: i32 = Register %vreg41 [ORD=830] [ID=1]

    0xf99250: i32 = Constant<-83941> [ORD=831] [ID=3]

 

This is the code that is triggering the error:

typedef struct{ uint x; uint c; } mwc64x_state_t;

 

enum{ MWC64X_A = 4294883355U };

enum{ MWC64X_M = 18446383549859758079UL };

 

void MWC64X_Step(mwc64x_state_t *s)

{

    uint X=s->x, C=s->c;

 

    uint Xn=MWC64X_A*X+C;

    uint carry=(uint)(Xn<C);

                     

    // Replacing this line with a constant makes the error go away.  

    uint Cn=mad_hi(MWC64X_A,X,carry);

 

    s->x=Xn;

 

    // As does replacing this assignment with a constant.

    s->c=Cn;

}

  • Re: LLVM error when compiling kernel
    himanshu.gautam Master
    Currently Being Moderated

    Hi,

    Can you let us your system setup details: CPU, GPU, SDK , Driver, OS (Window/Linux) (32/64)

    I copied your function MXC64X_step inside a opencl sample, and ran the sample. So I assume, I am able to compile your code with out any issues. I used SDK 2.8, Win8 64 bit, 13.1 driver, AMD A10 5800K APU with HD 7660D iGPU.

    Regards

    Himanshu , Bruhaspati

    --------------------------------

    The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied

    • Re: LLVM error when compiling kernel
      AlexMcLeod Newbie
      Currently Being Moderated

      I realise this is an old thread, but this issue still exists with APP SDK 2.9.

       

      MWC64X is a useful random number generator, and it works with earlier versions of the APP SDK - but it triggers an LLVM error in the current version. It comes from here:

      http://cas.ee.ic.ac.uk/people/dt10/research/rngs-gpu-mwc64x.html

       

      Building my program with -cl-opt-disable fixes the error and the program runs, but overall performance is reduced. It'd be really nice if MWC64X wasn't broken by the system's attempts to optimise it.

       

      My hardware:

      CPU: FX-8350

      GPU: Radeon R9 290

      APP SDK: 2.9

      OS: Windows 8.1, 64-bit

      • Re: LLVM error when compiling kernel
        gordonei Newbie
        Currently Being Moderated

        Ja, also seeing this issue with Firepro W5000, SDK 2.9, running on CentOS 6.3.

         

        Using the -cl-opt-disable flag results in a x20 performance hit for my code (the runtime is in the order of seconds, so this really hurts). I've found a better solution is to use the CPU to seed the Random Number Generator (i.e. create a separate context, device, command queue, etc.). The RNG state struct is only 2 uints, so if you overlap your communication and compute carefully, then you can still use the MWC RNG. Its a pain, but depending on your application it might be worth it.

  • Re: LLVM error when compiling kernel
    endoerner Newbie
    Currently Being Moderated

    Hi, I am sorry to revive this old post, but I (still) have the same problem... but in my case I have also problems in the function MWC_SeedImpl_Mod64:

     

    uint2 MWC_SeedImpl_Mod64(ulong A, ulong M, uint vecSize, uint vecOffset, ulong streamBase, ulong streamGap)

    {

      enum{ MWC_BASEID = 4077358422479273989UL };

     

      ulong dist=streamBase + (get_global_id(0)*vecSize+vecOffset)*streamGap;

      ulong m=MWC_PowMod64(A, dist, M);

     

      ulong x=MWC_MulMod64(MWC_BASEID, m, M);

      return (uint2)((uint)(x/A), (uint)(x%A)); // ED: this line generates the error... commenting it "fixes" the problem

    }

     

    I attached the code, if someone is able to run it without problems... I am really stuck with this and I would not like to throw away my $2000 AMD PC and buy an

    intel CPU just to compile this kernel.... the RNG is very nice and I would like to be able to use it in other projects...

     

    PS: I added the absolute path to the kernel files on clBuildProgram() (C++ Wrapper).

    • Re: LLVM error when compiling kernel
      sudarshan Moderator
      Currently Being Moderated

      Hi,

      We would try to compile and run the code. Meanwhile it would be of great help if you could give us rest of the system parameters (OS, driver and SDK version).

       

      Thanks,

      -Sudarshan

      ----------------------------------------

      The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

      • Re: Re: LLVM error when compiling kernel
        endoerner Newbie
        Currently Being Moderated

        Hi sudarshan... I am using the AMD APP SDK v2.9 integrated under Visual Studio 2012, Windows 7 Professional 64-bit

         

        I always forget to put my system specs, here they come (from CodeXL menu inside VS2012, main settings)

         

        System

        Operating System Version (name), Windows 7 Professional, 64bit

        Operating System Version (number), 6.1.7601

        Number Of Processors, 6

        System Type, x86

        CPU Family, Family 0x10, Model 0xa, Stepping 0x0

        Total Physical Memory, 16383 MB

         

        OpenCL Platforms

        Vendor, Advanced Micro Devices, Inc.

        Name, AMD Accelerated Parallel Processing

        Profile, FULL_PROFILE

        Version, OpenCL 1.2 AMD-APP (1348.5)

        Extensions, cl_khr_icd, cl_amd_event_callback, cl_amd_offline_devices, cl_khr_d3d10_sharing, cl_khr_d3d11_sharing, cl_khr_dx9_media_sharing

         

        OpenCL Devices

        Platform ID, 1, 1, 1

        Device Type, GPU, GPU, CPU

        Device Name, Tahiti (AMD Radeon HD7970), Tahiti (AMD Radeon HD7970), AMD Phenom(tm) II X6 1100T Processor

        (...)

         

        Software information (Catalyst Control Center)

        Driver Packaging Version13.251-131206a-166389C-ATI
        Catalyst Version13.12
        ProviderAdvanced Micro Devices, Inc.
        2D Driver Version8.01.01.1360
        2D Driver File Path/REGISTRY/MACHINE/SYSTEM/ControlSet001/Control/CLASS/{4D36E968-E325-11CE-BFC1-08002BE10318}/0000
        Direct3D Version9.14.10.01001
        OpenGL Version6.14.10.12618
        AMD Catalyst Control Center Version2013.1206.1603.28764
        AMD Audio Driver Version7.12.0.7717

         

        Thanks for your help!

        • Re: LLVM error when compiling kernel
          dipak Moderator
          Currently Being Moderated

          Hi,

          We are able to reproduce the error. We're looking into it and will give you the update shortly.

           

          Thanks,

          Dipak

          - Dipak

          ----------------------------------------

          The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

          • Re: LLVM error when compiling kernel
            dipak Moderator
            Currently Being Moderated

            Hi endoerner,

            Here is the update.

            We're able to write a sample code which captures similar scenario and generates the same compilation error. We've identified the possible block of codes which is causing the error. We're looking into it in details and if find it's a really compiler optimization issue, we'll forward the bug to our compilation team.

             

            Thanks,

            Dipak

            - Dipak

            ----------------------------------------

            The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

            • Re: LLVM error when compiling kernel
              endoerner Newbie
              Currently Being Moderated

              Dear dipak,

               

              thanks for your help, I hope that it will be found the source (and hopefully the solution) to this issue.

               

              Best regards!

              • Re: LLVM error when compiling kernel
                dipak Moderator
                Currently Being Moderated

                Hi endoerner,

                 

                An internal bug report has been filed to the compiler team, and we will keep you updated.

                 

                Regards,

                - Dipak

                ----------------------------------------

                The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

                • Re: LLVM error when compiling kernel
                  endoerner Newbie
                  Currently Being Moderated

                  Hi dipak,

                   

                  have had been any progress on this matter?... I would like to ask if this is a problem with the gpu driver or the amd sdk?... because maybe one could try to install an older driver/sdk in the meanwhile...

                   

                  thanks a lot for your help!

                  • Re: LLVM error when compiling kernel
                    dipak Moderator
                    Currently Being Moderated

                    Hi endoerner,

                     

                    Sorry, I forgot to give you the update.

                    As per the latest status report, the reported bug has been successfully reproduced by the verification team with the latest driver (an internal one) and it'll be forwarded to corresponding dev team.

                    AFAIK, this problem is related to driver (more specifically compiler), not SDK. So, if you want to try out some older versions, you can but I'm not sure. Please let us know if any older version works for you.

                     

                    Regards,

                    - Dipak

                    ----------------------------------------

                    The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

  • Re: LLVM error when compiling kernel
    endoerner Newbie
    Currently Being Moderated

    To all that may be interested:

     

    I realized several tests with different AMD drivers and OpenCL SDK suites and I was able to run again this RNG without problems, the hardware was:

     

    CPU: AMD Phenom II X6 1100T

    GPU: AMD Radeon HD 7970 (Tahiti)

    OS: Windows 7 Professional, x64

    OpenCL SDK´s : AMD APP v2.9 and Intel OpenCL SDK 2014 beta + Intel CPU only runtime 2013 (yup, I was able to install the Intel SDK on a AMD only machine... hehehe, you need the CPU only runtime first)

    IDE: VS2012

     

    I tested different driver versions, from 12.1 up to 12.10:

    * For AMD Catalyst 12.10 the mwc64x RNG works flawlessly as used to be...

    * from 12.1 to 12.8 the GPU is not able to execute this piece of code:

     

    uint2 MWC_SeedImpl_Mod64(ulong A, ulong M, uint vecSize, uint vecOffset, ulong streamBase, ulong streamGap)

    {

      enum{ MWC_BASEID = 4077358422479273989UL };

     

      ulong dist=streamBase + (get_global_id(0)*vecSize+vecOffset)*streamGap;

      ulong m=MWC_PowMod64(A, dist, M);

     

      ulong x=MWC_MulMod64(MWC_BASEID, m, M);

      return (uint2)((uint)(x/A), (uint)(x%A)); // ED: this line triggers an exception

    }

     

    it seems that the older drivers are not able to handle the division or modulo operation between two ulong constants, therefore one has to convert x and A constants to uint before the operations.

     

    So as a resume the winner combination is:

    AMD Catalyst v12.10

    AMD APP SDK v2.9

     

    I did not adventured installing more recent drivers, I did not want to take the risk to have again the execution error, and, at least for me, this drivers are not sooo old... if someone wants to test be welcome. At least in the meanwhile I can work with this RNG...

More Like This

Legend

  • Correct Answers - 4 points
  • Helpful Answers - 2 points