cancel
Showing results for 
Search instead for 
Did you mean: 

OpenCL

drallan
Challenger

Very simple kernel crashes compiler

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

0 Likes
Reply
13 Replies
antzrhere
Adept III

Very simple kernel crashes compiler

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) 

 

0 Likes
Reply
antzrhere
Adept III

Very simple kernel crashes compiler

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)

0 Likes
Reply
dovalec
Staff
Staff

Re: Very simple kernel crashes compiler

Hi

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

Thx Dov

0 Likes
Reply
MicahVillmow
Staff
Staff

Re: Very simple kernel crashes compiler

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

0 Likes
Reply
drallan
Challenger

Re: Very simple kernel crashes compiler

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

0 Likes
Reply
MicahVillmow
Staff
Staff

Re: Very simple kernel crashes compiler

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.

0 Likes
Reply
drallan
Challenger

Re: Very simple kernel crashes compiler

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

0 Likes
Reply
MicahVillmow
Staff
Staff

Re: Very simple kernel crashes compiler

drallan,

Thank you very much for the help!

0 Likes
Reply
Meteorhead
Challenger

Re: Very simple kernel crashes compiler

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.

0 Likes
Reply