13 Replies Latest reply on Nov 19, 2012 7:28 PM by ticktock

    Very simple kernel crashes compiler

    drallan

      Crash occurs in application code and Kernel Analyzer

       

       

      It's a bit suprising, but the simple one line kernel in the code window crashs during compile in both application code and in the Kernel Analyzer. Windows crash info is attached. I am using Windows 7 x64 with SDK 2.6 cal version 11.12 (upgraded from 2.4 and 2.5).

      Not sure what the cause is. My system is relatively stable. Changing the exact form of the code can prevent the crash, e.g., a larger loop counter or a slighly more complex expression. The problem might occur when the compiler tries to unroll a long loop.

      Any feedback is appreciated.

      Thanks.

       

      //Kernel crashes aticaldd.bin in Kernel Analyzer using cal 11.12, 11.11 and earlier versions 
      //APP SDK 2.6, cal 11.12, Windows 7 x64 
      //-------------------------------------------------------------------------------------------------------
       __kernel void systest(__global uint *A) 
      { 
          int gid,gx=get_global_id(0),gy=get_global_id(1);
          uint i; gid=64*gy+gx;
      
         for(i=0;i<0x8000;i++)A[(i<<10)]=5; 
      }
      //------------------------------------------------------------------------------------------------------- 
      //Other forms of the loop may or may not crash: 
      //NOTE: This line prevents crash: for(i=0;i<0x20000;i++)A[(i<<10)]=5; 
      //NOTE: This line prevents crash: for(i=0;i<0x8000;i++)A[((i<<10)+gid)&0xfffffffc]=5; 
      
      /* Crash Report from windows 
          Problem Event Name: APPCRASH
         Application Name: AMDAPPKernelAnalyzer.exe
         Application Version: 1.10.0.1149
         Application Timestamp: 4ec4bb5b 
         Fault Module Name: aticaldd.dll 
         Fault Module Version: 6.14.10.1646 
         Fault Module Timestamp: 4ebb3717 
         Exception Code: c0000005 
         Exception Offset: 00241e54 
         OS Version: 6.1.7601.2.1.0.256.1 
         Locale ID 33 
         Additional Information 1: 0a9e 
         Additional Information 2: 0a9e372d3b4ad19135b953a78882e789 
         Additional Information 3: 0a9e 
         Additional Information 4: 0a9e372d3b4ad19135b953a78882e789 */
      

       

      Message edited to reformat attachment

        • Very simple kernel crashes compiler
          antzrhere

          I get something similar - haven't tested runtime, but kernelanalyzer spits out "Unknown exception. Please save your shader & restart AMD APP KernelAnalyzer."

           

          This is a problem for 5870 but some of the older 4xxx series seem to be able to output code...but what target ISA object code works seems to change randomly (i managed to get 5870 to compile once, then the next time i pressed the button it threw an error) 

           

          • Re: Very simple kernel crashes compiler
            MicahVillmow

            drallan,

            I am not able to reproduce this isuse, even with forced unrolling. Most likely this is an issue that was fixed in our upcoming release, can you try Catalyst 12.1 preview, http://support.amd.com/us/kbarticles/Pages/Catalyst121Previewdriver.aspx, to see if it fixes your issue?

             

            Thanks,

            Micah

            • Re: Very simple kernel crashes compiler
              MicahVillmow

              drallan,

              We are trying to fix a possible issue with the forums, would it be possible to attach as a file the source code you posted in the original email and also explain how you added it to your post? The source code should have been formatted nicely, instead of put on a single line.

                • Re: Very simple kernel crashes compiler
                  drallan

                  MicahVillmow wrote:

                   

                  drallan,

                  We are trying to fix a possible issue with the forums, would it be possible to attach as a file the source code you posted in the original email and also explain how you added it to your post? The source code should have been formatted nicely, instead of put on a single line.

                   

                  Hi Micah,

                   

                  I'm not sure I have done exactly as requested, but I have:

                   

                  1. Edited the code window in the original message to add lines and spacing to make it look good.

                  2. Attached the original source file as is to this message named systest_bug.cl.

                  3. The original code was added in the old forum using whatever method was made easily available

                      to the user, either "insert file as code window" or "insert code window" followed by a manual WinXP

                     cut and paste.The original file contained 0x0D0A (crlf), 0x20 (space), and 0x09 (tab) sequences,

                     all of which were changed to single spaces during the migration to the new forum.

                   

                  Allan

                • Re: Very simple kernel crashes compiler
                  Meteorhead

                  Let me reuse this thread for another compiler crash issue. This kernel used to work, and now it crashes compilers 12.8 through 12.11beta. I have tried commenting out differnt parts of the kernel to see which command triggers the crash, and it seems that when *state = ...; is which causes the problem. Naturally I need this update to keep the PRNG running, so this is a must. Please, give me some advice how to find a way around that does not cause the compiler to crash.

                    • Re: Very simple kernel crashes compiler
                      drallan

                      Meteorhead wrote:

                       

                      Let me reuse this thread for another compiler crash issue. This kernel used to work, and now it crashes compilers 12.8 through 12.11beta. I have tried commenting out differnt parts of the kernel to see which command triggers the crash, and it seems that when *state = ...; is which causes the problem. Naturally I need this update to keep the PRNG running, so this is a must. Please, give me some advice how to find a way around that does not cause the compiler to crash.

                       

                      Most definitely a compiler bug. It crashes writing c to the upper dword of (uint2)  *state but the problem occurs because the compiler gets confused by the expression (x<c) in the RNG.  Most any other expression works but not the reverse (c>x) so you might  use something like:

                        c = hi + (x<=c);

                         or c = hi + (x<=c && x!= c);   which is logically the same with the original expression.

                       

                      Versions 12_8 through 12_10 work fine for me and only 12_11 crashes. If you re-installed 12_8 - 12_10 after installing 12_11, parts of 12_11 can remain, which a new install can't remove, including the compiler. That happen to me and I had to manually remove several of the driver files.

                       

                      Interestingly, when it crashes it dumps piles of LLVM debug code that I've never seen before so maybe it's a new LLVM compiler. That might explain why the 12_11 "Don't Settle"  drivers are significantly faster that older drivers for gcn.   ...... when it works.  .

                       

                      LLVM ERROR: Cannot select: 0x2b7cb580: i8 = setcc 0x2b0aca40,0x2b0ac840, 0x2a07f590 [ID=102]

                        0x2b0aca40: i32 = AMDILISD::ADD 0x2b0ac840, 0x2a07fc90 [ID=100]

                          0x2b0ac840: i32 = mul 0x2b0ae660, 0x2b0b2eb0 [ORD=44] [ID=95]

                              0x2b0b36b0: i32 = mul 0x2b7cb880, 0x2b0b2eb0 [ORD=35] [ID=86]

                                0x2b7cb880: i32 = AMDILISD::ADD 0x2b0ac540, 0x2b0ac440 [ID=81]

                                 0x2b0b2eb0: i32 = Constant<-83941> [ORD=16] [ID=4]

                       

                      drallan

                        • Re: Very simple kernel crashes compiler
                          Meteorhead

                          Thanks for your reply drallan. I went and tried out 12.11 beta7 to see if perhaps this issue has been addressed, but no. Infact practically no OpenCL program works that I've developed. There are two running projects, and neither of them compile. These beta drivers sure accelerate games, but they sure wreck havoc in OpenCL.

                            • Re: Very simple kernel crashes compiler
                              Meteorhead

                              Can someone tell me how to effectively roll back AMD drivers on a Win 8 system, because I cannot develop OpenCL programs since roughly 2 weeks now and it DRIVING ME NUTS! I have tried uninstaller programs aiming at removing ATi drivers after regular uninstallation, I have also tried

                               

                              http://www.brightsideofnews.com/news/2012/4/4/guide-how-to-completely-uninstall-amd-graphics-drivers.aspx?pageid=1

                               

                              to no avail. Is it really that complicated to wipe a driver off of a Windows machine? After uninstalling ALL AMD software from the computer, I uninstall the driver by force from Device Manager, and after a restart, the long awaited SW rendered 800*600 resolution awaits me, but before I could install Catalyst, Windows finds some driver named ATi FirePro M7820, which effectively renders desktop. Even after having Catalyst 12.8 installed on my machine, the same device string remains. I wouldn't care about the device string, but all kernel compilations fail. Kernel Analyzer crashes all my ongoing project kernel's compilations.

                               

                              Someone help me cause it's really becoming uncool that reverting drivers is so freaking hard.

                               

                              Just out of curiosity: why do so many registry entries remain after a full uninstall?

                                • Re: Very simple kernel crashes compiler
                                  ticktock

                                  We went rounds with removing the 12.11 Beta driver, but i found 2 ways to completely remove it in order to revert to 12.8.  On one Windows 7 x64 system, without an SSD, I was able to use a restore point from before I installed the 12.11 Beta.  On my Windows 8 system I used the steps I outlined here:

                                   

                                  http://www.overclock.net/t/1323729/amd-gpu-folding-on-12-11-beta-drivers/0_50

                                   

                                  Windows 8 will still run some sort of an AMD driver, but you should be able to install the driver of your choice without any remnants of the 12.11 Beta drivers.

                                   

                                  I hope this helps, and I have to agree that the beta drivers are great for gaming, but any compute tasks are totally broken with the beta OpenCL version.