OpenCL

[s23ba32]

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;

}

[himanshu_gautam]

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.

[AlexMcLeod]

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

[gordonei]

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.

[endoerner]

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

[sudarshan]

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,

[endoerner]

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 Version	13.251-131206a-166389C-ATI
Catalyst Version	13.12
Provider	Advanced Micro Devices, Inc.
2D Driver Version	8.01.01.1360
2D Driver File Path	/REGISTRY/MACHINE/SYSTEM/ControlSet001/Control/CLASS/{4D36E968-E325-11CE-BFC1-08002BE10318}/0000
Direct3D Version	9.14.10.01001
OpenGL Version	6.14.10.12618
AMD Catalyst Control Center Version	2013.1206.1603.28764
AMD Audio Driver Version	7.12.0.7717

Thanks for your help!

[dipak]

Hi,

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

Thanks,

Dipak

[dipak]

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

[endoerner]

Dear dipak,

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

Best regards!