[Libclc-dev] LIBCLC and AMD RADEON HD R9 290 (Hawaii)

[Ricardo Nobre via Libclc-dev]

Hi,

I'm using Clang/LLVM 3.7 and libclc from the repository to compile and
OpenCL kernel.

typedef float DATA_TYPE;

__kernel void gemm(__global DATA_TYPE *a, __global DATA_TYPE *b, __global
DATA_TYPE *c, DATA_TYPE alpha, DATA_TYPE beta, int ni, int nj, int nk)
{
        int j = get_global_id(0);
        int i = get_global_id(1);

        if ((i < ni) && (j < nj))
        {
                c[i * nj + j] *= beta;
                int k;
                for(k=0; k < nk; k++)
                {
                        c[i * nj + j] += alpha * a[i * nk + k] * b[k * nj
+j];
                }
        }
}

I'm doing it as follows:
clang -Dcl_clang_storage_class_specifiers -isystem libclc/generic/include
-include clc/clc.h -target amdgcn -S -emit-llvm -xcl -o gemm.ll gemm.cl
llvm-link /usr/local/lib/clc/hawaii-amdgcn--.bc gemm.ll -o gemm.linked.bc
clang -target amdgcn gemm.linked.bc -S -o gemm.hawaii.s

When I execute the third command, I get the following error:

warning: overriding the module target triple with amdgcn [-Woverride-module]
Pass 'Structurize control flow' is not initialized.
Verify if there is a pass dependency cycle.
Required Passes:
        Lower SwitchInst's to branches
        Dominator Tree Construction
#0 0x13e06b8 llvm::sys::PrintStackTrace(llvm::raw_ostream&)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x13e06b8)
#1 0x13e1a1b SignalHandler(int)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x13e1a1b)
#2 0x7fcb0f850340 __restore_rt
(/lib/x86_64-linux-gnu/libpthread.so.0+0x10340)
#3 0x1122128 llvm::PMTopLevelManager::schedulePass(llvm::Pass*)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x1122128)
#4 0xf7015e llvm::TargetPassConfig::addPass(llvm::Pass*, bool, bool)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0xf7015e)
#5 0x7ffd69 (anonymous namespace)::GCNPassConfig::addPreISel()
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x7ffd69)
#6 0xf709f5 llvm::TargetPassConfig::addISelPrepare()
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0xf709f5)
#7 0xef6807 addPassesToGenerateCode(llvm::LLVMTargetMachine*,
llvm::legacy::PassManagerBase&, bool, void const*, void const*, void
const*, llvm::MachineFunctionInitializer*)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0xef6807)
#8 0xef60cd
llvm::LLVMTargetMachine::addPassesToEmitFile(llvm::legacy::PassManagerBase&,
llvm::raw_pwrite_stream&, llvm::TargetMachine::CodeGenFileType, bool, void
const*, void const*, void const*, llvm::MachineFunctionInitializer*)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0xef60cd)
#9 0x14b5c79 clang::EmitBackendOutput(clang::DiagnosticsEngine&,
clang::CodeGenOptions const&, clang::TargetOptions const&,
clang::LangOptions const&, llvm::StringRef, llvm::Module*,
clang::BackendAction, llvm::raw_pwrite_stream*)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x14b5c79)
#10 0x197634d clang::CodeGenAction::ExecuteAction()
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x197634d)
#11 0x173c9e9 clang::FrontendAction::Execute()
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x173c9e9)
#12 0x170d4d3
clang::CompilerInstance::ExecuteAction(clang::FrontendAction&)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x170d4d3)
#13 0x17a9de3 clang::ExecuteCompilerInvocation(clang::CompilerInstance*)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x17a9de3)
#14 0x6f6634 cc1_main(llvm::ArrayRef<char const*>, char const*, void*)
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x6f6634)
#15 0x6f586f main
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x6f586f)
#16 0x7fcb0ec7cec5 __libc_start_main
(/lib/x86_64-linux-gnu/libc.so.6+0x21ec5)
#17 0x6f2893 _start
(/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7+0x6f2893)
Stack dump:
0.      Program arguments:
/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/clang-3.7 -cc1
-triple amdgcn -S -disable-free -disable-llvm-verifier -main-file-name
gemm.linked.bc -mrelocation-model static -mthread-model posix
-mdisable-fp-elim -fmath-errno -no-integrated-as -mconstructor-aliases
-dwarf-column-info -coverage-file
/home/rjfn/Desktop/polybench-gpu-1.0/OpenCL/GEMM/gemm.hawaii.s
-resource-dir
/opt/clang+llvm-3.7.0-x86_64-linux-gnu-ubuntu-14.04/bin/../lib/clang/3.7.0
-fno-dwarf-directory-asm -fdebug-compilation-dir
/home/rjfn/Desktop/polybench-gpu-1.0/OpenCL/GEMM -ferror-limit 19
-fmessage-length 0 -mstackrealign -fobjc-runtime=gcc
-fdiagnostics-show-option -o gemm.hawaii.s -x ir gemm.linked.bc
clang-3.7: error: unable to execute command: Segmentation fault (core
dumped)
clang-3.7: error: clang frontend command failed due to signal (use -v to
see invocation)
clang version 3.7.0 (tags/RELEASE_370/final)
Target: amdgcn
Thread model: posix
clang-3.7: note: diagnostic msg: PLEASE submit a bug report to
http://llvm.org/bugs/ and include the crash backtrace, preprocessed source,
and associated run script.
clang-3.7: note: diagnostic msg: Error generating preprocessed source(s) -
no preprocessable inputs.

What I'm I doing wrong?
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/libclc-dev/attachments/20151026/801556f8/attachment.html>