1 // RUN: echo -n "GPU binary would be here." > %t
2 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
3 // RUN: -target-sdk-version=11.0 -fcuda-include-gpubinary %t -o - \
4 // RUN: | FileCheck %s --check-prefixes CUDA
5 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x hip \
6 // RUN: -fcuda-include-gpubinary %t -o - \
7 // RUN: | FileCheck %s --check-prefixes HIP
8
9 #include "Inputs/cuda.h"
10
11 template <typename T>
12 struct S { T t; };
13
14 template <typename T>
Kernel(S<T>)15 static __global__ void Kernel(S<T>) {}
16
17 // For some reason it takes three or more instantiations of Kernel to trigger a
18 // crash during CUDA compilation.
19 auto x = &Kernel<double>;
20 auto y = &Kernel<float>;
21 auto z = &Kernel<int>;
22
23 // This triggers HIP-specific code path.
func()24 void func (){
25 Kernel<short><<<1,1>>>({1});
26 }
27
28 // CUDA-LABEL: @__cuda_register_globals(
29 // CUDA: call i32 @__cudaRegisterFunction(ptr %0, ptr @_ZL21__device_stub__KernelIdEv1SIT_E
30 // CUDA: call i32 @__cudaRegisterFunction(ptr %0, ptr @_ZL21__device_stub__KernelIfEv1SIT_E
31 // CUDA: call i32 @__cudaRegisterFunction(ptr %0, ptr @_ZL21__device_stub__KernelIiEv1SIT_E
32 // CUDA: ret void
33
34 // HIP-LABEL: @__hip_register_globals(
35 // HIP: call i32 @__hipRegisterFunction(ptr %0, ptr @_ZL6KernelIdEv1SIT_E
36 // HIP: call i32 @__hipRegisterFunction(ptr %0, ptr @_ZL6KernelIfEv1SIT_E
37 // HIP: call i32 @__hipRegisterFunction(ptr %0, ptr @_ZL6KernelIiEv1SIT_E
38 // HIP: ret void
39