yaxunl created this revision.
yaxunl added a reviewer: tra.
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.
https://reviews.llvm.org/D54183
Files:
lib/Frontend/CompilerInvocation.cpp
test/CodeGenCUDA/device-stub.cu
test/CodeGenCUDA/flush-denormals.cu
Index: test/CodeGenCUDA/flush-denormals.cu
===================================================================
--- test/CodeGenCUDA/flush-denormals.cu
+++ test/CodeGenCUDA/flush-denormals.cu
@@ -25,7 +25,7 @@
// 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"
Index: test/CodeGenCUDA/device-stub.cu
===================================================================
--- test/CodeGenCUDA/device-stub.cu
+++ test/CodeGenCUDA/device-stub.cu
@@ -11,16 +11,16 @@
// 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"
Index: lib/Frontend/CompilerInvocation.cpp
===================================================================
--- lib/Frontend/CompilerInvocation.cpp
+++ lib/Frontend/CompilerInvocation.cpp
@@ -123,6 +123,8 @@
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))
Index: test/CodeGenCUDA/flush-denormals.cu
===================================================================
--- test/CodeGenCUDA/flush-denormals.cu
+++ test/CodeGenCUDA/flush-denormals.cu
@@ -25,7 +25,7 @@
// 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"
Index: test/CodeGenCUDA/device-stub.cu
===================================================================
--- test/CodeGenCUDA/device-stub.cu
+++ test/CodeGenCUDA/device-stub.cu
@@ -11,16 +11,16 @@
// 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"
Index: lib/Frontend/CompilerInvocation.cpp
===================================================================
--- lib/Frontend/CompilerInvocation.cpp
+++ lib/Frontend/CompilerInvocation.cpp
@@ -123,6 +123,8 @@
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))
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
