Author: Ayokunle Amodu Date: 2026-09-10T21:49:56-04:00 New Revision: 7c3849ae3bc162f383c54e38b1ed32b733a31bd8
URL: https://github.com/llvm/llvm-project/commit/7c3849ae3bc162f383c54e38b1ed32b733a31bd8 DIFF: https://github.com/llvm/llvm-project/commit/7c3849ae3bc162f383c54e38b1ed32b733a31bd8.diff LOG: [CIR][AMDGPU] Add support for AMDGCN global/DS load builtins (#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. Added: clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-load-tr.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950-read-tr.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w32.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-global-load-tr-w64.hip Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 3494d1363d70f..b7efaed0e0871 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -442,45 +442,66 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned 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: + return emitBuiltinWithOneOverloadedType<1>( + expr, "amdgcn.global.load.tr.b64", convertType(expr->getType())) + .getValue(); 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_tr4_b64_v2i32: - case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32: - 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: + return emitBuiltinWithOneOverloadedType<1>( + expr, "amdgcn.global.load.tr.b128", convertType(expr->getType())) + .getValue(); + case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32: + return emitBuiltinWithOneOverloadedType<1>( + expr, "amdgcn.global.load.tr4.b64", convertType(expr->getType())) + .getValue(); + case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32: + return emitBuiltinWithOneOverloadedType<1>( + expr, "amdgcn.global.load.tr6.b96", convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32: - case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr4.b64", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr6.b96", + convertType(expr->getType())) + .getValue(); + case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr8.b64", + convertType(expr->getType())) + .getValue(); 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: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr16.b128", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr4.b64", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr8.b64", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr6.b96", + convertType(expr->getType())) + .getValue(); + 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: - 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{}; - } + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr16.b64", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: 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..449f80762444a --- /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 -target-feature +wavefrontsize32 -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 -target-feature +wavefrontsize32 -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 -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 +// diff erence 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 +// diff erence 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 +// diff erence 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 +// diff erence 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
