[HIP] Change default optimization level to -O3

[yaxunl]

Summary

The default optimization level of nvcc is -O3. There are HIP applications which expect
the default optimization level to be -O3. Most HIP applications use -O3, therefore
making it default.

[tra]

I'm not convinced that nvcc's behavior is a good guide for clang's defaults.

Considering that clang is not compatible with nvcc when it comes to command line options, whoever is using clang to compile CUDA already has to adjust command line options. Explicitly adding -O3 during the process is, IMO, better than requiring everyone to be aware that if the input file is CUDA (or HIP), then clang would default to -O3.

lib/Frontend/CompilerInvocation.cpp

 // Deserialization (from args)
 //===----------------------------------------------------------------------===//
 
 static unsigned getOptimizationLevel(ArgList &Args, InputKind IK,
                                      DiagnosticsEngine &Diags) {
   unsigned DefaultOpt = 0;
   if (IK.getLanguage() == InputKind::OpenCL && !Args.hasArg(OPT_cl_opt_disable))
     DefaultOpt = 2;
+  else if (IK.getLanguage() == InputKind::HIP)
+    DefaultOpt = 3;
 
   if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
     if (A->getOption().matches(options::OPT_O0))
       return 0;
 
     if (A->getOption().matches(options::OPT_Ofast))
       return 3;
 

test/CodeGenCUDA/device-stub.cu

 // RUN: echo "GPU binary would be here" > %t
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fcuda-include-gpubinary %t -o - \
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,CUDA,CUDANORDC
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,CUDANOGLOBALS
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fgpu-rdc -fcuda-include-gpubinary %t -o - \
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,RDC,CUDA,CUDARDC
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - \
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
 
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: %clang_cc1 -O0 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP,HIPEF
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: %clang_cc1 -O0 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS -x hip \
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: %clang_cc1 -O0 -triple x86_64-linux-gnu -emit-llvm %s \
 // RUN:     -fgpu-rdc -fcuda-include-gpubinary %t -o - -x hip \
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP,HIPEF
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
+// RUN: %clang_cc1 -O0 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=ALL,NORDC,HIP,HIPNEF
 
 #include "Inputs/cuda.h"
 
 #ifndef NOGLOBALS
 // ALL-DAG: @device_var = internal global i32
 __device__ int device_var;
 

test/CodeGenCUDA/flush-denormals.cu

 // or absence of -fcuda-flush-denormals-to-zero in a module flag.
 
 // AMDGCN targets always have +fp64-fp16-denormals.
 // AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals.
 // For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals
 // by default and -fp32-denormals when there is option
 // -fcuda-flush-denormals-to-zero.
 
-// CHECK-LABEL: define void @foo() #0
+// CHECK-LABEL: define void @foo(){{.*}}#0
 extern "C" __device__ void foo() {}
 
 // FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
 // NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
 // AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals
 // AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals
 
 // FTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
 // FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
 
 // NOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
 // NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}