Skip to content

Commit 78caf4f

Browse files
ranapratap55ranapratap55
committed
[AMDGPU] Adds EmitAMDGCNImageOverloadedReturnType for amdgcn_image_load/store and adds 'image-insts' feature
1 parent 21c6829 commit 78caf4f

File tree

3 files changed

+19
-78
lines changed

3 files changed

+19
-78
lines changed

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,9 +10,12 @@
1010
//
1111
//===----------------------------------------------------------------------===//
1212

13+
#include "CodeGenFunction.h"
1314
#include "CGBuiltin.h"
1415
#include "clang/Basic/TargetBuiltins.h"
16+
#include "clang/Frontend/FrontendDiagnostic.h"
1517
#include "llvm/Analysis/ValueTracking.h"
18+
#include "llvm/CodeGen/MachineFunction.h"
1619
#include "llvm/IR/IntrinsicsAMDGPU.h"
1720
#include "llvm/IR/IntrinsicsR600.h"
1821
#include "llvm/IR/MemoryModelRelaxationAnnotations.h"

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 1 addition & 78 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@
1313
#include "clang/Sema/SemaAMDGPU.h"
1414
#include "clang/Basic/DiagnosticSema.h"
1515
#include "clang/Basic/TargetBuiltins.h"
16+
#include "clang/Frontend/FrontendDiagnostic.h"
1617
#include "clang/Sema/Ownership.h"
1718
#include "clang/Sema/Sema.h"
1819
#include "llvm/Support/AMDGPUAddrSpace.h"
@@ -211,84 +212,6 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
211212
(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
212213
(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)));
213214
}
214-
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
215-
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
216-
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
217-
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
218-
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
219-
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
220-
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
221-
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
222-
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
223-
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
224-
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
225-
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
226-
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
227-
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
228-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
229-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
230-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
231-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
232-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
233-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
234-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
235-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
236-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
237-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
238-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
239-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
240-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
241-
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
242-
unsigned ArgCount = TheCall->getNumArgs() - 1;
243-
llvm::APSInt Result;
244-
bool isImmArg =
245-
(!(SemaRef.BuiltinConstantArg(TheCall, 0, Result)) &&
246-
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
247-
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
248-
? false
249-
: true;
250-
251-
return isImmArg;
252-
}
253-
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
254-
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
255-
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
256-
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
257-
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
258-
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
259-
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
260-
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
261-
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
262-
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
263-
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
264-
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
265-
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
266-
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
267-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
268-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
269-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
270-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
271-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
272-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
273-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
274-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
275-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
276-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
277-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
278-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
279-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
280-
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
281-
unsigned ArgCount = TheCall->getNumArgs() - 1;
282-
llvm::APSInt Result;
283-
bool isImmArg =
284-
(!(SemaRef.BuiltinConstantArg(TheCall, 1, Result)) &&
285-
!(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) &&
286-
!(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)))
287-
? false
288-
: true;
289-
290-
return isImmArg;
291-
}
292215
default:
293216
return false;
294217
}

‎clang/test/CodeGen/builtins-image-load.c‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,7 @@ float4 test_builtin_image_load_2d_1(float4 v4f32, int i32, int8 vec8i32) {
5656

5757
return __builtin_amdgcn_image_load_2d_v4f32_i32(100, i32, i32, vec8i32, 120, 110);
5858
}
59+
5960
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_2d_2(
6061
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
6162
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -106,6 +107,7 @@ float test_builtin_image_load_2darray(float f32, int i32, int8 vec8i32) {
106107

107108
return __builtin_amdgcn_image_load_2darray_f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
108109
}
110+
109111
// CHECK-LABEL: define dso_local <4 x float> @test_builtin_image_load_2darray_1(
110112
// CHECK-SAME: <4 x float> noundef [[V4F32:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
111113
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -131,6 +133,7 @@ float4 test_builtin_image_load_2darray_1(float4 v4f32, int i32, int8 vec8i32) {
131133

132134
return __builtin_amdgcn_image_load_2darray_v4f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
133135
}
136+
134137
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_2darray_2(
135138
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
136139
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -180,6 +183,7 @@ float4 test_builtin_image_load_1d_1(float4 v4f32, int i32, int8 vec8i32) {
180183

181184
return __builtin_amdgcn_image_load_1d_v4f32_i32(100, i32, vec8i32, 120, 110);
182185
}
186+
183187
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_1d_2(
184188
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
185189
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -228,6 +232,7 @@ float4 test_builtin_image_load_1darray_1(float4 v4f32, int i32, int8 vec8i32) {
228232

229233
return __builtin_amdgcn_image_load_1darray_v4f32_i32(100, i32, i32, vec8i32, 120, 110);
230234
}
235+
231236
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_1darray_2(
232237
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
233238
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -278,6 +283,7 @@ float4 test_builtin_image_load_3d_1(float4 v4f32, int i32, int8 vec8i32) {
278283

279284
return __builtin_amdgcn_image_load_3d_v4f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
280285
}
286+
281287
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_3d_2(
282288
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
283289
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -329,6 +335,7 @@ float4 test_builtin_image_load_cube_1(float4 v4f32, int i32, int8 vec8i32) {
329335

330336
return __builtin_amdgcn_image_load_cube_v4f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
331337
}
338+
332339
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_cube_2(
333340
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
334341
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -379,6 +386,7 @@ float4 test_builtin_image_load_mip_1d_1(float4 v4f32, int i32, int8 vec8i32) {
379386

380387
return __builtin_amdgcn_image_load_mip_1d_v4f32_i32(100, i32, i32, vec8i32, 120, 110);
381388
}
389+
382390
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_mip_1d_2(
383391
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
384392
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -429,6 +437,7 @@ float4 test_builtin_image_load_mip_1darray_1(float4 v4f32, int i32, int8 vec8i32
429437

430438
return __builtin_amdgcn_image_load_mip_1darray_v4f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
431439
}
440+
432441
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_mip_1darray_2(
433442
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
434443
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -480,6 +489,7 @@ float test_builtin_image_load_mip_2d(float f32, int i32, int8 vec8i32) {
480489

481490
return __builtin_amdgcn_image_load_mip_2d_f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
482491
}
492+
483493
// CHECK-LABEL: define dso_local <4 x float> @test_builtin_image_load_mip_2d_1(
484494
// CHECK-SAME: <4 x float> noundef [[V4F32:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
485495
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -505,6 +515,7 @@ float4 test_builtin_image_load_mip_2d_1(float4 v4f32, int i32, int8 vec8i32) {
505515

506516
return __builtin_amdgcn_image_load_mip_2d_v4f32_i32(100, i32, i32, i32, vec8i32, 120, 110);
507517
}
518+
508519
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_mip_2d_2(
509520
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
510521
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -557,6 +568,7 @@ float test_builtin_image_load_mip_2darray(float f32, int i32, int8 vec8i32) {
557568

558569
return __builtin_amdgcn_image_load_mip_2darray_f32_i32(100, i32, i32, i32, i32, vec8i32, 120, 110);
559570
}
571+
560572
// CHECK-LABEL: define dso_local <4 x float> @test_builtin_image_load_mip_2darray_1(
561573
// CHECK-SAME: <4 x float> noundef [[V4F32:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
562574
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -583,6 +595,7 @@ float4 test_builtin_image_load_mip_2darray_1(float4 v4f32, int i32, int8 vec8i32
583595

584596
return __builtin_amdgcn_image_load_mip_2darray_v4f32_i32(100, i32, i32, i32, i32, vec8i32, 120, 110);
585597
}
598+
586599
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_mip_2darray_2(
587600
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
588601
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -636,6 +649,7 @@ float4 test_builtin_image_load_mip_3d_1(float4 v4f32, int i32, int8 vec8i32) {
636649

637650
return __builtin_amdgcn_image_load_mip_3d_v4f32_i32(100, i32, i32, i32, i32, vec8i32, 120, 110);
638651
}
652+
639653
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_mip_3d_2(
640654
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
641655
// CHECK-NEXT: [[ENTRY:.*:]]
@@ -689,6 +703,7 @@ float4 test_builtin_image_load_mip_cube_1(float4 v4f32, int i32, int8 vec8i32) {
689703

690704
return __builtin_amdgcn_image_load_mip_cube_v4f32_i32(100, i32, i32, i32, i32, vec8i32, 120, 110);
691705
}
706+
692707
// CHECK-LABEL: define dso_local <4 x half> @test_builtin_image_load_mip_cube_2(
693708
// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], i32 noundef [[I32:%.*]], <8 x i32> noundef [[VEC8I32:%.*]]) #[[ATTR0]] {
694709
// CHECK-NEXT: [[ENTRY:.*:]]

0 commit comments

Comments
 (0)