Mercurial > hg > CbC > CbC_llvm

annotate clang/test/CodeGenCUDA/flush-denormals.cu @ 266:00f31e85ec16 default tip

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Added tag current for changeset 31d058e83c98
author Shinji KONO <kono@ie.u-ryukyu.ac.jp>
date Sat, 14 Oct 2023 10:13:55 +0900
parents 2e18cbf3894f
children
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
150
anatofuz
parents:
diff changeset
1 // RUN: %clang_cc1 -fcuda-is-device \
anatofuz
parents:
diff changeset
2 // RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
3 // RUN: FileCheck -check-prefixes=NOFTZ,PTXNOFTZ %s
150
anatofuz
parents:
diff changeset
4
anatofuz
parents:
diff changeset
5 // RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=ieee \
anatofuz
parents:
diff changeset
6 // RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
7 // RUN: FileCheck -check-prefixes=NOFTZ,PTXNOFTZ %s
150
anatofuz
parents:
diff changeset
8
anatofuz
parents:
diff changeset
9 // RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \
anatofuz
parents:
diff changeset
10 // RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
11 // RUN: FileCheck -check-prefixes=FTZ,PTXFTZ %s
150
anatofuz
parents:
diff changeset
12
anatofuz
parents:
diff changeset
13 // RUN: %clang_cc1 -fcuda-is-device -x hip \
anatofuz
parents:
diff changeset
14 // RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
15 // RUN: FileCheck -check-prefix=NOFTZ %s
150
anatofuz
parents:
diff changeset
16
anatofuz
parents:
diff changeset
17 // RUN: %clang_cc1 -fcuda-is-device -x hip \
anatofuz
parents:
diff changeset
18 // RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -fdenormal-fp-math-f32=ieee -emit-llvm -o - %s | \
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
19 // RUN: FileCheck -check-prefix=NOFTZ %s
150
anatofuz
parents:
diff changeset
20
anatofuz
parents:
diff changeset
21 // RUN: %clang_cc1 -fcuda-is-device -x hip -fdenormal-fp-math-f32=preserve-sign \
anatofuz
parents:
diff changeset
22 // RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
23 // RUN: FileCheck -check-prefix=FTZ %s
150
anatofuz
parents:
diff changeset
24
anatofuz
parents:
diff changeset
25 #include "Inputs/cuda.h"
anatofuz
parents:
diff changeset
26
anatofuz
parents:
diff changeset
27 // Checks that device function calls get emitted with the "denormal-fp-math-f32"
anatofuz
parents:
diff changeset
28 // attribute set when we compile CUDA device code with
anatofuz
parents:
diff changeset
29 // -fdenormal-fp-math-f32. Further, check that we reflect the presence or
207
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 173
diff changeset
30 // absence of -fgpu-flush-denormals-to-zero in a module flag.
150
anatofuz
parents:
diff changeset
31
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
32 // AMDGCN targets always have f64/f16 denormals enabled.
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
33 //
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
34 // AMDGCN targets without fast FMAF (e.g. gfx803) always have f32 denormal
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
35 // flushing by default.
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
36 //
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
37 // For AMDGCN target with fast FMAF (e.g. gfx900), it has ieee denormals by
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
38 // default and preserve-sign when there with the option
207
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 173
diff changeset
39 // -fgpu-flush-denormals-to-zero.
150
anatofuz
parents:
diff changeset
40
anatofuz
parents:
diff changeset
41 // CHECK-LABEL: define void @foo() #0
anatofuz
parents:
diff changeset
42 extern "C" __device__ void foo() {}
anatofuz
parents:
diff changeset
43
anatofuz
parents:
diff changeset
44 // FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
45 // NOFTZ-NOT: "denormal-fp-math-f32"
150
anatofuz
parents:
diff changeset
46
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
47 // PTXFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
48 // PTXFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
150
anatofuz
parents:
diff changeset
49
173
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
50 // PTXNOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
0572611fdcc8 reorgnization done
Shinji KONO <kono@ie.u-ryukyu.ac.jp>
parents: 150
diff changeset
51 // PTXNOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}