Skip to content

Commit 5ee5dd8

Browse files
skc7skc7
committed
[CIR][AMDGPU] Add lowering for amdgcn div fixup builtins
1 parent 4916f0e commit 5ee5dd8

File tree

5 files changed

+178
-1
lines changed

5 files changed

+178
-1
lines changed

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -159,7 +159,8 @@ mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
159159
case AMDGPU::BI__builtin_amdgcn_div_fixup:
160160
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
161161
case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
162-
llvm_unreachable("div_fixup_* NYI");
162+
return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.div.fixup")
163+
.getScalarVal();
163164
}
164165
case AMDGPU::BI__builtin_amdgcn_trig_preop:
165166
case AMDGPU::BI__builtin_amdgcn_trig_preopf: {

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

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
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 tonga -fcuda-is-device -emit-cir %s -o %t.cir
6+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
7+
8+
// REQUIRES: amdgpu-registered-target
9+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
10+
// RUN: -target-cpu gfx900 -fcuda-is-device -emit-cir %s -o %t.cir
11+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
12+
13+
// REQUIRES: amdgpu-registered-target
14+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
15+
// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir
16+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
17+
18+
// REQUIRES: amdgpu-registered-target
19+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
20+
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-cir %s -o %t.cir
21+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
22+
23+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
24+
// RUN: -target-cpu tonga -fcuda-is-device -emit-llvm %s -o %t.ll
25+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
26+
27+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
28+
// RUN: -target-cpu gfx900 -fcuda-is-device -emit-llvm %s -o %t.ll
29+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
30+
31+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
32+
// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-llvm %s -o %t.ll
33+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
34+
35+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
36+
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
37+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
38+
39+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
40+
// RUN: -target-cpu tonga -fcuda-is-device -emit-llvm %s -o %t.ll
41+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
42+
43+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
44+
// RUN: -target-cpu gfx900 -fcuda-is-device -emit-llvm %s -o %t.ll
45+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
46+
47+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
48+
// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-llvm %s -o %t.ll
49+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
50+
51+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
52+
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
53+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
54+
55+
//===----------------------------------------------------------------------===//
56+
// Test AMDGPU builtins
57+
//===----------------------------------------------------------------------===//
58+
59+
// CIR-LABEL: @_Z18test_div_fixup_f16PDF16_DF16_DF16_DF16_
60+
// CIR: cir.llvm.intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
61+
// LLVM: define{{.*}} void @_Z18test_div_fixup_f16PDF16_DF16_DF16_DF16_
62+
// LLVM: call{{.*}} half @llvm.amdgcn.div.fixup.f16(half %{{.+}}, half %{{.+}}, half %{{.+}})
63+
// OGCG: define{{.*}} void @_Z18test_div_fixup_f16PDF16_DF16_DF16_DF16_
64+
// OGCG: call{{.*}} half @llvm.amdgcn.div.fixup.f16(half %{{.+}}, half %{{.+}}, half %{{.+}})
65+
__device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) {
66+
*out = __builtin_amdgcn_div_fixuph(a, b, c);
67+
}

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

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -305,3 +305,23 @@ __device__ void test_readlane(int* out, int a, int b) {
305305
__device__ void test_readfirstlane(int* out, int a) {
306306
*out = __builtin_amdgcn_readfirstlane(a);
307307
}
308+
309+
// CIR-LABEL: @_Z18test_div_fixup_f32Pffff
310+
// CIR: cir.llvm.intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.float, !cir.float, !cir.float) -> !cir.float
311+
// LLVM: define{{.*}} void @_Z18test_div_fixup_f32Pffff
312+
// LLVM: call{{.*}} float @llvm.amdgcn.div.fixup.f32(float %{{.+}}, float %{{.+}}, float %{{.+}})
313+
// OGCG: define{{.*}} void @_Z18test_div_fixup_f32Pffff
314+
// OGCG: call{{.*}} float @llvm.amdgcn.div.fixup.f32(float %{{.+}}, float %{{.+}}, float %{{.+}})
315+
__device__ void test_div_fixup_f32(float* out, float a, float b, float c) {
316+
*out = __builtin_amdgcn_div_fixupf(a, b, c);
317+
}
318+
319+
// CIR-LABEL: @_Z18test_div_fixup_f64Pdddd
320+
// CIR: cir.llvm.intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.double, !cir.double, !cir.double) -> !cir.double
321+
// LLVM: define{{.*}} void @_Z18test_div_fixup_f64Pdddd
322+
// LLVM: call{{.*}} double @llvm.amdgcn.div.fixup.f64(double %{{.+}}, double %{{.+}}, double %{{.+}})
323+
// OGCG: define{{.*}} void @_Z18test_div_fixup_f64Pdddd
324+
// OGCG: call{{.*}} double @llvm.amdgcn.div.fixup.f64(double %{{.+}}, double %{{.+}}, double %{{.+}})
325+
__device__ void test_div_fixup_f64(double* out, double a, double b, double c) {
326+
*out = __builtin_amdgcn_div_fixup(a, b, c);
327+
}

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

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
1+
#include "../Inputs/cuda.h"
2+
3+
// REQUIRES: amdgpu-registered-target
4+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
5+
// RUN: -target-cpu tonga -emit-cir %s -o %t.cir
6+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
7+
8+
// REQUIRES: amdgpu-registered-target
9+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
10+
// RUN: -target-cpu gfx900 -emit-cir %s -o %t.cir
11+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
12+
13+
// REQUIRES: amdgpu-registered-target
14+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
15+
// RUN: -target-cpu gfx1010 -emit-cir %s -o %t.cir
16+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
17+
18+
// REQUIRES: amdgpu-registered-target
19+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
20+
// RUN: -target-cpu gfx1012 -emit-cir %s -o %t.cir
21+
// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
22+
23+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
24+
// RUN: -target-cpu tonga -emit-llvm %s -o %t.ll
25+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
26+
27+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
28+
// RUN: -target-cpu gfx900 -emit-llvm %s -o %t.ll
29+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
30+
31+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
32+
// RUN: -target-cpu gfx1010 -emit-llvm %s -o %t.ll
33+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
34+
35+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -fclangir \
36+
// RUN: -target-cpu gfx1012 -emit-llvm %s -o %t.ll
37+
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
38+
39+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
40+
// RUN: -target-cpu tonga -emit-llvm %s -o %t.ll
41+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
42+
43+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
44+
// RUN: -target-cpu gfx900 -emit-llvm %s -o %t.ll
45+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
46+
47+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
48+
// RUN: -target-cpu gfx1010 -emit-llvm %s -o %t.ll
49+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
50+
51+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 \
52+
// RUN: -target-cpu gfx1012 -emit-llvm %s -o %t.ll
53+
// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s
54+
55+
//===----------------------------------------------------------------------===//
56+
// Test AMDGPU builtins
57+
//===----------------------------------------------------------------------===//
58+
59+
// CIR-LABEL: @test_div_fixup_f16
60+
// CIR: cir.llvm.intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
61+
// LLVM: define{{.*}} void @test_div_fixup_f16
62+
// LLVM: call{{.*}} half @llvm.amdgcn.div.fixup.f16(half %{{.+}}, half %{{.+}}, half %{{.+}})
63+
// OGCG: define{{.*}} void @test_div_fixup_f16
64+
// OGCG: call{{.*}} half @llvm.amdgcn.div.fixup.f16(half %{{.+}}, half %{{.+}}, half %{{.+}})
65+
__device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) {
66+
*out = __builtin_amdgcn_div_fixuph(a, b, c);
67+
}

‎clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -318,3 +318,25 @@ void test_readlane(global int* out, int a, int b) {
318318
void test_readfirstlane(global int* out, int a) {
319319
*out = __builtin_amdgcn_readfirstlane(a);
320320
}
321+
322+
// CIR-LABEL: @test_div_fixup_f32
323+
// CIR: cir.llvm.intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.float, !cir.float, !cir.float) -> !cir.float
324+
// LLVM: define{{.*}} void @test_div_fixup_f32
325+
// LLVM: call{{.*}} float @llvm.amdgcn.div.fixup.f32(float %{{.+}}, float %{{.+}}, float %{{.+}})
326+
// OGCG: define{{.*}} void @test_div_fixup_f32
327+
// OGCG: call{{.*}} float @llvm.amdgcn.div.fixup.f32(float %{{.+}}, float %{{.+}}, float %{{.+}})
328+
void test_div_fixup_f32(global float* out, float a, float b, float c)
329+
{
330+
*out = __builtin_amdgcn_div_fixupf(a, b, c);
331+
}
332+
333+
// CIR-LABEL: @test_div_fixup_f64
334+
// CIR: cir.llvm.intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.double, !cir.double, !cir.double) -> !cir.double
335+
// LLVM: define{{.*}} void @test_div_fixup_f64
336+
// LLVM: call{{.*}} double @llvm.amdgcn.div.fixup.f64(double %{{.+}}, double %{{.+}}, double %{{.+}})
337+
// OGCG: define{{.*}} void @test_div_fixup_f64
338+
// OGCG: call{{.*}} double @llvm.amdgcn.div.fixup.f64(double %{{.+}}, double %{{.+}}, double %{{.+}})
339+
void test_div_fixup_f64(global double* out, double a, double b, double c)
340+
{
341+
*out = __builtin_amdgcn_div_fixup(a, b, c);
342+
}

0 commit comments

Comments
 (0)