| ;; struct Node; |
| ;; typedef struct { |
| ;; __global struct Node* pNext; |
| ;; } Node; |
| ;; |
| ;; __kernel void verify_linked_lists(__global Node* pNodes) |
| ;; { |
| ;; __global Node *pNode = pNodes; |
| ;; |
| ;; for(int j=0; j < 10; j++) { |
| ;; pNode = pNode->pNext; |
| ;; } |
| ;; } |
| |
| ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV |
| ; XFAIL: * |
| |
| %struct.Node = type { %struct.Node.0 addrspace(1)* } |
| %struct.Node.0 = type opaque |
| |
| define spir_kernel void @verify_linked_lists(%struct.Node addrspace(1)* %pNodes) { |
| entry: |
| br label %for.cond |
| |
| for.cond: ; preds = %for.inc, %entry |
| %pNode.0 = phi %struct.Node addrspace(1)* [ %pNodes, %entry ], [ %1, %for.inc ] |
| %j.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ] |
| ; CHECK-SPIRV: %[[#]] = OpPhi %[[#]] %[[#]] %[[#]] %[[#BitcastResultId:]] %[[#]] |
| ; CHECK-SPIRV-NEXT: OpPhi |
| |
| %cmp = icmp slt i32 %j.0, 10 |
| br i1 %cmp, label %for.body, label %for.end |
| |
| for.body: ; preds = %for.cond |
| %pNext = getelementptr inbounds %struct.Node, %struct.Node addrspace(1)* %pNode.0, i32 0, i32 0 |
| |
| %0 = load %struct.Node.0 addrspace(1)*, %struct.Node.0 addrspace(1)* addrspace(1)* %pNext, align 4 |
| %1 = bitcast %struct.Node.0 addrspace(1)* %0 to %struct.Node addrspace(1)* |
| ; CHECK-SPIRV: %[[#LoadResultId:]] = OpLoad %[[#]] |
| ; CHECK-SPIRV: %[[#BitcastResultId]] = OpBitcast %[[#]] %[[#LoadResultId]] |
| |
| br label %for.inc |
| |
| for.inc: ; preds = %for.body |
| %inc = add nsw i32 %j.0, 1 |
| br label %for.cond |
| |
| for.end: ; preds = %for.cond |
| ret void |
| } |