Shared memory + multiple function exits cause invalid results

[maleadt]

Cause seems to be an added checkbounds, if that even makes sense.

Repro:

using CUDAdrv, CUDAnative

@target ptx function kernel(arr::Ptr{Int32})
    temp = @cuStaticSharedMem(Int32, (2, 1))
    tx = Int(threadIdx().x)

    if tx == 1
        for i = 1:2
            # THIS BREAKS STUFF: checkbounds(temp, i)
            Base.pointerset(temp.ptr, 1, i, 8)
        end
    end
    sync_threads()

    Base.pointerset(arr, Base.pointerref(temp.ptr, tx, 8), tx, 8)

    return nothing
end

dev = CuDevice(0)
ctx = CuContext(dev)

d_arr = CuArray(Int32, (2, 1))
@cuda (1,2) kernel(d_arr.ptr)
println(Array(d_arr))

destroy(ctx)

Result without checkbounds: [1; 1]. With: [1; 0].

cc @cfoket

[maleadt]

Not reproducible with this code anymore, but rodinia/lud.jl still fails with --check-bounds=yes probably still caused by the same underlying issue.

[maleadt]

New repro, again using shared memory + bounds checking, but this time the invalid value is the result of __shfl_down (not touching shared memory at all):

using CUDAdrv, CUDAnative

function kernel(ptr::Ptr{Cint})
    shared = @cuStaticSharedMem(Cint, 4)

    lane = (threadIdx().x-1) % warpsize

    if lane == 0
        @boundscheck Base.checkbounds(shared, threadIdx().x)
        unsafe_store!(shared.ptr, 0, threadIdx().x)
    end

    sync_threads()

    val = shfl_down(Cint(32), 1, 4)
    if lane == 0
        unsafe_store!(ptr, val)
    end

    return
end

dev = CuDevice(0)
ctx = CuContext(dev)

gpu_val = CuArray(Cint, 1)
@cuda dev (1,4) kernel(gpu_val.ptr)
val = Array(gpu_val)[1]
println(val)

destroy(ctx)

Returns 0 with checkbounds, 32 without.

[maleadt]

Managed to reduce to two sets of LLVM IR, executed using the following snippet:

using CUDAdrv, CUDAnative, LLVM

dev = CuDevice(0)
ctx = CuContext(dev)

for ir_fn in ["bug-working.ll", "bug-broken.ll"]
    gpu_val = CuArray(Cint[42])

    ir = readstring(ir_fn)
    mod = parse(LLVM.Module, ir)
    fn = "kernel"
    entry = get(functions(mod), "kernel")
    ptx = CUDAnative.mcgen(mod, entry, v"3.0")

    cuda_mod = CuModule(ptx)
    cuda_fun = CuFunction(cuda_mod, fn)

    cudacall(cuda_fun, 1, 4, (Ptr{Cint},), gpu_val.ptr)

    val = Array(gpu_val)[1]
    println(val)
end

destroy(ctx)

Working IR:

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

@shmem = internal addrspace(3) global [4 x i32] zeroinitializer, align 4

define void @kernel(i32*) {
top:
  %1 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %2 = and i32 %1, 31
  %3 = icmp eq i32 %2, 0
  br i1 %3, label %lane0_boundscheck, label %sync_shfl

lane0_boundscheck:
  %4 = icmp ugt i32 %1, 3
  br i1 %4, label %lane0_oob, label %lane0_shmem

lane0_oob:
  tail call void @llvm.trap()
  unreachable

sync_shfl:
  tail call void @llvm.nvvm.barrier0()
  %5 = tail call i32 @llvm.nvvm.shfl.down.i32(i32 32, i32 1, i32 7199)
  br i1 %3, label %lane0_writeback, label %end

lane0_shmem:
  %6 = getelementptr [4 x i32], [4 x i32] addrspace(3)* @shmem, i32 0, i32 %1
  store i32 0, i32 addrspace(3)* %6, align 8
  br label %sync_shfl

lane0_writeback:
  store i32 %5, i32* %0, align 8
  br label %end

end:
  ret void
}

declare void @llvm.trap()
declare i32 @llvm.nvvm.shfl.down.i32(i32, i32, i32)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
declare void @llvm.nvvm.barrier0()

Broken IR:

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

@shmem = internal addrspace(3) global [4 x i32] zeroinitializer, align 4

define void @kernel(i32*) {
top:
  %1 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
  %2 = and i32 %1, 31
  %3 = icmp eq i32 %2, 0
  br i1 %3, label %lane0_boundscheck, label %sync_shfl

lane0_boundscheck:
  %4 = icmp ugt i32 %1, 3
  br i1 %4, label %lane0_oob, label %lane0_shmem

sync_shfl:
  tail call void @llvm.nvvm.barrier0()
  %5 = tail call i32 @llvm.nvvm.shfl.down.i32(i32 32, i32 1, i32 7199)
  br i1 %3, label %lane0_writeback, label %end

lane0_oob:
  tail call void @llvm.trap()
  unreachable

lane0_shmem:
  %6 = getelementptr [4 x i32], [4 x i32] addrspace(3)* @shmem, i32 0, i32 %1
  store i32 0, i32 addrspace(3)* %6, align 8
  br label %sync_shfl

lane0_writeback:
  store i32 %5, i32* %0, align 8
  br label %end

end:
  ret void
}

declare void @llvm.trap()
declare i32 @llvm.nvvm.shfl.down.i32(i32, i32, i32)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
declare void @llvm.nvvm.barrier0()

That's right, the only difference between those two is the placement of the oob BB...
cc @cfoket

[maleadt]

One layer deeper...

Working PTX:

.version 3.2
.target sm_30
.address_size 64

.visible .entry kernel(
        .param .u64 output  // single int output
)
{
        .reg .pred      %p<4>;
        .reg .b32       %r<6>;
        .reg .b64       %rd<6>;
        .shared .align 4 .b8 shmem[16];     // 4 integers
        ld.param.u64    %rd1, [output];

        // calculate lane, check if 0
        mov.u32 %r1, %tid.x;
        and.b32         %r2, %r1, 31;
        setp.ne.s32     %p1, %r2, 0;
        @%p1 bra        BB_SHFL;

        // bounds check for shmem access
        setp.lt.u32     %p2, %r1, 4;
        @%p2 bra        BB_SHMEM;
        bra.uni         BB_OOB;
BB_SHMEM:
        mul.wide.s32    %rd2, %r1, 4;
        mov.u64         %rd3, shmem;
        add.s64         %rd4, %rd3, %rd2;
        mov.u32         %r4, 0;
        st.shared.u32   [%rd4], %r4;
BB_SHFL:
        setp.eq.s32     %p3, %r2, 0;
        bar.sync        0;
        mov.u32         %r5, 32;
        shfl.down.b32 %r3, %r5, 1, 7199;
        @%p3 bra        BB_WRITEBACK;
        bra.uni         BB_END;
BB_WRITEBACK:
        cvta.to.global.u64      %rd5, %rd1;
        st.global.u32   [%rd5], %r3;
BB_END:
        ret;
BB_OOB:
        trap;
}

Broken PTX:

.version 3.2
.target sm_30
.address_size 64

.visible .entry kernel(
        .param .u64 output  // single int output
)
{
        .reg .pred      %p<4>;
        .reg .b32       %r<6>;
        .reg .b64       %rd<6>;
        .shared .align 4 .b8 shmem[16];     // 4 integers
        ld.param.u64    %rd1, [output];

        // calculate lane, check if 0
        mov.u32 %r1, %tid.x;
        and.b32         %r2, %r1, 31;
        setp.ne.s32     %p1, %r2, 0;
        @%p1 bra        BB_SHFL;

        // bounds check for shmem access
        setp.gt.u32     %p2, %r1, 3;
        @%p2 bra        BB_OOB;
        bra.uni         BB_SHMEM;
BB_SHMEM:
        mul.wide.s32    %rd2, %r1, 4;
        mov.u64         %rd3, shmem;
        add.s64         %rd4, %rd3, %rd2;
        mov.u32         %r4, 0;
        st.shared.u32   [%rd4], %r4;
BB_SHFL:
        setp.eq.s32     %p3, %r2, 0;
        bar.sync        0;
        mov.u32         %r5, 32;
        shfl.down.b32 %r3, %r5, 1, 7199;
        @%p3 bra        BB_WRITEBACK;
        bra.uni         BB_END;
BB_WRITEBACK:
        cvta.to.global.u64      %rd5, %rd1;
        st.global.u32   [%rd5], %r3;
BB_END:
        ret;
BB_OOB:
        trap;
}

Loader:

using CUDAdrv

dev = CuDevice(0)
ctx = CuContext(dev)

fn = "kernel"

for name in ["bug-working", "bug-broken"]
    gpu_val = CuArray(Cint[42])

    ptx = readstring("$name.ptx")

    cuda_mod = CuModule(ptx)
    cuda_fun = CuFunction(cuda_mod, fn)

    cudacall(cuda_fun, 1, 4, (Ptr{Cint},), gpu_val.ptr)

    val = Array(gpu_val)[1]
    println(val)
end

destroy(ctx)

Only difference: the bounds-check branch (>3 or <4):

$ diff bug-working.ptx bug-broken.ptx                                                                                                                                            *[master] 
22,24c22,24
<         setp.lt.u32     %p2, %r1, 4;
<         @%p2 bra        BB_SHMEM;
<         bra.uni         BB_OOB;
---
>         setp.gt.u32     %p2, %r1, 3;
>         @%p2 bra        BB_OOB;
>         bra.uni         BB_SHMEM;

Probably an assembler bug.

[maleadt]

Alternative loader, using ptxas to generate a cubin (in order to play with ptxas optimization flags, but doesn't seem to matter):

using CUDAdrv

dev = CuDevice(0)
ctx = CuContext(dev)

fn = "kernel"

for name in ["kernel-working", "kernel-broken"]
    gpu_val = CuArray(Cint[42])

    run(`ptxas -arch=sm_61 -o $name.cubin $name.ptx`)

    cuda_mod = CuModule(read("$name.cubin"))
    cuda_fun = CuFunction(cuda_mod, fn)

    cudacall(cuda_fun, 1, 4, (Ptr{Cint},), gpu_val.ptr)

    val = Array(gpu_val)[1]
    println(val)
end

destroy(ctx)

[maleadt]

Almost definitely looks like an assembler bug. See the following annotated & prettified Pascal SASS (sm_61):

Working version:

kernel:
.text.kernel:
        MOV R1, c[0x0][0x20];
        S2R R2, SR_TID.X;
        SSY `(BB_SHFL);         // push BB_SHFL on reconvergence stack

        // calculate lane, check if 0
        LOP32I.AND R0, R2, 0x1f;
        ISETP.NE.AND P0, PT, R0, RZ, PT;
    @P0 SYNC                    // not lane 0, pop BB_SHFL from reconvergence stack

        // bounds check for shmem access
        ISETP.LT.U32.AND P0, PT, R2, 0x4, PT;
    @P0 BRA `(BB_SHMEM);

//BB_OOB:
        BPT.TRAP 0x1;
        EXIT;
BB_SHMEM:
        SHL R2, R2, 0x2;
        STS [R2], RZ;
        SYNC                    // pop BB_SHFL from reconvergence stack
BB_SHFL:
        // check if lane 0
  {     ISETP.EQ.AND P0, PT, R0, RZ, PT;
        BAR.SYNC 0x0;        }
        // shuffle unconditionally
        MOV32I R0, 0x20;
        SHFL.DOWN PT, R0, R0, 0x1, 0x1c1f;
   @!P0 EXIT;                  // not lane 0, exit
//BB_WRITEBACK:
        MOV R2, c[0x0][0x140];
        MOV R3, c[0x0][0x144];
        STG.E [R2], R0;
        EXIT;
.BB_END:
        BRA `(.BB_END);

Broken version:

kernel:
.text.kernel:
        MOV R1, c[0x0][0x20];
        S2R R2, SR_TID.X;

        // calculate lane, check if 0
        LOP32I.AND R0, R2, 0x1f;
        ISETP.NE.AND P0, PT, R0, RZ, PT;
    @P0 BRA `(BB_SHFL);         // not lane 0, branch to BB_SHFL

        // bounds check for shmem access
        ISETP.GT.U32.AND P0, PT, R2, 0x3, PT;
    @P0 BRA `(BB_OOB);

//BB_SHMEM:
        SHL R2, R2, 0x2;
        STS [R2], RZ;
BB_SHFL:
        // check if lane 0
 {      ISETP.EQ.AND P0, PT, R0, RZ, PT;
        BAR.SYNC 0x0;        }
        // shuffle unconditionally
        MOV32I R0, 0x20;
        SHFL.DOWN PT, R0, R0, 0x1, 0x1c1f;
   @!P0 EXIT;                  // not lane 0, exit
//BB_WRITEBACK:
        MOV R2, c[0x0][0x140];
        MOV R3, c[0x0][0x144];
        STG.E [R2], R0;
        EXIT;
BB_OOB:
        BPT.TRAP 0x1;
        EXIT;
.L_3:
        BRA `(.L_3);
.L_18:

The broken version clearly messes up its reconvergence stack, not pushing anything on it despite multiple conditional branches (for some info on how this works, see this paper by Bialas and Strzelecki)...

[maleadt]

And a C++ loader, for reporting purposes.

#include <stdio.h>

#include <cuda.h>

#define CHECK(err) __check(err, __FILE__, __LINE__)
inline void __check(CUresult err, const char *file, const int line) {
  if (CUDA_SUCCESS != err) {
    const char *name, *descr;
    cuGetErrorName(err, &name);
    cuGetErrorString(err, &name);
    fprintf(stderr, "CUDA error #%s: %s at %s:%i\n", name, descr, file, line);
    abort();
  }
}

int test(const char *path) {
  CUmodule mod;
  cuModuleLoad(&mod, path);

  CUfunction fun;
  CHECK(cuModuleGetFunction(&fun, mod, "kernel"));

  int *gpu_val;
  CHECK(cuMemAlloc((CUdeviceptr*) &gpu_val, sizeof(int)));

  void *args[1] = {&gpu_val};
  cuLaunchKernel(fun, 1, 1, 1, 4, 1, 1, 0, NULL, args, NULL);

  int val;
  CHECK(cuMemcpyDtoH(&val, (CUdeviceptr) gpu_val, sizeof(int)));

  CHECK(cuModuleUnload(mod));

  return val;
}

int main() {
  CHECK(cuInit(0));

  CUdevice dev;
  CHECK(cuDeviceGet(&dev, 0));

  CUcontext ctx;
  CHECK(cuCtxCreate(&ctx, 0, dev));

  printf("working: %d\n", test("kernel-working.ptx"));
  printf("broken: %d\n", test("kernel-broken.ptx"));

  CHECK(cuCtxDestroy(ctx));

  return 0;
}

Will probably submit this to NVIDIA soon, unless anybody still spots us doing something wrong.

[maleadt]

Reported this repro to NVIDIA, bug #1833004. Will disable bounds checking for the time being.

[vchuravy]

Could we fix this on the LLVM side? Any bugfix to the assembler is going to be deployed slowly.

[maleadt]

I haven't figured out what PTX pattern exactly triggers the SASS emission bug. Probably the branch to a trap BB. I've asked NVIDIA for some background on the bug, if they deem it a bug, so I'm going to wait for them to respond before sinking more time into this.