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