Skip to content

Commit dd6a9f6

Browse files
lucas-ramilucas-rami
committed
Remove verifier code and add codegen tests
1 parent d08d49d commit dd6a9f6

File tree

5 files changed

+55
-64
lines changed

5 files changed

+55
-64
lines changed

‎clang/lib/CodeGen/Targets/AMDGPU.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -751,7 +751,7 @@ void CodeGenModule::handleAMDGPUWavesPerEUAttr(
751751
"Min must be non-zero when Max is non-zero");
752752
assert(Min <= Max && "Min must be less than or equal to Max");
753753
// Do not add the attribute if min,max=0,0.
754-
if (Min != 0) {
754+
if (Max != 0) {
755755
std::string AttrVal = llvm::utostr(Min) + "," + llvm::utostr(Max);
756756
F->addFnAttr("amdgpu-waves-per-eu", AttrVal);
757757
}

‎clang/test/CodeGenHIP/amdgpu-waves-per-eu.hip

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,53 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-attributes --check-globals
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fno-ident -fcuda-is-device -emit-llvm -o - %s | FileCheck %s
4+
5+
// COM: Most tests are in the OpenCL semastics, this is just a verification for HIP
6+
7+
#define __global__ __attribute__((global))
8+
9+
//.
10+
// CHECK: @__hip_cuid_ = addrspace(1) global i8 0
11+
// CHECK: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
12+
//.
13+
// CHECK: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
14+
// CHECK-LABEL: define {{[^@]+}}@_Z21kernel_waves_per_eu_0v
15+
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
16+
// CHECK-NEXT: entry:
17+
// CHECK-NEXT: ret void
18+
//
19+
__global__ __attribute__((amdgpu_waves_per_eu(0))) void kernel_waves_per_eu_0() {}
20+
21+
// Equivalent to kernel_waves_per_eu_0.
22+
// CHECK: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
23+
// CHECK-LABEL: define {{[^@]+}}@_Z23kernel_waves_per_eu_0_0v
24+
// CHECK-SAME: () #[[ATTR0]] {
25+
// CHECK-NEXT: entry:
26+
// CHECK-NEXT: ret void
27+
//
28+
__global__ __attribute__((amdgpu_waves_per_eu(0, 0))) void kernel_waves_per_eu_0_0() {}
29+
30+
// CHECK: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
31+
// CHECK-LABEL: define {{[^@]+}}@_Z23kernel_waves_per_eu_0_4v
32+
// CHECK-SAME: () #[[ATTR1:[0-9]+]] {
33+
// CHECK-NEXT: entry:
34+
// CHECK-NEXT: ret void
35+
//
36+
__global__ __attribute__((amdgpu_waves_per_eu(0, 4))) void kernel_waves_per_eu_0_4() {}
37+
38+
// CHECK: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
39+
// CHECK-LABEL: define {{[^@]+}}@_Z23kernel_waves_per_eu_1_4v
40+
// CHECK-SAME: () #[[ATTR2:[0-9]+]] {
41+
// CHECK-NEXT: entry:
42+
// CHECK-NEXT: ret void
43+
//
44+
__global__ __attribute__((amdgpu_waves_per_eu(1, 4))) void kernel_waves_per_eu_1_4() {}
45+
//.
46+
// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
47+
// CHECK: attributes #[[ATTR1]] = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "amdgpu-waves-per-eu"="0,4" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
48+
// CHECK: attributes #[[ATTR2]] = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "amdgpu-waves-per-eu"="1,4" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
49+
//.
50+
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
51+
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
52+
// CHECK: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
53+
//.

‎clang/test/SemaOpenCL/amdgpu-attrs.cl

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,7 @@ __attribute__((amdgpu_num_vgpr(0))) kernel void kernel_num_vgpr_0() {}
6060

6161
kernel __attribute__((amdgpu_flat_work_group_size(32, 64))) void kernel_flat_work_group_size_32_64() {}
6262
kernel __attribute__((amdgpu_waves_per_eu(2))) void kernel_waves_per_eu_2() {}
63+
kernel __attribute__((amdgpu_waves_per_eu(0, 4))) kernel void kernel_waves_per_eu_0_4() {}
6364
kernel __attribute__((amdgpu_waves_per_eu(2, 4))) void kernel_waves_per_eu_2_4() {}
6465
kernel __attribute__((amdgpu_num_sgpr(32))) void kernel_num_sgpr_32() {}
6566
kernel __attribute__((amdgpu_num_vgpr(64))) void kernel_num_vgpr_64() {}

‎llvm/lib/IR/Verifier.cpp

Lines changed: 0 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -2517,29 +2517,6 @@ void Verifier::verifyFunctionAttrs(FunctionType *FT, AttributeList Attrs,
25172517
CheckFailed("invalid value for 'denormal-fp-math-f32' attribute: " + S,
25182518
V);
25192519
}
2520-
2521-
if (TT.isAMDGPU()) {
2522-
if (auto A = Attrs.getFnAttr("amdgpu-waves-per-eu"); A.isValid()) {
2523-
std::pair<StringRef, StringRef> Strs = A.getValueAsString().split(',');
2524-
unsigned Min = 0;
2525-
StringRef MinStr = Strs.first.trim();
2526-
Check(!MinStr.getAsInteger(0, Min),
2527-
"minimum for 'amdgpu-waves-per-eu' must be integer: " + MinStr);
2528-
if (!Strs.second.empty()) {
2529-
unsigned Max = 0;
2530-
StringRef MaxStr = Strs.second.trim();
2531-
Check(!MaxStr.getAsInteger(0, Max),
2532-
"maximum for 'amdgpu-waves-per-eu' must be integer: " + MaxStr);
2533-
Check(Max, "maximum for 'amdgpu-waves-per-eu' must be non-zero");
2534-
Check(Min <= Max, "minimum must be less than or equal to maximum for "
2535-
"'amdgpu-waves-per-eu': " +
2536-
MinStr + " > " + MaxStr);
2537-
} else {
2538-
Check(Min, "minimum for 'amdgpu-waves-per-eu' must be non-zero when "
2539-
"maximum is not provided");
2540-
}
2541-
}
2542-
}
25432520
}
25442521

25452522
void Verifier::verifyFunctionMetadata(

‎llvm/test/Verifier/AMDGPU/amdgpu-waves-per-eu.ll

Lines changed: 0 additions & 40 deletions
This file was deleted.

0 commit comments

Comments
 (0)