51 lines · plain
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 %s16; 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 opaque20 21define spir_kernel void @verify_linked_lists(%struct.Node addrspace(1)* %pNodes) {22entry:23 br label %for.cond24 25for.cond: ; preds = %for.inc, %entry26 %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: OpPhi30 31 %cmp = icmp slt i32 %j.0, 1032 br i1 %cmp, label %for.body, label %for.end33 34for.body: ; preds = %for.cond35 %pNext = getelementptr inbounds %struct.Node, %struct.Node addrspace(1)* %pNode.0, i32 0, i32 036 37 %0 = load %struct.Node.0 addrspace(1)*, %struct.Node.0 addrspace(1)* addrspace(1)* %pNext, align 438 %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.inc43 44for.inc: ; preds = %for.body45 %inc = add nsw i32 %j.0, 146 br label %for.cond47 48for.end: ; preds = %for.cond49 ret void50}51