1;; struct Node; 2;; typedef struct { 3;; __global struct Node* pNext; 4;; } Node; 5;; 6;; __kernel void verify_linked_lists(__global Node* pNodes) 7;; { 8;; __global Node *pNode = pNodes; 9;; 10;; for(int j=0; j < 10; j++) { 11;; pNode = pNode->pNext; 12;; } 13;; } 14 15; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s 16; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} 17 18%struct.Node = type { %struct.Node.0 addrspace(1)* } 19%struct.Node.0 = type opaque 20 21define spir_kernel void @verify_linked_lists(%struct.Node addrspace(1)* %pNodes) { 22entry: 23 br label %for.cond 24 25for.cond: ; preds = %for.inc, %entry 26 %pNode.0 = phi %struct.Node addrspace(1)* [ %pNodes, %entry ], [ %1, %for.inc ] 27 %j.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ] 28; CHECK: %[[#]] = OpPhi %[[#]] %[[#]] %[[#]] %[[#BitcastResultId:]] %[[#]] 29; CHECK-NEXT: OpPhi 30 31 %cmp = icmp slt i32 %j.0, 10 32 br i1 %cmp, label %for.body, label %for.end 33 34for.body: ; preds = %for.cond 35 %pNext = getelementptr inbounds %struct.Node, %struct.Node addrspace(1)* %pNode.0, i32 0, i32 0 36 37 %0 = load %struct.Node.0 addrspace(1)*, %struct.Node.0 addrspace(1)* addrspace(1)* %pNext, align 4 38 %1 = bitcast %struct.Node.0 addrspace(1)* %0 to %struct.Node addrspace(1)* 39; CHECK: %[[#LoadResultId:]] = OpLoad %[[#]] 40; CHECK: %[[#BitcastResultId]] = OpBitcast %[[#]] %[[#LoadResultId]] 41 42 br label %for.inc 43 44for.inc: ; preds = %for.body 45 %inc = add nsw i32 %j.0, 1 46 br label %for.cond 47 48for.end: ; preds = %for.cond 49 ret void 50} 51