47 lines · plain
1// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device \2// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCN %s3// RUN: %clang_cc1 -x hip -triple spirv64-amd-amdhsa -fcuda-is-device \4// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCNSPIRV %s5// RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda- -fcuda-is-device \6// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=NVPTX %s7#include "Inputs/cuda.h"8 9struct A {10 int a[32];11 float *p;12};13 14// AMDGCN: define{{.*}} amdgpu_kernel void @_Z6kernel1A(ptr addrspace(4) noundef byref(%struct.A) align 8 %{{.+}})15// AMDGCNSPIRV: define{{.*}} spir_kernel void @_Z6kernel1A(%struct.A %{{.+}})16// NVPTX: define{{.*}} void @_Z6kernel1A(ptr noundef byval(%struct.A) align 8 %x)17__global__ void kernel(A x) {18}19 20class Kernel {21public:22 // AMDGCN: define{{.*}} amdgpu_kernel void @_ZN6Kernel12memberKernelE1A(ptr addrspace(4) noundef byref(%struct.A) align 8 %{{.+}})23 // AMDGCNSPIRV: define{{.*}} spir_kernel void @_ZN6Kernel12memberKernelE1A(%struct.A %{{.+}})24 // NVPTX: define{{.*}} void @_ZN6Kernel12memberKernelE1A(ptr noundef byval(%struct.A) align 8 %x)25 static __global__ void memberKernel(A x){}26 template<typename T> static __global__ void templateMemberKernel(T x) {}27};28 29 30template <typename T>31__global__ void templateKernel(T x) {}32 33void launch(void*);34 35void test() {36 Kernel K;37 // AMDGCN: define{{.*}} amdgpu_kernel void @_Z14templateKernelI1AEvT_(ptr addrspace(4) noundef byref(%struct.A) align 8 %{{.+}}38 // AMDGCNSPIRV: define{{.*}} spir_kernel void @_Z14templateKernelI1AEvT_(%struct.A %{{.+}})39 // NVPTX: define{{.*}} void @_Z14templateKernelI1AEvT_(ptr noundef byval(%struct.A) align 8 %x)40 launch((void*)templateKernel<A>);41 42 // AMDGCN: define{{.*}} amdgpu_kernel void @_ZN6Kernel20templateMemberKernelI1AEEvT_(ptr addrspace(4) noundef byref(%struct.A) align 8 %{{.+}}43 // AMDGCNSPIRV: define{{.*}} spir_kernel void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A %{{.+}}44 // NVPTX: define{{.*}} void @_ZN6Kernel20templateMemberKernelI1AEEvT_(ptr noundef byval(%struct.A) align 8 %x)45 launch((void*)Kernel::templateMemberKernel<A>);46}47