[clang] [Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic (PR #95276)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 12 10:51:18 PDT 2024
================
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p5(ptr addrspace(5) [[P:%.*]], i16 4, i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+ __buffer_rsrc_t test_amdgcn_make_buffer_rsrc(void *p, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, 4, num, flags);
----------------
arsenm wrote:
Test more combinations of arguments. All arguments can be variable, but you have a constant here
Also test with different address space pointers. We also probably should have a HIP test. Ideally we would error for private/local inputs.
Also test null and literal pointer inputs?
https://github.com/llvm/llvm-project/pull/95276
More information about the cfe-commits
mailing list