Skip to content

Commit 2e6caeb

Browse files
skc7skc7
committed
[CIR][AMDGPU] Add lowering for amdgcn permlane builtins
1 parent e0e0345 commit 2e6caeb

File tree

6 files changed

+364
-5
lines changed

6 files changed

+364
-5
lines changed

‎clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp‎

Lines changed: 10 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -98,9 +98,17 @@ mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
9898
llvm_unreachable("mov_dpp_* NYI");
9999
}
100100
case AMDGPU::BI__builtin_amdgcn_permlane16:
101-
case AMDGPU::BI__builtin_amdgcn_permlanex16:
101+
case AMDGPU::BI__builtin_amdgcn_permlanex16: {
102+
llvm::StringRef intrinsicName =
103+
builtinId == AMDGPU::BI__builtin_amdgcn_permlane16
104+
? "amdgcn.permlane16"
105+
: "amdgcn.permlanex16";
106+
return emitBuiltinWithOneOverloadedType<6>(expr, intrinsicName)
107+
.getScalarVal();
108+
}
102109
case AMDGPU::BI__builtin_amdgcn_permlane64: {
103-
llvm_unreachable("permlane_* NYI");
110+
return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64")
111+
.getScalarVal();
104112
}
105113
case AMDGPU::BI__builtin_amdgcn_readlane:
106114
case AMDGPU::BI__builtin_amdgcn_readfirstlane: {

‎clang/test/CIR/CodeGen/HIP/builtins-amdgcn-gfx10.hip‎

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
#include "../Inputs/cuda.h"
2+
3+
// REQUIRES: amdgpu-registered-target
4+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
5+
// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir
6+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
7+
8+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
9+
// RUN: -target-cpu gfx1011 -fcuda-is-device -emit-cir %s -o %t.cir
10+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
11+
12+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
13+
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-cir %s -o %t.cir
14+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
15+
16+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
17+
// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-llvm %s -o %t.ll
18+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
19+
20+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
21+
// RUN: -target-cpu gfx1011 -fcuda-is-device -emit-llvm %s -o %t.ll
22+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
23+
24+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
25+
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
26+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
27+
28+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
29+
// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-llvm %s -o %t.ll
30+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
31+
32+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
33+
// RUN: -target-cpu gfx1011 -fcuda-is-device -emit-llvm %s -o %t.ll
34+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
35+
36+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
37+
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
38+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
39+
40+
//===----------------------------------------------------------------------===//
41+
// Test AMDGPU builtins
42+
//===----------------------------------------------------------------------===//
43+
44+
// CIR-LABEL: @_Z15test_permlane16Pjjjjj
45+
// CIR: cir.llvm.intrinsic "amdgcn.permlane16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i
46+
// LLVM: define{{.*}} void @_Z15test_permlane16Pjjjjj
47+
// LLVM: call i32 @llvm.amdgcn.permlane16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
48+
// OGCG: define{{.*}} void @_Z15test_permlane16Pjjjjj
49+
// OGCG: call i32 @llvm.amdgcn.permlane16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
50+
__device__ void test_permlane16(unsigned int* out, unsigned int a, unsigned int b, unsigned int c, unsigned int d) {
51+
*out = __builtin_amdgcn_permlane16(a, b, c, d, 0, 0);
52+
}
53+
54+
55+
// CIR-LABEL: @_Z16test_permlanex16Pjjjjj
56+
// CIR: cir.llvm.intrinsic "amdgcn.permlanex16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i
57+
// LLVM: define{{.*}} void @_Z16test_permlanex16Pjjjjj
58+
// LLVM: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
59+
// OGCG: define{{.*}} void @_Z16test_permlanex16Pjjjjj
60+
// OGCG: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
61+
__device__ void test_permlanex16(unsigned int* out, unsigned int a, unsigned int b, unsigned int c, unsigned int d) {
62+
*out = __builtin_amdgcn_permlanex16(a, b, c, d, 0, 0);
63+
}

‎clang/test/CIR/CodeGen/HIP/builtins-amdgcn-gfx11.hip‎

Lines changed: 108 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,108 @@
1+
#include "../Inputs/cuda.h"
2+
3+
// REQUIRES: amdgpu-registered-target
4+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
5+
// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-cir %s -o %t.cir
6+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
7+
8+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
9+
// RUN: -target-cpu gfx1101 -fcuda-is-device -emit-cir %s -o %t.cir
10+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
11+
12+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
13+
// RUN: -target-cpu gfx1102 -fcuda-is-device -emit-cir %s -o %t.cir
14+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
15+
16+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
17+
// RUN: -target-cpu gfx1103 -fcuda-is-device -emit-cir %s -o %t.cir
18+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
19+
20+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
21+
// RUN: -target-cpu gfx1150 -fcuda-is-device -emit-cir %s -o %t.cir
22+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
23+
24+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
25+
// RUN: -target-cpu gfx1151 -fcuda-is-device -emit-cir %s -o %t.cir
26+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
27+
28+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
29+
// RUN: -target-cpu gfx1152 -fcuda-is-device -emit-cir %s -o %t.cir
30+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
31+
32+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
33+
// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll
34+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
35+
36+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
37+
// RUN: -target-cpu gfx1101 -fcuda-is-device -emit-llvm %s -o %t.ll
38+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
39+
40+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
41+
// RUN: -target-cpu gfx1102 -fcuda-is-device -emit-llvm %s -o %t.ll
42+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
43+
44+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
45+
// RUN: -target-cpu gfx1103 -fcuda-is-device -emit-llvm %s -o %t.ll
46+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
47+
48+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
49+
// RUN: -target-cpu gfx1150 -fcuda-is-device -emit-llvm %s -o %t.ll
50+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
51+
52+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
53+
// RUN: -target-cpu gfx1151 -fcuda-is-device -emit-llvm %s -o %t.ll
54+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
55+
56+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
57+
// RUN: -target-cpu gfx1152 -fcuda-is-device -emit-llvm %s -o %t.ll
58+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
59+
60+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
61+
// RUN: -target-cpu gfx1153 -fcuda-is-device -emit-llvm %s -o %t.ll
62+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
63+
64+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
65+
// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll
66+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
67+
68+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
69+
// RUN: -target-cpu gfx1101 -fcuda-is-device -emit-llvm %s -o %t.ll
70+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
71+
72+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
73+
// RUN: -target-cpu gfx1102 -fcuda-is-device -emit-llvm %s -o %t.ll
74+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
75+
76+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
77+
// RUN: -target-cpu gfx1103 -fcuda-is-device -emit-llvm %s -o %t.ll
78+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
79+
80+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
81+
// RUN: -target-cpu gfx1150 -fcuda-is-device -emit-llvm %s -o %t.ll
82+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
83+
84+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
85+
// RUN: -target-cpu gfx1151 -fcuda-is-device -emit-llvm %s -o %t.ll
86+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
87+
88+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
89+
// RUN: -target-cpu gfx1152 -fcuda-is-device -emit-llvm %s -o %t.ll
90+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
91+
92+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
93+
// RUN: -target-cpu gfx1153 -fcuda-is-device -emit-llvm %s -o %t.ll
94+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
95+
96+
//===----------------------------------------------------------------------===//
97+
// Test AMDGPU builtins
98+
//===----------------------------------------------------------------------===//
99+
100+
// CIR-LABEL: @_Z15test_permlane64Pjj
101+
// CIR: cir.llvm.intrinsic "amdgcn.permlane64" {{.*}} : (!u32i) -> !u32i
102+
// LLVM: define{{.*}} void @_Z15test_permlane64Pjj
103+
// LLVM: call i32 @llvm.amdgcn.permlane64.i32(i32 %{{.*}})
104+
// OGCG: define{{.*}} void @_Z15test_permlane64Pjj
105+
// OGCG: call i32 @llvm.amdgcn.permlane64.i32(i32 %{{.*}})
106+
__device__ void test_permlane64(unsigned int* out, unsigned int a) {
107+
*out = __builtin_amdgcn_permlane64(a);
108+
}

‎clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,15 +2,15 @@
22

33
// REQUIRES: amdgpu-registered-target
44
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
5-
// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
5+
// RUN: -target-cpu tahiti -fcuda-is-device -emit-cir %s -o %t.cir
66
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
77

88
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
9-
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
9+
// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t.ll
1010
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
1111

1212
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
13-
// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
13+
// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t.ll
1414
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
1515

1616
//===----------------------------------------------------------------------===//

‎clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-gfx10.cl‎

Lines changed: 65 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,65 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
3+
// RUN: -target-cpu gfx1010 -emit-cir %s -o %t.cir
4+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
5+
6+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
7+
// RUN: -target-cpu gfx1011 -emit-cir %s -o %t.cir
8+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
9+
10+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
11+
// RUN: -target-cpu gfx1012 -emit-cir %s -o %t.cir
12+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
13+
14+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
15+
// RUN: -target-cpu gfx1010 -emit-llvm %s -o %t.ll
16+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
17+
18+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
19+
// RUN: -target-cpu gfx1011 -emit-llvm %s -o %t.ll
20+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
21+
22+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
23+
// RUN: -target-cpu gfx1012 -emit-llvm %s -o %t.ll
24+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
25+
26+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
27+
// RUN: -target-cpu gfx1010 -emit-llvm %s -o %t.ll
28+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
29+
30+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
31+
// RUN: -target-cpu gfx1011 -emit-llvm %s -o %t.ll
32+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
33+
34+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
35+
// RUN: -target-cpu gfx1012 -emit-llvm %s -o %t.ll
36+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
37+
38+
//===----------------------------------------------------------------------===//
39+
// Test AMDGPU builtins
40+
//===----------------------------------------------------------------------===//
41+
42+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
43+
44+
typedef unsigned int uint;
45+
typedef unsigned long ulong;
46+
47+
// CIR-LABEL: @test_permlane16
48+
// CIR: cir.llvm.intrinsic "amdgcn.permlane16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i
49+
// LLVM: define{{.*}} void @test_permlane16
50+
// LLVM: call i32 @llvm.amdgcn.permlane16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
51+
// OGCG: define{{.*}} void @test_permlane16
52+
// OGCG: call i32 @llvm.amdgcn.permlane16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
53+
void test_permlane16(global uint* out, uint a, uint b, uint c, uint d) {
54+
*out = __builtin_amdgcn_permlane16(a, b, c, d, 0, 0);
55+
}
56+
57+
// CIR-LABEL: @test_permlanex16
58+
// CIR: cir.llvm.intrinsic "amdgcn.permlanex16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i
59+
// LLVM: define{{.*}} void @test_permlanex16
60+
// LLVM: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
61+
// OGCG: define{{.*}} void @test_permlanex16
62+
// OGCG: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
63+
void test_permlanex16(global uint* out, uint a, uint b, uint c, uint d) {
64+
*out = __builtin_amdgcn_permlanex16(a, b, c, d, 0, 0);
65+
}

‎clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-gfx11.cl‎

Lines changed: 115 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,115 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
3+
// RUN: -target-cpu gfx1100 -emit-cir %s -o %t.cir
4+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
5+
6+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
7+
// RUN: -target-cpu gfx1101 -emit-cir %s -o %t.cir
8+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
9+
10+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
11+
// RUN: -target-cpu gfx1102 -emit-cir %s -o %t.cir
12+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
13+
14+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
15+
// RUN: -target-cpu gfx1103 -emit-cir %s -o %t.cir
16+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
17+
18+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
19+
// RUN: -target-cpu gfx1150 -emit-cir %s -o %t.cir
20+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
21+
22+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
23+
// RUN: -target-cpu gfx1151 -emit-cir %s -o %t.cir
24+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
25+
26+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
27+
// RUN: -target-cpu gfx1152 -emit-cir %s -o %t.cir
28+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
29+
30+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
31+
// RUN: -target-cpu gfx1153 -emit-cir %s -o %t.cir
32+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
33+
34+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
35+
// RUN: -target-cpu gfx1100 -emit-llvm %s -o %t.ll
36+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
37+
38+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
39+
// RUN: -target-cpu gfx1101 -emit-llvm %s -o %t.ll
40+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
41+
42+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
43+
// RUN: -target-cpu gfx1102 -emit-llvm %s -o %t.ll
44+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
45+
46+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
47+
// RUN: -target-cpu gfx1103 -emit-llvm %s -o %t.ll
48+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
49+
50+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
51+
// RUN: -target-cpu gfx1150 -emit-llvm %s -o %t.ll
52+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
53+
54+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
55+
// RUN: -target-cpu gfx1151 -emit-llvm %s -o %t.ll
56+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
57+
58+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
59+
// RUN: -target-cpu gfx1152 -emit-llvm %s -o %t.ll
60+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
61+
62+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
63+
// RUN: -target-cpu gfx1153 -emit-llvm %s -o %t.ll
64+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
65+
66+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
67+
// RUN: -target-cpu gfx1100 -emit-llvm %s -o %t.ll
68+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
69+
70+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
71+
// RUN: -target-cpu gfx1101 -emit-llvm %s -o %t.ll
72+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
73+
74+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
75+
// RUN: -target-cpu gfx1102 -emit-llvm %s -o %t.ll
76+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
77+
78+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
79+
// RUN: -target-cpu gfx1103 -emit-llvm %s -o %t.ll
80+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
81+
82+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
83+
// RUN: -target-cpu gfx1150 -emit-llvm %s -o %t.ll
84+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
85+
86+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
87+
// RUN: -target-cpu gfx1151 -emit-llvm %s -o %t.ll
88+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
89+
90+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
91+
// RUN: -target-cpu gfx1152 -emit-llvm %s -o %t.ll
92+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
93+
94+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
95+
// RUN: -target-cpu gfx1153 -emit-llvm %s -o %t.ll
96+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
97+
98+
//===----------------------------------------------------------------------===//
99+
// Test AMDGPU builtins
100+
//===----------------------------------------------------------------------===//
101+
102+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
103+
104+
typedef unsigned int uint;
105+
typedef unsigned long ulong;
106+
107+
// CIR-LABEL: @test_permlanex16
108+
// CIR: cir.llvm.intrinsic "amdgcn.permlanex16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i
109+
// LLVM: define{{.*}} void @test_permlanex16
110+
// LLVM: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
111+
// OGCG: define{{.*}} void @test_permlanex16
112+
// OGCG: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false)
113+
void test_permlanex16(global uint* out, uint a, uint b, uint c, uint d) {
114+
*out = __builtin_amdgcn_permlanex16(a, b, c, d, 0, 0);
115+
}

0 commit comments

Comments
 (0)