https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/220611
Adds codegen support for AMDGCN global and DS load/read builtins with transpose variants. These builtins are lowered to their corresponding `llvm.amdgcn.*` intrinsics. >From c3e6642cda72ca67215b3f2797f7ed8e0f0f97b3 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 2 Sep 2026 16:21:09 +0200 Subject: [PATCH] add global-load builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 74 +++++++--- .../builtins-amdgcn-gfx1250-load-tr.hip | 131 ++++++++++++++++++ .../builtins-amdgcn-gfx950-read-tr.hip | 78 +++++++++++ .../builtins-amdgcn-global-load-tr-w32.hip | 60 ++++++++ .../builtins-amdgcn-global-load-tr-w64.hip | 59 ++++++++ 5 files changed, 386 insertions(+), 16 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-load-tr.hip create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950-read-tr.hip create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w32.hip create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w64.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index e3a12d825434a..c8f072050d0b5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -428,33 +428,75 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32: case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16: case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16: - case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32: case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32: case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16: case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16: - case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32: case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32: case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32: case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16: case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16: case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + llvm::StringRef intrinsicName; + switch (builtinId) { + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32: + case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32: + intrinsicName = "amdgcn.global.load.tr.b64"; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16: + case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: + intrinsicName = "amdgcn.global.load.tr.b128"; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32: + intrinsicName = "amdgcn.global.load.tr4.b64"; + break; + case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32: + intrinsicName = "amdgcn.global.load.tr6.b96"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32: + intrinsicName = "amdgcn.ds.load.tr4.b64"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32: + intrinsicName = "amdgcn.ds.load.tr6.b96"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: + intrinsicName = "amdgcn.ds.load.tr8.b64"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16: + case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16: + case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: + intrinsicName = "amdgcn.ds.load.tr16.b128"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32: + intrinsicName = "amdgcn.ds.read.tr4.b64"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32: + intrinsicName = "amdgcn.ds.read.tr8.b64"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32: + intrinsicName = "amdgcn.ds.read.tr6.b96"; + break; + case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: + case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16: + case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16: + intrinsicName = "amdgcn.ds.read.tr16.b64"; + break; + } + mlir::Type loadTy = convertType(expr->getType()); + return emitBuiltinWithOneOverloadedType<1>(expr, intrinsicName, loadTy) + .getValue(); } case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-load-tr.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-load-tr.hip new file mode 100644 index 0000000000000..13942e99245dc --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-load-tr.hip @@ -0,0 +1,131 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR +// calls yet, so the LLVM check lines use call{{.*}} to tolerate the +// difference between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef int v3i __attribute__((ext_vector_type(3))); +typedef __fp16 v8h __attribute__((ext_vector_type(8))); +typedef short v8s __attribute__((ext_vector_type(8))); +typedef __bf16 v8y __attribute__((ext_vector_type(8))); + +typedef __attribute__((address_space(1))) v2i *v2i_as1_ptr; +typedef __attribute__((address_space(3))) v2i *v2i_as3_ptr; +typedef __attribute__((address_space(1))) v3i *v3i_as1_ptr; +typedef __attribute__((address_space(3))) v3i *v3i_as3_ptr; +typedef __attribute__((address_space(1))) v8h *v8h_as1_ptr; +typedef __attribute__((address_space(3))) v8h *v8h_as3_ptr; +typedef __attribute__((address_space(1))) v8s *v8s_as1_ptr; +typedef __attribute__((address_space(3))) v8s *v8s_as3_ptr; +typedef __attribute__((address_space(1))) v8y *v8y_as1_ptr; +typedef __attribute__((address_space(3))) v8y *v8y_as3_ptr; + +// CIR-LABEL: @_Z30test_global_load_tr4_b64_v2i32PU3AS1Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr4.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z30test_global_load_tr4_b64_v2i32PU3AS1Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.global.load.tr4.b64.v2i32(ptr addrspace(1) %{{.*}}) +__device__ v2i test_global_load_tr4_b64_v2i32(v2i_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr4_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z30test_global_load_tr8_b64_v2i32PU3AS1Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z30test_global_load_tr8_b64_v2i32PU3AS1Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.global.load.tr.b64.v2i32(ptr addrspace(1) %{{.*}}) +__device__ v2i test_global_load_tr8_b64_v2i32(v2i_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr8_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z30test_global_load_tr6_b96_v3i32PU3AS1Dv3_i +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr6.b96" {{.*}} : (!cir.ptr<!cir.vector<3 x !s32i>, target_address_space(1)>) -> !cir.vector<3 x !s32i> +// LLVM-LABEL: @_Z30test_global_load_tr6_b96_v3i32PU3AS1Dv3_i +// LLVM: call{{.*}} <3 x i32> @llvm.amdgcn.global.load.tr6.b96.v3i32(ptr addrspace(1) %{{.*}}) +__device__ v3i test_global_load_tr6_b96_v3i32(v3i_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr6_b96_v3i32(inptr); +} + +// CIR-LABEL: @_Z32test_global_load_tr16_b128_v8i16PU3AS1Dv8_s +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !s16i>, target_address_space(1)>) -> !cir.vector<8 x !s16i> +// LLVM-LABEL: @_Z32test_global_load_tr16_b128_v8i16PU3AS1Dv8_s +// LLVM: call{{.*}} <8 x i16> @llvm.amdgcn.global.load.tr.b128.v8i16(ptr addrspace(1) %{{.*}}) +__device__ v8s test_global_load_tr16_b128_v8i16(v8s_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr16_b128_v8i16(inptr); +} + +// CIR-LABEL: @_Z32test_global_load_tr16_b128_v8f16PU3AS1Dv8_Dh +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !cir.f16>, target_address_space(1)>) -> !cir.vector<8 x !cir.f16> +// LLVM-LABEL: @_Z32test_global_load_tr16_b128_v8f16PU3AS1Dv8_Dh +// LLVM: call{{.*}} <8 x half> @llvm.amdgcn.global.load.tr.b128.v8f16(ptr addrspace(1) %{{.*}}) +__device__ v8h test_global_load_tr16_b128_v8f16(v8h_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr16_b128_v8f16(inptr); +} + +// CIR-LABEL: @_Z33test_global_load_tr16_b128_v8bf16PU3AS1Dv8_DF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !cir.bf16>, target_address_space(1)>) -> !cir.vector<8 x !cir.bf16> +// LLVM-LABEL: @_Z33test_global_load_tr16_b128_v8bf16PU3AS1Dv8_DF16b +// LLVM: call{{.*}} <8 x bfloat> @llvm.amdgcn.global.load.tr.b128.v8bf16(ptr addrspace(1) %{{.*}}) +__device__ v8y test_global_load_tr16_b128_v8bf16(v8y_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr16_b128_v8bf16(inptr); +} + +// CIR-LABEL: @_Z26test_ds_load_tr4_b64_v2i32PU3AS3Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.load.tr4.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(3)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z26test_ds_load_tr4_b64_v2i32PU3AS3Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.ds.load.tr4.b64.v2i32(ptr addrspace(3) %{{.*}}) +__device__ v2i test_ds_load_tr4_b64_v2i32(v2i_as3_ptr inptr) { + return __builtin_amdgcn_ds_load_tr4_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z26test_ds_load_tr8_b64_v2i32PU3AS3Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.load.tr8.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(3)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z26test_ds_load_tr8_b64_v2i32PU3AS3Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.ds.load.tr8.b64.v2i32(ptr addrspace(3) %{{.*}}) +__device__ v2i test_ds_load_tr8_b64_v2i32(v2i_as3_ptr inptr) { + return __builtin_amdgcn_ds_load_tr8_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z26test_ds_load_tr6_b96_v3i32PU3AS3Dv3_i +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.load.tr6.b96" {{.*}} : (!cir.ptr<!cir.vector<3 x !s32i>, target_address_space(3)>) -> !cir.vector<3 x !s32i> +// LLVM-LABEL: @_Z26test_ds_load_tr6_b96_v3i32PU3AS3Dv3_i +// LLVM: call{{.*}} <3 x i32> @llvm.amdgcn.ds.load.tr6.b96.v3i32(ptr addrspace(3) %{{.*}}) +__device__ v3i test_ds_load_tr6_b96_v3i32(v3i_as3_ptr inptr) { + return __builtin_amdgcn_ds_load_tr6_b96_v3i32(inptr); +} + +// CIR-LABEL: @_Z28test_ds_load_tr16_b128_v8i16PU3AS3Dv8_s +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.load.tr16.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !s16i>, target_address_space(3)>) -> !cir.vector<8 x !s16i> +// LLVM-LABEL: @_Z28test_ds_load_tr16_b128_v8i16PU3AS3Dv8_s +// LLVM: call{{.*}} <8 x i16> @llvm.amdgcn.ds.load.tr16.b128.v8i16(ptr addrspace(3) %{{.*}}) +__device__ v8s test_ds_load_tr16_b128_v8i16(v8s_as3_ptr inptr) { + return __builtin_amdgcn_ds_load_tr16_b128_v8i16(inptr); +} + +// CIR-LABEL: @_Z28test_ds_load_tr16_b128_v8f16PU3AS3Dv8_Dh +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.load.tr16.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !cir.f16>, target_address_space(3)>) -> !cir.vector<8 x !cir.f16> +// LLVM-LABEL: @_Z28test_ds_load_tr16_b128_v8f16PU3AS3Dv8_Dh +// LLVM: call{{.*}} <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) %{{.*}}) +__device__ v8h test_ds_load_tr16_b128_v8f16(v8h_as3_ptr inptr) { + return __builtin_amdgcn_ds_load_tr16_b128_v8f16(inptr); +} + +// CIR-LABEL: @_Z29test_ds_load_tr16_b128_v8bf16PU3AS3Dv8_DF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.load.tr16.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !cir.bf16>, target_address_space(3)>) -> !cir.vector<8 x !cir.bf16> +// LLVM-LABEL: @_Z29test_ds_load_tr16_b128_v8bf16PU3AS3Dv8_DF16b +// LLVM: call{{.*}} <8 x bfloat> @llvm.amdgcn.ds.load.tr16.b128.v8bf16(ptr addrspace(3) %{{.*}}) +__device__ v8y test_ds_load_tr16_b128_v8bf16(v8y_as3_ptr inptr) { + return __builtin_amdgcn_ds_load_tr16_b128_v8bf16(inptr); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950-read-tr.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950-read-tr.hip new file mode 100644 index 0000000000000..e976d66809f70 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950-read-tr.hip @@ -0,0 +1,78 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu9.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu9.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu9.50-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR +// calls yet, so the LLVM check lines use call{{.*}} to tolerate the +// difference between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef int v3i __attribute__((ext_vector_type(3))); +typedef __fp16 v4h __attribute__((ext_vector_type(4))); +typedef short v4s __attribute__((ext_vector_type(4))); +typedef __bf16 v4y __attribute__((ext_vector_type(4))); + +typedef __attribute__((address_space(3))) v2i *v2i_as3_ptr; +typedef __attribute__((address_space(3))) v3i *v3i_as3_ptr; +typedef __attribute__((address_space(3))) v4h *v4h_as3_ptr; +typedef __attribute__((address_space(3))) v4s *v4s_as3_ptr; +typedef __attribute__((address_space(3))) v4y *v4y_as3_ptr; + +// CIR-LABEL: @_Z26test_ds_read_tr4_b64_v2i32PU3AS3Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.read.tr4.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(3)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z26test_ds_read_tr4_b64_v2i32PU3AS3Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.ds.read.tr4.b64.v2i32(ptr addrspace(3) %{{.*}}) +__device__ v2i test_ds_read_tr4_b64_v2i32(v2i_as3_ptr inptr) { + return __builtin_amdgcn_ds_read_tr4_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z26test_ds_read_tr6_b96_v3i32PU3AS3Dv3_i +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.read.tr6.b96" {{.*}} : (!cir.ptr<!cir.vector<3 x !s32i>, target_address_space(3)>) -> !cir.vector<3 x !s32i> +// LLVM-LABEL: @_Z26test_ds_read_tr6_b96_v3i32PU3AS3Dv3_i +// LLVM: call{{.*}} <3 x i32> @llvm.amdgcn.ds.read.tr6.b96.v3i32(ptr addrspace(3) %{{.*}}) +__device__ v3i test_ds_read_tr6_b96_v3i32(v3i_as3_ptr inptr) { + return __builtin_amdgcn_ds_read_tr6_b96_v3i32(inptr); +} + +// CIR-LABEL: @_Z26test_ds_read_tr8_b64_v2i32PU3AS3Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.read.tr8.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(3)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z26test_ds_read_tr8_b64_v2i32PU3AS3Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.ds.read.tr8.b64.v2i32(ptr addrspace(3) %{{.*}}) +__device__ v2i test_ds_read_tr8_b64_v2i32(v2i_as3_ptr inptr) { + return __builtin_amdgcn_ds_read_tr8_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z27test_ds_read_tr16_b64_v4i16PU3AS3Dv4_s +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.read.tr16.b64" {{.*}} : (!cir.ptr<!cir.vector<4 x !s16i>, target_address_space(3)>) -> !cir.vector<4 x !s16i> +// LLVM-LABEL: @_Z27test_ds_read_tr16_b64_v4i16PU3AS3Dv4_s +// LLVM: call{{.*}} <4 x i16> @llvm.amdgcn.ds.read.tr16.b64.v4i16(ptr addrspace(3) %{{.*}}) +__device__ v4s test_ds_read_tr16_b64_v4i16(v4s_as3_ptr inptr) { + return __builtin_amdgcn_ds_read_tr16_b64_v4i16(inptr); +} + +// CIR-LABEL: @_Z27test_ds_read_tr16_b64_v4f16PU3AS3Dv4_Dh +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.read.tr16.b64" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.f16>, target_address_space(3)>) -> !cir.vector<4 x !cir.f16> +// LLVM-LABEL: @_Z27test_ds_read_tr16_b64_v4f16PU3AS3Dv4_Dh +// LLVM: call{{.*}} <4 x half> @llvm.amdgcn.ds.read.tr16.b64.v4f16(ptr addrspace(3) %{{.*}}) +__device__ v4h test_ds_read_tr16_b64_v4f16(v4h_as3_ptr inptr) { + return __builtin_amdgcn_ds_read_tr16_b64_v4f16(inptr); +} + +// CIR-LABEL: @_Z28test_ds_read_tr16_b64_v4bf16PU3AS3Dv4_DF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.read.tr16.b64" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.bf16>, target_address_space(3)>) -> !cir.vector<4 x !cir.bf16> +// LLVM-LABEL: @_Z28test_ds_read_tr16_b64_v4bf16PU3AS3Dv4_DF16b +// LLVM: call{{.*}} <4 x bfloat> @llvm.amdgcn.ds.read.tr16.b64.v4bf16(ptr addrspace(3) %{{.*}}) +__device__ v4y test_ds_read_tr16_b64_v4bf16(v4y_as3_ptr inptr) { + return __builtin_amdgcn_ds_read_tr16_b64_v4bf16(inptr); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w32.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w32.hip new file mode 100644 index 0000000000000..52be00c8ac0b6 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w32.hip @@ -0,0 +1,60 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -target-feature +wavefrontsize32 -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -target-feature +wavefrontsize32 -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -target-feature +wavefrontsize32 -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR +// calls yet, so the LLVM check lines use call{{.*}} to tolerate the +// difference between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef __fp16 v8h __attribute__((ext_vector_type(8))); +typedef short v8s __attribute__((ext_vector_type(8))); +typedef __bf16 v8y __attribute__((ext_vector_type(8))); + +typedef __attribute__((address_space(1))) v2i *v2i_as1_ptr; +typedef __attribute__((address_space(1))) v8h *v8h_as1_ptr; +typedef __attribute__((address_space(1))) v8s *v8s_as1_ptr; +typedef __attribute__((address_space(1))) v8y *v8y_as1_ptr; + +// CIR-LABEL: @_Z29test_global_load_tr_b64_v2i32PU3AS1Dv2_i +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)>) -> !cir.vector<2 x !s32i> +// LLVM-LABEL: @_Z29test_global_load_tr_b64_v2i32PU3AS1Dv2_i +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.global.load.tr.b64.v2i32(ptr addrspace(1) %{{.*}}) +__device__ v2i test_global_load_tr_b64_v2i32(v2i_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b64_v2i32(inptr); +} + +// CIR-LABEL: @_Z30test_global_load_tr_b128_v8i16PU3AS1Dv8_s +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !s16i>, target_address_space(1)>) -> !cir.vector<8 x !s16i> +// LLVM-LABEL: @_Z30test_global_load_tr_b128_v8i16PU3AS1Dv8_s +// LLVM: call{{.*}} <8 x i16> @llvm.amdgcn.global.load.tr.b128.v8i16(ptr addrspace(1) %{{.*}}) +__device__ v8s test_global_load_tr_b128_v8i16(v8s_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b128_v8i16(inptr); +} + +// CIR-LABEL: @_Z30test_global_load_tr_b128_v8f16PU3AS1Dv8_Dh +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !cir.f16>, target_address_space(1)>) -> !cir.vector<8 x !cir.f16> +// LLVM-LABEL: @_Z30test_global_load_tr_b128_v8f16PU3AS1Dv8_Dh +// LLVM: call{{.*}} <8 x half> @llvm.amdgcn.global.load.tr.b128.v8f16(ptr addrspace(1) %{{.*}}) +__device__ v8h test_global_load_tr_b128_v8f16(v8h_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b128_v8f16(inptr); +} + +// CIR-LABEL: @_Z31test_global_load_tr_b128_v8bf16PU3AS1Dv8_DF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<8 x !cir.bf16>, target_address_space(1)>) -> !cir.vector<8 x !cir.bf16> +// LLVM-LABEL: @_Z31test_global_load_tr_b128_v8bf16PU3AS1Dv8_DF16b +// LLVM: call{{.*}} <8 x bfloat> @llvm.amdgcn.global.load.tr.b128.v8bf16(ptr addrspace(1) %{{.*}}) +__device__ v8y test_global_load_tr_b128_v8bf16(v8y_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b128_v8bf16(inptr); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w64.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w64.hip new file mode 100644 index 0000000000000..b822d1b549b19 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w64.hip @@ -0,0 +1,59 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -target-feature +wavefrontsize64 -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -target-feature +wavefrontsize64 -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -target-feature +wavefrontsize64 -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR +// calls yet, so the LLVM check lines use call{{.*}} to tolerate the +// difference between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) + +typedef __fp16 v4h __attribute__((ext_vector_type(4))); +typedef short v4s __attribute__((ext_vector_type(4))); +typedef __bf16 v4y __attribute__((ext_vector_type(4))); + +typedef __attribute__((address_space(1))) int *int_as1_ptr; +typedef __attribute__((address_space(1))) v4h *v4h_as1_ptr; +typedef __attribute__((address_space(1))) v4s *v4s_as1_ptr; +typedef __attribute__((address_space(1))) v4y *v4y_as1_ptr; + +// CIR-LABEL: @_Z27test_global_load_tr_b64_i32PU3AS1i +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b64" {{.*}} : (!cir.ptr<!s32i, target_address_space(1)>) -> !s32i +// LLVM-LABEL: @_Z27test_global_load_tr_b64_i32PU3AS1i +// LLVM: call{{.*}} i32 @llvm.amdgcn.global.load.tr.b64.i32(ptr addrspace(1) %{{.*}}) +__device__ int test_global_load_tr_b64_i32(int_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b64_i32(inptr); +} + +// CIR-LABEL: @_Z30test_global_load_tr_b128_v4i16PU3AS1Dv4_s +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<4 x !s16i>, target_address_space(1)>) -> !cir.vector<4 x !s16i> +// LLVM-LABEL: @_Z30test_global_load_tr_b128_v4i16PU3AS1Dv4_s +// LLVM: call{{.*}} <4 x i16> @llvm.amdgcn.global.load.tr.b128.v4i16(ptr addrspace(1) %{{.*}}) +__device__ v4s test_global_load_tr_b128_v4i16(v4s_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b128_v4i16(inptr); +} + +// CIR-LABEL: @_Z30test_global_load_tr_b128_v4f16PU3AS1Dv4_Dh +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.f16>, target_address_space(1)>) -> !cir.vector<4 x !cir.f16> +// LLVM-LABEL: @_Z30test_global_load_tr_b128_v4f16PU3AS1Dv4_Dh +// LLVM: call{{.*}} <4 x half> @llvm.amdgcn.global.load.tr.b128.v4f16(ptr addrspace(1) %{{.*}}) +__device__ v4h test_global_load_tr_b128_v4f16(v4h_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b128_v4f16(inptr); +} + +// CIR-LABEL: @_Z31test_global_load_tr_b128_v4bf16PU3AS1Dv4_DF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.global.load.tr.b128" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.bf16>, target_address_space(1)>) -> !cir.vector<4 x !cir.bf16> +// LLVM-LABEL: @_Z31test_global_load_tr_b128_v4bf16PU3AS1Dv4_DF16b +// LLVM: call{{.*}} <4 x bfloat> @llvm.amdgcn.global.load.tr.b128.v4bf16(ptr addrspace(1) %{{.*}}) +__device__ v4y test_global_load_tr_b128_v4bf16(v4y_as1_ptr inptr) { + return __builtin_amdgcn_global_load_tr_b128_v4bf16(inptr); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
