1717 * launch_inlining.cpp
1818 *
1919 * Description:
20- * Ensure kernels are fully inlined
20+ * Ensure kernels are inlined
2121 **************************************************************************/
22-
2322// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fgpu-inline-threshold=0 %if cl_options %{/clang:-S /clang:-emit-llvm%} %else %{-S -emit-llvm%} %s -o - | FileCheck %s
24- // TODO(joe): update_cc_test_checks.py
23+ // We set -fgpu-inline-threshold=0 to disable heuristic inlining for the purposes of the test
2524#include < sycl/detail/core.hpp>
2625#include < sycl/group_barrier.hpp>
2726#include < syclcompat/launch.hpp>
@@ -33,27 +32,20 @@ namespace sycl_intel_exp = sycl::ext::intel::experimental;
3332
3433static constexpr int LOCAL_MEM_SIZE = 1024 ;
3534
36- template <typename T>
37- __syclcompat_inline__ T dummy_fn (T input){
38- return -input;
39- }
40-
4135// CHECK: define {{.*}}spir_kernel{{.*}}write_mem_kernel{{.*}} {
4236// CHECK-NOT: call {{.*}}write_mem_kernel
43- // CHECK-NOT: call {{.*}}dummy_fn
4437// CHECK: }
4538
4639template <typename T> void write_mem_kernel (T *data, int num_elements) {
4740 const int id =
4841 sycl::ext::oneapi::this_work_item::get_nd_item<1 >().get_global_id (0 );
4942 if (id < num_elements) {
50- data[id] = dummy_fn ( static_cast <T>(id) );
43+ data[id] = static_cast <T>(id);
5144 }
5245};
5346
5447// CHECK: define {{.*}}spir_kernel{{.*}}dynamic_local_mem_typed_kernel{{.*}} {
5548// CHECK-NOT: call {{.*}}dynamic_local_mem_typed_kernel
56- // CHECK-NOT: call {{.*}}dummy_fn
5749// CHECK: }
5850template <typename T>
5951void dynamic_local_mem_typed_kernel (T *data, char *local_mem) {
@@ -67,7 +59,7 @@ void dynamic_local_mem_typed_kernel(T *data, char *local_mem) {
6759 }
6860 sycl::group_barrier (sycl::ext::oneapi::this_work_item::get_work_group<1 >());
6961 if (id < num_elements) {
70- data[id] = dummy_fn ( typed_local_mem[num_elements - id - 1 ]) ;
62+ data[id] = typed_local_mem[num_elements - id - 1 ];
7163 }
7264};
7365
0 commit comments