Can't compile shared static array in dcompute

[gulugulubing]

Does current ldc's dcompute support shared memory?
I tried to compile this:

@compute(CompileFor.deviceOnly) module shared_static_array_kernel;

import ldc.dcompute;

private enum N = 16u;

@kernel()
void k(GlobalPointer!float g) {
    Shared!(float[N]) tile;
    tile[0] = g[0];
}

but it crashed:
/Users/qiugaofei/dlang/ldc/runtime/druntime/src/core/internal/array/equality.d(42): Error: can only call functions from other @compute modules in @compute code
return __equals(lhs[], rhs[]);
^
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var LLVM_SYMBOLIZER_PATH to point to it):
0 libLLVM.dylib 0x0000000112d0db48 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 56
1 libLLVM.dylib 0x0000000112d0e328 SignalHandler(int, __siginfo*, void*) + 272
2 libsystem_platform.dylib 0x000000019cb64624 _sigtramp + 56
3 ldc2 0x00000001043566a4 RecursiveWalker::visit(CallExp*) + 40
4 ldc2 0x00000001043566a4 RecursiveWalker::visit(CallExp*) + 40
5 ldc2 0x0000000104356088 RecursiveWalker::visit(CompoundStatement*) + 132
6 ldc2 0x0000000104356088 RecursiveWalker::visit(CompoundStatement*) + 132
7 ldc2 0x0000000104355338 RecursiveWalker::visit(AttribDeclaration*) + 132
8 ldc2 0x000000010439206c dcomputeSemanticAnalysis(Module*) + 144
9 ldc2 0x000000010406f314 extraLDCSpecificSemanticAnalysis(Array<Module*>&) + 80
10 ldc2 0x00000001041bcd28 mars_tryMain(Param&, Array<char const*>&) + 4308
11 ldc2 0x00000001043fafe0 cppmain() + 6764
12 ldc2 0x000000010455e4f0 _D2rt6dmain212_d_run_main2UAAamPUQgZiZ6runAllMFZv + 120
13 ldc2 0x000000010455e29c _d_run_main2 + 364
14 ldc2 0x000000010455e114 _d_run_main + 144
15 ldc2 0x00000001043f951c main + 596
16 dyld 0x000000019c78ab98 start + 6076