107 lines · c
1// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device -x c++ %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK642// RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK643// RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK324 5// CHECK: @{{.*}}test_num_workgroups{{.*}}(6// CHECK-NEXT: [[ENTRY:.*:]]7// CHECK64-NEXT: tail call i64 @llvm.spv.num.workgroups.i64(i32 0)8// CHECK32-NEXT: tail call i32 @llvm.spv.num.workgroups.i32(i32 0)9//10[[clang::sycl_external]] unsigned int test_num_workgroups() {11 return __builtin_spirv_num_workgroups(0);12}13 14// CHECK: @{{.*}}test_workgroup_size{{.*}}(15// CHECK-NEXT: [[ENTRY:.*:]]16// CHECK64-NEXT: tail call i64 @llvm.spv.workgroup.size.i64(i32 0)17// CHECK32-NEXT: tail call i32 @llvm.spv.workgroup.size.i32(i32 0)18//19[[clang::sycl_external]] unsigned int test_workgroup_size() {20 return __builtin_spirv_workgroup_size(0);21}22 23// CHECK: @{{.*}}test_workgroup_id{{.*}}(24// CHECK-NEXT: [[ENTRY:.*:]]25// CHECK64-NEXT: tail call i64 @llvm.spv.group.id.i64(i32 0)26// CHECK32-NEXT: tail call i32 @llvm.spv.group.id.i32(i32 0)27//28[[clang::sycl_external]] unsigned int test_workgroup_id() {29 return __builtin_spirv_workgroup_id(0);30}31 32// CHECK: @{{.*}}test_local_invocation_id{{.*}}(33// CHECK-NEXT: [[ENTRY:.*:]]34// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0)35// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0)36//37[[clang::sycl_external]] unsigned int test_local_invocation_id() {38 return __builtin_spirv_local_invocation_id(0);39}40 41// CHECK: @{{.*}}test_global_invocation_id{{.*}}(42// CHECK-NEXT: [[ENTRY:.*:]]43// CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.i64(i32 0)44// CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.i32(i32 0)45//46[[clang::sycl_external]] unsigned int test_global_invocation_id() {47 return __builtin_spirv_global_invocation_id(0);48}49 50// CHECK: @{{.*}}test_global_size{{.*}}(51// CHECK-NEXT: [[ENTRY:.*:]]52// CHECK64-NEXT: tail call i64 @llvm.spv.global.size.i64(i32 0)53// CHECK32-NEXT: tail call i32 @llvm.spv.global.size.i32(i32 0)54//55[[clang::sycl_external]] unsigned int test_global_size() {56 return __builtin_spirv_global_size(0);57}58 59// CHECK: @{{.*}}test_global_offset{{.*}}(60// CHECK-NEXT: [[ENTRY:.*:]]61// CHECK64-NEXT: tail call i64 @llvm.spv.global.offset.i64(i32 0)62// CHECK32-NEXT: tail call i32 @llvm.spv.global.offset.i32(i32 0)63//64[[clang::sycl_external]] unsigned int test_global_offset() {65 return __builtin_spirv_global_offset(0);66}67 68// CHECK: @{{.*}}test_subgroup_size{{.*}}(69// CHECK-NEXT: [[ENTRY:.*:]]70// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.size()71//72[[clang::sycl_external]] unsigned int test_subgroup_size() {73 return __builtin_spirv_subgroup_size();74}75 76// CHECK: @{{.*}}test_subgroup_max_size{{.*}}(77// CHECK-NEXT: [[ENTRY:.*:]]78// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.max.size()79//80[[clang::sycl_external]] unsigned int test_subgroup_max_size() {81 return __builtin_spirv_subgroup_max_size();82}83 84// CHECK: @{{.*}}test_num_subgroups{{.*}}(85// CHECK-NEXT: [[ENTRY:.*:]]86// CHECK-NEXT: tail call i32 @llvm.spv.num.subgroups()87//88[[clang::sycl_external]] unsigned int test_num_subgroups() {89 return __builtin_spirv_num_subgroups();90}91 92// CHECK: @{{.*}}test_subgroup_id{{.*}}(93// CHECK-NEXT: [[ENTRY:.*:]]94// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.id()95//96[[clang::sycl_external]] unsigned int test_subgroup_id() {97 return __builtin_spirv_subgroup_id();98}99 100// CHECK: @{{.*}}test_subgroup_local_invocation_id{{.*}}(101// CHECK-NEXT: [[ENTRY:.*:]]102// CHECK-NEXT: tail call i32 @llvm.spv.subgroup.local.invocation.id()103//104[[clang::sycl_external]] unsigned int test_subgroup_local_invocation_id() {105 return __builtin_spirv_subgroup_local_invocation_id();106}107