Removed always-inline attribute from stubs and made required littest changes

Author: anikelal

clang/lib/AST/MicrosoftMangle.cpp

@@ -1163,7 +1163,7 @@ void MicrosoftCXXNameMangler::mangleUnqualifiedName(GlobalDecl GD,
                   ->hasAttr<CUDAGlobalAttr>())) &&
             GD.getKernelReferenceKind() == KernelReferenceKind::Stub;
         bool IsOCLDeviceStub =
-            ND && (isa<FunctionDecl>(ND) && ND->hasAttr<OpenCLKernelAttr>()) &&
+            ND && isa<FunctionDecl>(ND) && ND->hasAttr<OpenCLKernelAttr>() &&
             GD.getKernelReferenceKind() == KernelReferenceKind::Stub;
         if (IsDeviceStub)
           mangleSourceName(

clang/lib/CodeGen/CGCall.cpp

@@ -2502,7 +2502,12 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
                                  NumElemsParam);
     }
 
-    if (TargetDecl->hasAttr<OpenCLKernelAttr>()) {
+    if (TargetDecl->hasAttr<OpenCLKernelAttr>() &&
+        CallingConv != CallingConv::CC_C &&
+        CallingConv !=
+            CallingConv::CC_SpirFunction) { // Check CallingConv to avoid adding
+                                            // uniform-work-group-size attribute
+                                            // to OpenCL Kernel Stub
       if (getLangOpts().OpenCLVersion <= 120) {
         // OpenCL v1.2 Work groups are always uniform
         FuncAttrs.addAttribute("uniform-work-group-size", "true");

clang/lib/CodeGen/CGExpr.cpp

@@ -6133,11 +6133,9 @@ RValue CodeGenFunction::EmitCall(QualType CalleeType,
 
   const auto *FnType = cast<FunctionType>(PointeeType);
 
-  if (auto FD = dyn_cast_or_null<FunctionDecl>(TargetDecl)) {
-    if (FD->hasAttr<OpenCLKernelAttr>()) {
-      CGM.getTargetCodeGenInfo().setOCLKernelStubCallingConvention(FnType);
-    }
-  }
+  if (const auto *FD = dyn_cast_or_null<FunctionDecl>(TargetDecl);
+      FD && FD->hasAttr<OpenCLKernelAttr>())
+    CGM.getTargetCodeGenInfo().setOCLKernelStubCallingConvention(FnType);
 
   // If we are checking indirect calls and this call is indirect, check that the
   // function pointer is a member of the bit set for the function type.

clang/lib/CodeGen/CodeGenModule.cpp

@@ -6151,16 +6151,6 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
 
   setNonAliasAttributes(GD, Fn);
 
-  if (D->hasAttr<OpenCLKernelAttr>()) {
-    if (GD.getKernelReferenceKind() == KernelReferenceKind::Stub) {
-      if (Fn->hasFnAttribute(llvm::Attribute::NoInline))
-        Fn->removeFnAttr(llvm::Attribute::NoInline);
-      if (Fn->hasFnAttribute(llvm::Attribute::InlineHint))
-        Fn->removeFnAttr(llvm::Attribute::InlineHint);
-      Fn->addFnAttr(llvm::Attribute::AlwaysInline);
-    }
-  }
-
   SetLLVMFunctionAttributesForDefinition(D, Fn);
 
   if (const ConstructorAttr *CA = D->getAttr<ConstructorAttr>())

clang/test/CodeGenOpenCL/addr-space-struct-arg.cl

@@ -5,6 +5,9 @@
 
 kernel void ker() {};
 // CHECK: define{{.*}}@ker() #[[ATTR0:[0-9]+]]
+// CHECK: call void @__clang_ocl_kern_imp_ker() #[[ATTR2:[0-9]+]]
+
+// CHECK: define{{.*}}@__clang_ocl_kern_imp_ker() #[[ATTR1:[0-9]+]]
 
 void foo() {};
 // CHECK: define{{.*}}@foo() #[[ATTR1:[0-9]+]]
@@ -15,3 +18,6 @@ void foo() {};
 
 // CHECK: attributes #[[ATTR1]]
 // CHECK-NOT: uniform-work-group-size
+
+// CHECK: attributes #[[ATTR2]]
+// CHECK-NOT: uniform-work-group-size

clang/test/CodeGenOpenCL/amdgpu-abi-struct-arg-byref.cl

@@ -1,12 +1,12 @@
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefix=CHECK-LIFETIMES
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR,TRIPLESPIR 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR,TRIPLESPIR 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=CHECK-LIFETIMES,TRIPLESPIR 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR,TRIPLESPIR 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR,TRIPLESPIR 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=CHECK-LIFETIMES,TRIPLESPIR 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86,TRIPLEX86 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86,TRIPLEX86 
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=CHECK-LIFETIMES,TRIPLEX86 
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
@@ -39,6 +39,12 @@ void callee(int id, __global int *out) {
   out[id] = id;
 }
 
+// TRIPLESPIR: define{{.*}} void @device_side_enqueue(ptr addrspace(1) align 4 %{{.*}}, ptr addrspace(1) align 4 %b, i32 %i)
+// TRIPLESPIR:    call spir_func void @__clang_ocl_kern_imp_device_side_enqueue({{.*}})
+
+// TRIPLEX86: define{{.*}} void @device_side_enqueue(ptr addrspace(1) align 4 %{{.*}}, ptr addrspace(1) align 4 %b, i32 %i)
+// TRIPLEX86:    call void @__clang_ocl_kern_imp_device_side_enqueue({{.*}})
+
 // COMMON-LABEL: define{{.*}} void @__clang_ocl_kern_imp_device_side_enqueue(ptr addrspace(1) align 4 %{{.*}}, ptr addrspace(1) align 4 %b, i32 %i)
 kernel void device_side_enqueue(global int *a, global int *b, int i) {
   // SPIR: %default_queue = alloca target("spirv.Queue")

clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl

@@ -127,7 +127,7 @@ void test_not_unroll() {
 // CHECK: declare spir_func void @nodupfun(){{[^#]*}} #[[attr3:[0-9]+]]
 
 // CHECK-LABEL: @assume_convergent_asm
-// CHECK: tail call void asm sideeffect "s_barrier", ""() #6
+// CHECK: tail call void asm sideeffect "s_barrier", ""() #5
 kernel void assume_convergent_asm()
 {
   __asm__ volatile("s_barrier");
@@ -138,7 +138,6 @@ kernel void assume_convergent_asm()
 // CHECK: attributes #2 = { {{[^}]*}}convergent{{[^}]*}} }
 // CHECK: attributes #3 = { {{[^}]*}}convergent noduplicate{{[^}]*}} }
 // CHECK: attributes #4 = { {{[^}]*}}convergent{{[^}]*}} }
-// CHECK: attributes #5 = { {{[^}]*}}alwaysinline convergent{{[^}]*}} }
-// CHECK: attributes #6 = { {{[^}]*}}convergent{{[^}]*}} }
-// CHECK: attributes #7 = { {{[^}]*}}nounwind{{[^}]*}} }
-// CHECK: attributes #8 = { {{[^}]*}}convergent noduplicate nounwind{{[^}]*}} }
+// CHECK: attributes #5 = { {{[^}]*}}convergent{{[^}]*}} }
+// CHECK: attributes #6 = { {{[^}]*}}nounwind{{[^}]*}} }
+// CHECK: attributes #7 = { {{[^}]*}}convergent noduplicate nounwind{{[^}]*}} }

clang/test/CodeGenOpenCL/cl-uniform-wg-size.cl

@@ -9,8 +9,15 @@
 typedef struct {int a;} ndrange_t;
 
 kernel void test(int i) {
+
 // AMDGPU-LABEL: define {{.*}} amdgpu_kernel void @test
+// AMDGPU-LABEL: call void @__clang_ocl_kern_imp_test(i32 noundef %0)
+
 // SPIR-LABEL: define {{.*}} spir_kernel void @test
+// SPIR-LABEL: call spir_func void @__clang_ocl_kern_imp_test(i32 noundef %0)
+
+// AMDGPU-LABEL: define {{.*}} void @__clang_ocl_kern_imp_test
+// SPIR-LABEL: define {{.*}} spir_func void @__clang_ocl_kern_imp_test
 
 // COMMON-LABEL: entry:
 // AMDGPU: %block_sizes = alloca [1 x i64]
@@ -37,5 +44,5 @@ kernel void test(int i) {
 
 // CHECK-DEBUG: ![[TESTFILE:[0-9]+]] = !DIFile(filename: "<stdin>"
 // CHECK-DEBUG: ![[TESTSCOPE:[0-9]+]] = distinct !DISubprogram(name: "test", linkageName: "__clang_ocl_kern_imp_test", {{.*}} file: ![[TESTFILE]]
-// CHECK-DEBUG: ![[IFSCOPE:[0-9]+]] = distinct !DILexicalBlock(scope: ![[TESTSCOPE]], file: ![[TESTFILE]], line: 26)
-// CHECK-DEBUG: ![[TEMPLOCATION]] = !DILocation(line: 27, scope: ![[IFSCOPE]])
+// CHECK-DEBUG: ![[IFSCOPE:[0-9]+]] = distinct !DILexicalBlock(scope: ![[TESTSCOPE]], file: ![[TESTFILE]], line: 33)
+// CHECK-DEBUG: ![[TEMPLOCATION]] = !DILocation(line: 34, scope: ![[IFSCOPE]])

clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl

@@ -57,7 +57,7 @@ __kernel void use_of_local_var()
 // CHECK-LABEL: define dso_local amdgpu_kernel void @use_of_local_var(
 // CHECK-SAME: ) #[[ATTR3:[0-9]+]] !kernel_arg_addr_space [[META8:![0-9]+]] !kernel_arg_access_qual [[META8]] !kernel_arg_type [[META8]] !kernel_arg_base_type [[META8]] !kernel_arg_type_qual [[META8]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
-// CHECK-NEXT:    call void @__clang_ocl_kern_imp_use_of_local_var() #[[ATTR7:[0-9]+]]
+// CHECK-NEXT:    call void @__clang_ocl_kern_imp_use_of_local_var() #[[ATTR6]]
 // CHECK-NEXT:    ret void
 //
 //