Archives Discussions

[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

[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) 

[antzrhere]

only the line "for(i=0;i<0x8000;i++)A[(i<<10)]=5;" seems  to matter.

Changing the number of iterations seems to allow the compiler to partially unroll the loop (i.e. 0x7998 results in 24 iterations within a single loop, whereas 0x7999 results in no unrolling). I guess maybe your right about unrolling causing the crash - or ony my system it simply reports an error (same as you but using Cat 12.1a preview)

[dovalec]

Hi

Please contact Yousef ( Yousef.Shajrawi@amd.com ) regarding compiler issues.

Thx Dov

[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

[drallan]

Hi Micah,

Yes, it seems fixed in the latest version. The code compiles correctly when I select version 12.1 (preview) in the Kernel Analyzer, and still fails when I select the older versions 11.11 and earlier, which give an "Unknown Exception, please save and restart" error.

Many thanks,

Allan

[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.

[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

[MicahVillmow]

drallan,

Thank you very much for the help!

[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.

[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

[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.

[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...

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?

[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.