cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

chris9871
Journeyman III

CodeXLAnalyzer core dump

I have a kernel which causes opencl to core dump inside libamdocl64.so.  It was causing my program to crash so I narrowed down the problem to the compile step, and down to a single line.  I've reproduced this issue in a simple kernel:

typedef struct test_struct test_t;

typedef __global test_t * test_p;

struct test_struct {

    test_p next;

};

typedef struct btest_struct btest_t;

struct btest_struct {

    size_t sz;

    test_p dt[];

};

__kernel void compilerCoreDump(__global const int* in, int n) {

    int i = get_global_id(0);

    if (i >= n)

        return;

    int val = in;

    btest_t btt;

    test_p tp = btt.dt[val]; // <-- Compile core dump

}

I'm aware that the code above is nonsense - the actual usage is more complex but the (compile) result is the same.  I assume that the compiler should be reporting an error instead of core dumping, so I have two questions:

  1. What am I doing wrong and what can I do to work around the problem?
  2. Is there a fix in the pipe (or already released) for this?

My versions:

  • Ubuntu 14.04
  • /opt/AMD/CodeXL_1.4.5728/CodeXLAnalyzer --asic Hawaii  --versionCodeXL KernelAnalyzer CLI Version 1,4,5728,0 OpenCL 1.2 AMD-APP (1411.4)
  • fglrx-updates: 13.350.1  kernel-3.13.0-34-generic-x86_64

Thanks,

Chris...

Core Dump stack trace (bottom bit):

#28 0x00007ffff1d47fc4 in aclCompile () from /usr/lib/fglrx/libamdocl64.so

#29 0x00007ffff159ec8e in ?? () from /usr/lib/fglrx/libamdocl64.so

#30 0x00007ffff1559365 in ?? () from /usr/lib/fglrx/libamdocl64.so

#31 0x00007ffff156dd1c in ?? () from /usr/lib/fglrx/libamdocl64.so

#32 0x00007ffff154e12d in clBuildProgram () from /usr/lib/fglrx/libamdocl64.so

#33 0x00007ffff7b9e762 in beProgramBuilderOpenCL::BuildOpenCLProgramWrapper (this=<optimized out>, status=@0x7fffffffd0f8: 6780480, program=<optimized out>,

    num_devices=<optimized out>, device_list=<optimized out>, options=<optimized out>, pfn_notify=0x0, user_data=0x0)

    at /data/jenkins/workspace/CodeXL-Full-Linux-Release-1.4/release/CodeXL/1.4/CommonProjects/AMDTBackEnd/src/beProgramBuilderOpenCL.cpp:1674

#34 0x00007ffff7ba41b4 in beProgramBuilderOpenCL::CompileOpenCLInternal (this=0x6417a0, programSource=..., oclOptions=..., requestedDeviceId=0x677640,

    program=@0x7fffffffda18: 0xdd3990, definesAndOptions=..., iCompilationNo=0, errString=...)

    at /data/jenkins/workspace/CodeXL-Full-Linux-Release-1.4/release/CodeXL/1.4/CommonProjects/AMDTBackEnd/src/beProgramBuilderOpenCL.cpp:1377

#35 0x00007ffff7ba6604 in beProgramBuilderOpenCL::Compile (this=0x6417a0, programSource=..., oclOptions=..., pSourcePath=<optimized out>)

    at /data/jenkins/workspace/CodeXL-Full-Linux-Release-1.4/release/CodeXL/1.4/CommonProjects/AMDTBackEnd/src/beProgramBuilderOpenCL.cpp:1297

#36 0x000000000041a632 in kcCLICommanderCL::Compile (this=<optimized out>, config=...)

    at Components/KernelAnalyzer/AMDTKernelAnalyzerCLI/src/kcCLICommanderCL.cpp:111

#37 0x000000000041a80a in kcCLICommanderCL::RunCompileCommands (this=0x641710, config=..., callback=<optimized out>)

    at Components/KernelAnalyzer/AMDTKernelAnalyzerCLI/src/kcCLICommanderCL.cpp:361

#38 0x000000000040c2bc in main (argc=<optimized out>, argv=0x7fffffffded8) at Components/KernelAnalyzer/AMDTKernelAnalyzerCLI/src/kcMain.cpp:78

0 Likes
1 Solution

Hi,


struct btest_struct {


    size_t sz;


    test_p dt[];


};



As per OpenCL 1.2 spec, inside the kernel code variable length arrays and structures with flexible (or unsized) arrays are not supported. [see chapter 6.9 Restrictions]. So, you need to supply size of the array at compile time.

Now, regarding the compiler crashing, I found the same problem for above kernel code with driver 14.30 (but on Windows 7) using both CodeXL and VSC++. I may need to submit a internal bug report for this issue (thanks for reporting). But it worked fine when size of the array "test_p dt[]" was set.

To test the points 2 and 3 under your additional notes, I added the following lines to the end of the kernel code and tried to compile. It worked fine for both CodeXL and VSC++.


void ** vpp = btt.dt;


test_p * tpp = btt.dt;



So, my suggestion is please check with the latest driver and let us know your findings.

Regards,

View solution in original post

0 Likes
3 Replies
chris9871
Journeyman III

Additional notes:

  1. In the declaration: test_p dt[0] is the same, however test_p dt[1] compiles successfully.
  2. Even without the pointer math the compiler crashes: void ** vpp = btt.dt; (I would expect a compiler warning or error)
  3. Even without casting, the compiler crashes: test_p * tpp = btt.dt;

Chris...

0 Likes

Hi,


struct btest_struct {


    size_t sz;


    test_p dt[];


};



As per OpenCL 1.2 spec, inside the kernel code variable length arrays and structures with flexible (or unsized) arrays are not supported. [see chapter 6.9 Restrictions]. So, you need to supply size of the array at compile time.

Now, regarding the compiler crashing, I found the same problem for above kernel code with driver 14.30 (but on Windows 7) using both CodeXL and VSC++. I may need to submit a internal bug report for this issue (thanks for reporting). But it worked fine when size of the array "test_p dt[]" was set.

To test the points 2 and 3 under your additional notes, I added the following lines to the end of the kernel code and tried to compile. It worked fine for both CodeXL and VSC++.


void ** vpp = btt.dt;


test_p * tpp = btt.dt;



So, my suggestion is please check with the latest driver and let us know your findings.

Regards,

0 Likes

Hi Dipak,

Thanks for your response.  I will note one slight error with your response.  While the spec does not allow for an array with no specified size, it doesn't disallow a size of 0 which also fails.  I suppose that could be a choice - the AMD compiler could choose to fail on a size of 0 (since that usually means 'variable').  I do find that this part of the spec makes for slightly more messy code when the actual size of the structure is variable.


Cheers,

Chris...

0 Likes