100 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 52// REQUIRES: x86-registered-target3// REQUIRES: nvptx-registered-target4 5// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \6// RUN: -disable-llvm-passes -o - %s | FileCheck -allow-deprecated-dag-overlap -check-prefix DEVICE %s7 8// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \9// RUN: -disable-llvm-passes -o - %s | \10// RUN: FileCheck -allow-deprecated-dag-overlap -check-prefix HOST %s11 12#include "Inputs/cuda.h"13 14// DEVICE-LABEL: define dso_local void @_Z3foov(15// DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {16// DEVICE-NEXT: [[ENTRY:.*:]]17// DEVICE-NEXT: ret void18//19__device__ void foo() {}20// DEVICE-LABEL: define dso_local void @_Z3baxv(21// DEVICE-SAME: ) #[[ATTR1:[0-9]+]] {22// DEVICE-NEXT: [[ENTRY:.*:]]23// DEVICE-NEXT: ret void24//25[[clang::noconvergent]] __device__ void bax() {}26 27__host__ __device__ void baz();28 29__host__ __device__ float aliasf0(int) asm("something");30__host__ __device__ [[clang::noconvergent]] float aliasf1(int) asm("somethingelse");31 32// DEVICE-LABEL: define dso_local void @_Z3barv(33// DEVICE-SAME: ) #[[ATTR0]] {34// DEVICE-NEXT: [[ENTRY:.*:]]35// DEVICE-NEXT: [[X:%.*]] = alloca i32, align 436// DEVICE-NEXT: call void @_Z3bazv() #[[ATTR4:[0-9]+]]37// DEVICE-NEXT: [[TMP0:%.*]] = call i32 asm "trap", "=l"() #[[ATTR5:[0-9]+]], !srcloc [[META3:![0-9]+]]38// DEVICE-NEXT: store i32 [[TMP0]], ptr [[X]], align 439// DEVICE-NEXT: call void asm sideeffect "trap", ""() #[[ATTR4]], !srcloc [[META4:![0-9]+]]40// DEVICE-NEXT: call void asm sideeffect "nop", ""() #[[ATTR6:[0-9]+]], !srcloc [[META5:![0-9]+]]41// DEVICE-NEXT: [[TMP1:%.*]] = load i32, ptr [[X]], align 442// DEVICE-NEXT: [[CALL:%.*]] = call contract noundef float @something(i32 noundef [[TMP1]]) #[[ATTR4]]43// DEVICE-NEXT: [[TMP2:%.*]] = load i32, ptr [[X]], align 444// DEVICE-NEXT: [[CALL1:%.*]] = call contract noundef float @somethingelse(i32 noundef [[TMP2]]) #[[ATTR6]]45// DEVICE-NEXT: ret void46//47// HOST-LABEL: define dso_local void @_Z3barv(48// HOST-SAME: ) #[[ATTR0:[0-9]+]] {49// HOST-NEXT: [[ENTRY:.*:]]50// HOST-NEXT: [[X:%.*]] = alloca i32, align 451// HOST-NEXT: call void @_Z3bazv()52// HOST-NEXT: [[TMP0:%.*]] = call i32 asm "trap", "=l,~{dirflag},~{fpsr},~{flags}"() #[[ATTR2:[0-9]+]], !srcloc [[META2:![0-9]+]]53// HOST-NEXT: store i32 [[TMP0]], ptr [[X]], align 454// HOST-NEXT: call void asm sideeffect "trap", "~{dirflag},~{fpsr},~{flags}"() #[[ATTR3:[0-9]+]], !srcloc [[META3:![0-9]+]]55// HOST-NEXT: call void asm sideeffect "nop", "~{dirflag},~{fpsr},~{flags}"() #[[ATTR3]], !srcloc [[META4:![0-9]+]]56// HOST-NEXT: [[TMP1:%.*]] = load i32, ptr [[X]], align 457// HOST-NEXT: [[CALL:%.*]] = call contract noundef float @something(i32 noundef [[TMP1]])58// HOST-NEXT: [[TMP2:%.*]] = load i32, ptr [[X]], align 459// HOST-NEXT: [[CALL1:%.*]] = call contract noundef float @somethingelse(i32 noundef [[TMP2]])60// HOST-NEXT: ret void61//62__host__ __device__ void bar() {63 baz();64 int x;65 asm ("trap" : "=l"(x));66 asm volatile ("trap");67 [[clang::noconvergent]] { asm volatile ("nop"); }68 aliasf0(x);69 aliasf1(x);70}71 72 73//.74// DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }75// DEVICE: attributes #[[ATTR1]] = { mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }76// DEVICE: attributes #[[ATTR2:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }77// DEVICE: attributes #[[ATTR3:[0-9]+]] = { nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }78// DEVICE: attributes #[[ATTR4]] = { convergent nounwind }79// DEVICE: attributes #[[ATTR5]] = { convergent nounwind memory(none) }80// DEVICE: attributes #[[ATTR6]] = { nounwind }81//.82// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }83// HOST: attributes #[[ATTR1:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }84// HOST: attributes #[[ATTR2]] = { nounwind memory(none) }85// HOST: attributes #[[ATTR3]] = { nounwind }86//.87// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}88// DEVICE: [[META1:![0-9]+]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}89// DEVICE: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}90// DEVICE: [[META3]] = !{i64 3120}91// DEVICE: [[META4]] = !{i64 3155}92// DEVICE: [[META5]] = !{i64 3206}93//.94// HOST: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}95// HOST: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}96// HOST: [[META2]] = !{i64 3120}97// HOST: [[META3]] = !{i64 3155}98// HOST: [[META4]] = !{i64 3206}99//.100