llvm/test/CodeGen/SPIRV/transcoding/OpPhi_ArgumentsPlaceholders.ll

;; 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 -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}

%struct.Node = type { ptr addrspace(1) }
%struct.Node.0 = type opaque

define spir_kernel void @verify_linked_lists(ptr addrspace(1) %pNodes) {
entry:
  br label %for.cond

for.cond:                                         ; preds = %for.inc, %entry
  %pNode.0 = phi ptr addrspace(1) [ %pNodes, %entry ], [ %1, %for.inc ]
  %j.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ]
; CHECK:      %[[#]] = OpPhi %[[#]] %[[#]] %[[#]] %[[#BitcastResultId:]] %[[#]]
; CHECK-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, ptr addrspace(1) %pNode.0, i32 0, i32 0

  %0 = load ptr addrspace(1), ptr addrspace(1) %pNext, align 4
  %1 = bitcast ptr addrspace(1) %0 to ptr addrspace(1)
; CHECK: %[[#LoadResultId:]] = OpLoad %[[#]]
; CHECK: %[[#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
}