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: llvm-as %s -o %t.bc 16 ; RUN: llvm-spirv %t.bc -spirv-text -o %t.txt 17 ; RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV 18 ; RUN: llvm-spirv %t.bc -o %t.spv 19 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc 20 ; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM 21 22 target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" 23 target triple = "spir" 24 25 %struct.Node = type { %struct.Node.0 addrspace(1)* } 26 %struct.Node.0 = type opaque 27 28 ; Function Attrs: nounwind 29 define spir_kernel void @verify_linked_lists(%struct.Node addrspace(1)* %pNodes) #0 { 30 entry: 31 br label %for.cond 32 33 for.cond: ; preds = %for.inc, %entry 34 %pNode.0 = phi %struct.Node addrspace(1)* [ %pNodes, %entry ], [ %1, %for.inc ] 35 %j.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ] 36 ;CHECK-SPIRV: Phi {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[BitcastResultId:[0-9]+]] {{[0-9]+}} 37 ;CHECK-SPIRV-NEXT: Phi 38 ;CHECK-LLVM: phi %struct.Node addrspace(1)* [ %pNodes, %entry ], [ [[BitcastResult:%[0-9]+]], %for.inc ] 39 ;CHECK-LLVM-NEXT: phi 40 41 %cmp = icmp slt i32 %j.0, 10 42 br i1 %cmp, label %for.body, label %for.end 43 44 for.body: ; preds = %for.cond 45 %pNext = getelementptr inbounds %struct.Node, %struct.Node addrspace(1)* %pNode.0, i32 0, i32 0 46 47 %0 = load %struct.Node.0 addrspace(1)*, %struct.Node.0 addrspace(1)* addrspace(1)* %pNext, align 4 48 %1 = bitcast %struct.Node.0 addrspace(1)* %0 to %struct.Node addrspace(1)* 49 ;CHECK-SPIRV: Load {{[0-9]+}} [[LoadResultId:[0-9]+]] 50 ;CHECK-SPIRV: Bitcast {{[0-9]+}} [[BitcastResultId]] [[LoadResultId]] 51 ;CHECK-LLVM: [[LoadResult:%[0-9]+]] = load %struct.Node.0, %struct.Node.0 addrspace(1)* addrspace(1)* {{.*}} 52 ;CHECK-LLVM: [[BitcastResult]] = bitcast %struct.Node.0 addrspace(1)* [[LoadResult]] to %struct.Node addrspace(1)* 53 54 br label %for.inc 55 56 for.inc: ; preds = %for.body 57 %inc = add nsw i32 %j.0, 1 58 br label %for.cond 59 60 for.end: ; preds = %for.cond 61 ret void 62 } 63 64 attributes #0 = { nounwind "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } 65 66 !opencl.kernels = !{!0} 67 !opencl.enable.FP_CONTRACT = !{} 68 !opencl.spir.version = !{!6} 69 !opencl.ocl.version = !{!7} 70 !opencl.used.extensions = !{!8} 71 !opencl.used.optional.core.features = !{!8} 72 !opencl.compiler.options = !{!8} 73 74 !0 = !{void (%struct.Node addrspace(1)*)* @verify_linked_lists, !1, !2, !3, !4, !5} 75 !1 = !{!"kernel_arg_addr_space", i32 1} 76 !2 = !{!"kernel_arg_access_qual", !"none"} 77 !3 = !{!"kernel_arg_type", !"Node*"} 78 !4 = !{!"kernel_arg_base_type", !"struct __Node*"} 79 !5 = !{!"kernel_arg_type_qual", !""} 80 !6 = !{i32 1, i32 2} 81 !7 = !{i32 2, i32 0} 82 !8 = !{} 83 84
