136 lines · plain
1// RUN: %clang_cc1 -emit-llvm -x hip %s -o - -triple x86_64-linux-gnu \2// RUN: | FileCheck -check-prefixes=CHECK,HOST %s3// RUN: %clang_cc1 -emit-llvm -x hip %s -o - -triple amdgcn-amd-amdhsa -fcuda-is-device \4// RUN: | FileCheck -check-prefixes=CHECK,DEV %s5 6#include "Inputs/cuda.h"7 8// CHECK: %class.anon = type { ptr, float, ptr, ptr }9// CHECK: %class.anon.0 = type { ptr, float, ptr, ptr }10// CHECK: %class.anon.1 = type { ptr, ptr, ptr }11// CHECK: %class.anon.2 = type { ptr, float, ptr, ptr }12 13// HOST: call void @_ZN8DevByVal21__device_stub__kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr noundef byval(%class.anon)14// DEV: define amdgpu_kernel void @_ZN8DevByVal6kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr addrspace(4) noundef byref(%class.anon)15 16// Only the device function passes arugments by value.17namespace DevByVal {18__device__ float fun(float x, float y) {19 return x;20}21 22float fun(const float &x, const float &y) {23 return x;24}25 26template<typename F>27void __global__ kernel(F f)28{29 f(1);30}31 32void test(float const * fl, float const * A, float * Vf)33{34 float constexpr small(1.0e-25);35 36 auto lambda = [=] __device__ __host__ (unsigned int n) {37 float const value = fun(small, fl[0]);38 Vf[0] = value * A[0];39 };40 kernel<<<1, 1>>>(lambda);41}42}43 44// HOST: call void @_ZN9HostByVal21__device_stub__kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr noundef byval(%class.anon.0)45// DEV: define amdgpu_kernel void @_ZN9HostByVal6kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr addrspace(4) noundef byref(%class.anon.0)46 47// Only the host function passes arugments by value.48namespace HostByVal {49float fun(float x, float y) {50 return x;51}52 53__device__ float fun(const float &x, const float &y) {54 return x;55}56 57template<typename F>58void __global__ kernel(F f)59{60 f(1);61}62 63void test(float const * fl, float const * A, float * Vf)64{65 float constexpr small(1.0e-25);66 67 auto lambda = [=] __device__ __host__ (unsigned int n) {68 float const value = fun(small, fl[0]);69 Vf[0] = value * A[0];70 };71 kernel<<<1, 1>>>(lambda);72}73}74 75// HOST: call void @_ZN9BothByVal21__device_stub__kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr noundef byval(%class.anon.1)76// DEV: define amdgpu_kernel void @_ZN9BothByVal6kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr addrspace(4) noundef byref(%class.anon.1)77 78// Both the host and device functions pass arugments by value.79namespace BothByVal {80float fun(float x, float y) {81 return x;82}83 84__device__ float fun(float x, float y) {85 return x;86}87 88template<typename F>89void __global__ kernel(F f)90{91 f(1);92}93 94void test(float const * fl, float const * A, float * Vf)95{96 float constexpr small(1.0e-25);97 98 auto lambda = [=] __device__ __host__ (unsigned int n) {99 float const value = fun(small, fl[0]);100 Vf[0] = value * A[0];101 };102 kernel<<<1, 1>>>(lambda);103}104}105 106// HOST: call void @_ZN12NeitherByVal21__device_stub__kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr noundef byval(%class.anon.2)107// DEV: define amdgpu_kernel void @_ZN12NeitherByVal6kernelIZNS_4testEPKfS2_PfEUljE_EEvT_(ptr addrspace(4) noundef byref(%class.anon.2)108 109// Neither the host nor device function passes arugments by value.110namespace NeitherByVal {111float fun(const float& x, const float& y) {112 return x;113}114 115__device__ float fun(const float& x, const float& y) {116 return x;117}118 119template<typename F>120void __global__ kernel(F f)121{122 f(1);123}124 125void test(float const * fl, float const * A, float * Vf)126{127 float constexpr small(1.0e-25);128 129 auto lambda = [=] __device__ __host__ (unsigned int n) {130 float const value = fun(small, fl[0]);131 Vf[0] = value * A[0];132 };133 kernel<<<1, 1>>>(lambda);134}135}136