3030// RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll
3131// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s
3232
33+ // Per Thread Stream test cases:
34+
35+ // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
36+ // RUN: -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \
37+ // RUN: -x hip -fhip-new-launch-api \
38+ // RUN: -I%S/../Inputs/ -emit-cir %s -o %t.cir
39+ // RUN: FileCheck --check-prefix=CIR-HOST-PTH --input-file=%t.cir %s
40+
41+ // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
42+ // RUN: -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \
43+ // RUN: -x hip -emit-llvm -fhip-new-launch-api \
44+ // RUN: -I%S/../Inputs/ %s -o %t.ll
45+ // RUN: FileCheck --check-prefix=LLVM-HOST-PTH --input-file=%t.ll %s
46+
47+ // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
48+ // RUN: -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \
49+ // RUN: -x hip -emit-llvm -fhip-new-launch-api \
50+ // RUN: -I%S/../Inputs/ %s -o %t.ll
51+ // RUN: FileCheck --check-prefix=OGCG-HOST-PTH --input-file=%t.ll %s
52+
3353
3454// Attribute for global_fn
3555// CIR-HOST: [[Kernel:#[a-zA-Z_0-9]+]] = {{.*}}#cir.cu.kernel_name<_Z9global_fni>{{.*}}
@@ -53,20 +73,22 @@ __global__ void global_fn(int a) {}
5373// CIR-HOST: %[[#Decayed:]] = cir.cast array_to_ptrdecay %[[#CIRKernelArgs]]
5474// CIR-HOST: cir.call @__hipPopCallConfiguration
5575// CIR-HOST: cir.get_global @_Z9global_fni : !cir.ptr<!cir.ptr<!cir.func<(!s32i)>>>
56- // CIR-HOST: cir.call @hipLaunchKernel
76+ // CIR-HOST-PTH : cir.call @hipLaunchKernel_spt
5777
5878// LLVM-HOST: void @_Z24__device_stub__global_fni
5979// LLVM-HOST: %[[#KernelArgs:]] = alloca [1 x ptr], i64 1, align 16
6080// LLVM-HOST: %[[#GEP1:]] = getelementptr ptr, ptr %[[#KernelArgs]], i32 0
6181// LLVM-HOST: %[[#GEP2:]] = getelementptr [1 x ptr], ptr %[[#KernelArgs]], i32 0, i64 0
6282// LLVM-HOST: call i32 @__hipPopCallConfiguration
6383// LLVM-HOST: call i32 @hipLaunchKernel(ptr @_Z9global_fni
84+ // LLVM-HOST-PTH: call i32 @hipLaunchKernel_spt(ptr @_Z9global_fni
6485//
6586// OGCG-HOST: define dso_local void @_Z24__device_stub__global_fni
6687// OGCG-HOST: %kernel_args = alloca ptr, i64 1, align 16
6788// OGCG-HOST: getelementptr ptr, ptr %kernel_args, i32 0
6889// OGCG-HOST: call i32 @__hipPopCallConfiguration
6990// OGCG-HOST: %call = call noundef i32 @hipLaunchKernel(ptr noundef @_Z9global_fni
91+ // OGCG-HOST-PTH: %call = call noundef i32 @hipLaunchKernel_spt(ptr noundef @_Z9global_fni
7092
7193
7294int main () {
0 commit comments