This is an archive of the discontinued LLVM Phabricator instance.

Differential D73651

[OpenCL][CUDA][HIP][SYCL] Add norecurse
ClosedPublic

Authored by yaxunl on Jan 29 2020, 11:21 AM.

Details

Reviewers
Anastasia
tra
bader
keryell
Commits
rGfb44b9db95a3: [OpenCL][CUDA][HIP][SYCL] Add norecurse
Summary

norecurse function attr indicates the function is not called recursively directly or indirectly.

Add norecurse to OpenCL functions, SYCL functions in device compilation and CUDA/HIP kernels.

Although there is LLVM pass adding norecurse to functions, it only works for whole-program compilation. Also FE adding norecurse can make that pass run faster since functions with norecurse do not need to be checked again.

Diff Detail

Event Timeline

yaxunl created this revision.Jan 29 2020, 11:21 AM
tra accepted this revision.Jan 29 2020, 11:29 AM

LGTM for CUDA.

clang/lib/CodeGen/CodeGenFunction.cpp
929

I believe dynamic parallelism is more like spawning a separate process, not recursion, so we should be OK.

This revision is now accepted and ready to land.Jan 29 2020, 11:29 AM
bader added a subscriber: bader.Jan 29 2020, 11:41 AM
bader added inline comments.
clang/lib/CodeGen/CodeGenFunction.cpp
934

SYCL doesn't support recursion neither. Does it make sense to add it too?

yaxunl updated this revision to Diff 241256.Jan 29 2020, 12:40 PM
yaxunl retitled this revision from [OpenCL][CUDA][HIP] Add norecurse to [OpenCL][CUDA][HIP][SYCL] Add norecurse.
yaxunl edited the summary of this revision. (Show Details)

Added handling of SYCL kernels by Alexey's comments. I cannot add a codegen test for SYCL since I cannot find a way to instantiate a SYCL kernel.

Herald added a subscriber: ebevhan. · View Herald TranscriptJan 29 2020, 12:40 PM
bader added a comment.Jan 29 2020, 12:49 PM
In D73651#1847862, @yaxunl wrote:

Added handling of SYCL kernels by Alexey's comments. I cannot add a codegen test for SYCL since I cannot find a way to instantiate a SYCL kernel.

I think SYCL implementation should follow OpenCL logic and apply attribute to any function, not only kernels.

yaxunl updated this revision to Diff 241265.Jan 29 2020, 1:01 PM
yaxunl edited the summary of this revision. (Show Details)

revised by Alexey's comments.

t-tye added a subscriber: t-tye.Jan 29 2020, 3:02 PM
t-tye added inline comments.
clang/lib/CodeGen/CodeGenFunction.cpp
929

Right, dynamic parallelism is equivalent to spawn a new thread to execute the function. Recursion is only about the same thread calling the function. So I agree that the TODO about it can be removed.

yaxunl updated this revision to Diff 241335.Jan 29 2020, 6:31 PM

Remove the ToDo for CUDA/HIP.

bader accepted this revision.Jan 29 2020, 11:53 PM

Thanks!

rampitec added a subscriber: rampitec.Jan 30 2020, 10:12 AM
keryell requested changes to this revision.Jan 31 2020, 7:05 PM
keryell added a subscriber: keryell.
keryell added inline comments.
clang/lib/CodeGen/CodeGenFunction.cpp
929

Can you change the SYCL version to "SYCL 1.2.1 s3.10"?
SYCL 2.2 is non normative: it was just a provisional version which is now deprecated and the SYCL committee is working on a newer version.
Thanks.

This revision now requires changes to proceed.Jan 31 2020, 7:05 PM
yaxunl updated this revision to Diff 241877.Feb 1 2020, 6:26 AM

Fix SYCL version.

Anastasia added inline comments.Feb 4 2020, 3:04 AM
clang/lib/CodeGen/CodeGenFunction.cpp
925

Can we remove reference to OpenCL C++ since it's not implemented. Feel free to add C++ for OpenCL instead. The comment explanation still makes sense. You can just remove references to docs.

yaxunl updated this revision to Diff 242300.Feb 4 2020, 5:09 AM
yaxunl marked 5 inline comments as done.

revised by Anastasia's comments.

Anastasia accepted this revision.Feb 5 2020, 5:19 AM

LGTM! Thanks!

This revision was not accepted when it landed; it landed in state Needs Review.Feb 16 2020, 5:53 PM
Closed by commit rGfb44b9db95a3: [OpenCL][CUDA][HIP][SYCL] Add norecurse (authored by yaxunl). · Explain Why
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptFeb 16 2020, 5:53 PM
Herald added subscribers: kerbowa, nhaehnle, jvesely. · View Herald Transcript

Revision Contents

PathSize
clang/
lib/
CodeGen/
18 lines
test/
CodeGenCUDA/
15 lines
25 lines
CodeGenOpenCL/
60 lines
6 lines
SemaCUDA/
9 lines

Diff 244906

clang/lib/CodeGen/CodeGenFunction.cpp

Show First 20 Lines • Show All 912 Lines • ▼ Show 20 Lines if (Nullability && *Nullability == NullabilityKind::NonNull) {
RetValNullabilityPrecondition = RetValNullabilityPrecondition =
llvm::ConstantInt::getTrue(getLLVMContext()); llvm::ConstantInt::getTrue(getLLVMContext());
} }
} }
// If we're in C++ mode and the function name is "main", it is guaranteed // If we're in C++ mode and the function name is "main", it is guaranteed
// to be norecurse by the standard (3.6.1.3 "The function main shall not be // to be norecurse by the standard (3.6.1.3 "The function main shall not be
// used within a program"). // used within a program").
if (getLangOpts().CPlusPlus) //
if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D)) // OpenCL C 2.0 v2.2-11 s6.9.i:
if (FD->isMain()) // Recursion is not supported.
//
// SYCL v1.2.1 s3.10:
AnastasiaUnsubmitted
Done
ReplyInline Actions

Can we remove reference to OpenCL C++ since it's not implemented. Feel free to add C++ for OpenCL instead. The comment explanation still makes sense. You can just remove references to docs.

Anastasia: Can we remove reference to OpenCL C++ since it's not implemented. Feel free to add C++ for…
// kernels cannot include RTTI information, exception classes,
// recursive code, virtual functions or make use of C++ libraries that
// are not compiled for the device.
if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D)) {
traUnsubmitted
Done
ReplyInline Actions

I believe dynamic parallelism is more like spawning a separate process, not recursion, so we should be OK.

tra: I believe dynamic parallelism is more like spawning a separate process, not recursion, so we…
t-tyeUnsubmitted
Done
ReplyInline Actions

Right, dynamic parallelism is equivalent to spawn a new thread to execute the function. Recursion is only about the same thread calling the function. So I agree that the TODO about it can be removed.

t-tye: Right, dynamic parallelism is equivalent to spawn a new thread to execute the function.
keryellUnsubmitted
Done
ReplyInline Actions

Can you change the SYCL version to "SYCL 1.2.1 s3.10"?
SYCL 2.2 is non normative: it was just a provisional version which is now deprecated and the SYCL committee is working on a newer version.
Thanks.

keryell: Can you change the SYCL version to "SYCL 1.2.1 s3.10"? SYCL 2.2 is non normative: it was just a…
if ((getLangOpts().CPlusPlus && FD->isMain()) || getLangOpts().OpenCL ||
getLangOpts().SYCLIsDevice ||
(getLangOpts().CUDA && FD->hasAttr<CUDAGlobalAttr>()))
Fn->addFnAttr(llvm::Attribute::NoRecurse); Fn->addFnAttr(llvm::Attribute::NoRecurse);
}
baderUnsubmitted
Done
ReplyInline Actions

SYCL doesn't support recursion neither. Does it make sense to add it too?

bader: SYCL doesn't support recursion neither. Does it make sense to add it too?
if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D)) if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
if (FD->usesFPIntrin()) if (FD->usesFPIntrin())
Fn->addFnAttr(llvm::Attribute::StrictFP); Fn->addFnAttr(llvm::Attribute::StrictFP);
// If a custom alignment is used, force realigning to this alignment on // If a custom alignment is used, force realigning to this alignment on
// any main function which certainly will need it. // any main function which certainly will need it.
if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D)) if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
▲ Show 20 Lines • Show All 1,587 Lines • Show Last 20 Lines

clang/test/CodeGenCUDA/norecurse.cu

  • This file was added.
// REQUIRES: amdgpu-registered-target
// REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device \
// RUN: -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
// RUN: -emit-llvm -disable-llvm-passes -o - -x hip %s | FileCheck %s
#include "Inputs/cuda.h"
__global__ void kernel1(int a) {}
// CHECK: define{{.*}}@_Z7kernel1i{{.*}}#[[ATTR:[0-9]*]]
// CHECK: attributes #[[ATTR]] = {{.*}}norecurse

clang/test/CodeGenCUDA/propagate-metadata.cu

Show First 20 Lines • Show All 42 Lines • ▼ Show 20 Lines
#include "Inputs/cuda.h" #include "Inputs/cuda.h"
__device__ void lib_fn(); __device__ void lib_fn();
__global__ void kernel() { lib_fn(); } __global__ void kernel() { lib_fn(); }
#endif #endif
} }
// The kernel and lib function should have the same attributes. // The kernel and lib function should have the same attributes.
// CHECK: define void @kernel() [[attr:#[0-9]+]] // CHECK: define void @kernel() [[kattr:#[0-9]+]]
// CHECK: define internal void @lib_fn() [[attr]] // CHECK: define internal void @lib_fn() [[fattr:#[0-9]+]]
// FIXME: These -NOT checks do not work as intended and do not check on the same // FIXME: These -NOT checks do not work as intended and do not check on the same
// line. // line.
// Check the attribute list. // Check the attribute list for kernel.
// CHECK: attributes [[attr]] = { // CHECK: attributes [[kattr]] = {
// CHECK-SAME: convergent // CHECK-SAME: convergent
// CHECK-SAME: norecurse
// FTZ-NOT: "denormal-fp-math"
// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
// CHECK-SAME: "no-trapping-math"="true"
// FAST-SAME: "unsafe-fp-math"="true"
// NOFAST-NOT: "unsafe-fp-math"="true"
// Check the attribute list for lib_fn.
// CHECK: attributes [[fattr]] = {
// CHECK-SAME: convergent
// CHECK-NOT: norecurse
// FTZ-NOT: "denormal-fp-math" // FTZ-NOT: "denormal-fp-math"
// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign" // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee" // NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
// CHECK-SAME: "no-trapping-math"="true" // CHECK-SAME: "no-trapping-math"="true"
// FAST-SAME: "unsafe-fp-math"="true" // FAST-SAME: "unsafe-fp-math"="true"
// NOFAST-NOT: "unsafe-fp-math"="true" // NOFAST-NOT: "unsafe-fp-math"="true"

clang/test/CodeGenOpenCL/amdgpu-attrs.cl

Show First 20 Lines • Show All 156 Lines • ▼ Show 20 Lines
// NONAMDHSA-NOT: "amdgpu-implicitarg-num-bytes" // NONAMDHSA-NOT: "amdgpu-implicitarg-num-bytes"
// CHECK-NOT: "amdgpu-flat-work-group-size"="0,0" // CHECK-NOT: "amdgpu-flat-work-group-size"="0,0"
// CHECK-NOT: "amdgpu-waves-per-eu"="0" // CHECK-NOT: "amdgpu-waves-per-eu"="0"
// CHECK-NOT: "amdgpu-waves-per-eu"="0,0" // CHECK-NOT: "amdgpu-waves-per-eu"="0,0"
// CHECK-NOT: "amdgpu-num-sgpr"="0" // CHECK-NOT: "amdgpu-num-sgpr"="0"
// CHECK-NOT: "amdgpu-num-vgpr"="0" // CHECK-NOT: "amdgpu-num-vgpr"="0"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_64_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="64,64" "amdgpu-implicitarg-num-bytes"="56" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_64_64]] = {{.*}} "amdgpu-flat-work-group-size"="64,64" "amdgpu-implicitarg-num-bytes"="56"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_16_128]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="16,128" "amdgpu-implicitarg-num-bytes"="56" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_16_128]] = {{.*}} "amdgpu-flat-work-group-size"="16,128" "amdgpu-implicitarg-num-bytes"="56"
// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[WAVES_PER_EU_2]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[WAVES_PER_EU_2_4]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[WAVES_PER_EU_2_4]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" // CHECK-DAG: attributes [[NUM_SGPR_32]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32"
// CHECK-DAG: attributes [[NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" // CHECK-DAG: attributes [[NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64"
// CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_SGPR_32]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_SGPR_32]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[NUM_SGPR_32_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" // CHECK-DAG: attributes [[NUM_SGPR_32_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4" // CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64]] = {{.*}} "amdgpu-flat-work-group-size"="32,64" "amdgpu-implicitarg-num-bytes"="56" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"
// CHECK-DAG: attributes [[A_FUNCTION]] = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" // CHECK-DAG: attributes [[A_FUNCTION]] = {{.*}} "correctly-rounded-divide-sqrt-fp-math"="false"
// CHECK-DAG: attributes [[DEFAULT_KERNEL_ATTRS]] = { convergent noinline nounwind optnone "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56" // CHECK-DAG: attributes [[DEFAULT_KERNEL_ATTRS]] = {{.*}} "amdgpu-flat-work-group-size"="1,256" "amdgpu-implicitarg-num-bytes"="56"

clang/test/CodeGenOpenCL/norecurse.cl

  • This file was added.
// RUN: %clang_cc1 -O0 -emit-llvm -o - %s | FileCheck %s
kernel void kernel1(int a) {}
// CHECK: define{{.*}}@kernel1{{.*}}#[[ATTR:[0-9]*]]
// CHECK: attributes #[[ATTR]] = {{.*}}norecurse

clang/test/SemaCUDA/call-kernel-from-kernel.cu

  • This file was added.
// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
// RUN: -verify -fsyntax-only -verify-ignore-unexpected=note
#include "Inputs/cuda.h"
__global__ void kernel1();
__global__ void kernel2() {
kernel1<<<1,1>>>(); // expected-error {{reference to __global__ function 'kernel1' in __global__ function}}
}