102 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// REQUIRES: amdgpu-registered-target3// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -target-cpu verde -emit-llvm -o - %s | FileCheck %s4 5// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0(6// CHECK-NEXT: entry:7// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i648// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[CONV]], i32 [[FLAGS:%.*]])9// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]10//11__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p, short stride, int num, int flags) {12 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);13}14 15// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_stride_constant(16// CHECK-NEXT: entry:17// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6418// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0(ptr [[P:%.*]], i16 4, i64 [[CONV]], i32 [[FLAGS:%.*]])19// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]20//21__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_stride_constant(void *p, int num, int flags) {22 return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);23}24 25// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_num_constant(26// CHECK-NEXT: entry:27// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 1234, i32 [[FLAGS:%.*]])28// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]29//30__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p, short stride, int flags) {31 return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);32}33 34// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_flags_constant(35// CHECK-NEXT: entry:36// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6437// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[CONV]], i32 5678)38// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]39//40__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p, short stride, int num) {41 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);42}43 44// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1(45// CHECK-NEXT: entry:46// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6447// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[CONV]], i32 [[FLAGS:%.*]])48// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]49//50__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1(global void *p, short stride, int num, int flags) {51 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);52}53 54// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_stride_constant(55// CHECK-NEXT: entry:56// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6457// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) [[P:%.*]], i16 4, i64 [[CONV]], i32 [[FLAGS:%.*]])58// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]59//60__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_stride_constant(global void *p, int num, int flags) {61 return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);62}63 64// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_num_constant(65// CHECK-NEXT: entry:66// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) [[P:%.*]], i16 [[STRIDE:%.*]], i64 1234, i32 [[FLAGS:%.*]])67// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]68//69__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_num_constant(global void *p, short stride, int flags) {70 return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);71}72 73// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_flags_constant(74// CHECK-NEXT: entry:75// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6476// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[CONV]], i32 5678)77// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]78//79__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_flags_constant(global void *p, short stride, int num) {80 return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);81}82 83// CHECK-LABEL: @test_amdgcn_make_buffer_p0_nullptr(84// CHECK-NEXT: entry:85// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6486// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0(ptr null, i16 [[STRIDE:%.*]], i64 [[CONV]], i32 [[FLAGS:%.*]])87// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]88//89__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p0_nullptr(short stride, int num, int flags) {90 return __builtin_amdgcn_make_buffer_rsrc((void *)0LL, stride, num, flags);91}92 93// CHECK-LABEL: @test_amdgcn_make_buffer_p1_nullptr(94// CHECK-NEXT: entry:95// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[NUM:%.*]] to i6496// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) null, i16 [[STRIDE:%.*]], i64 [[CONV]], i32 [[FLAGS:%.*]])97// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]98//99__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p1_nullptr(short stride, int num, int flags) {100 return __builtin_amdgcn_make_buffer_rsrc((global void *)0LL, stride, num, flags);101}102