Author: Ayokunle Amodu Date: 2026-09-04T07:26:51-04:00 New Revision: f26bfb40c9dad93578167859c6722305bef4cc52
URL: https://github.com/llvm/llvm-project/commit/f26bfb40c9dad93578167859c6722305bef4cc52 DIFF: https://github.com/llvm/llvm-project/commit/f26bfb40c9dad93578167859c6722305bef4cc52.diff LOG: [CIR][CUDA] Add support for __nvvm_ldg builtins (#213178) Adds CIRGen for the NVVM ldg builtins, which perform loads through the read-only data cache. These lower to the corresponding `llvm.nvvm.ldg.global.*` intrinsic. Added: clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu-ldg.cu Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp Removed: clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index 2220639876695..ae994005c588a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -16,6 +16,7 @@ #include "clang/Basic/TargetBuiltins.h" #include "clang/CIR/Dialect/IR/CIRDataLayout.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" +#include "llvm/Support/NVPTXAddrSpace.h" using namespace clang; using namespace clang::CIRGen; @@ -36,6 +37,25 @@ static mlir::Value makeLdu(CIRGenFunction &cgf, const CallExpr *expr, .getResult(); } +static mlir::Value makeLdg(CIRGenFunction &cgf, const CallExpr *expr) { + auto &builder = cgf.getBuilder(); + Address ptr = cgf.emitPointerWithAlignment(expr->getArg(0)); + QualType argType = expr->getArg(0)->getType(); + mlir::Type elemTy = cgf.convertTypeForMem(argType->getPointeeType()); + mlir::Location loc = cgf.getLoc(expr->getExprLoc()); + + // Use addrspace(1) for NVPTX ADDRESS_SPACE_GLOBAL. + mlir::Type globalPtrTy = cir::PointerType::get( + elemTy, cir::TargetAddressSpaceAttr::get( + builder.getContext(), llvm::NVPTXAS::ADDRESS_SPACE_GLOBAL)); + mlir::Value asc = + builder.createAddrSpaceCast(loc, ptr.getPointer(), globalPtrTy); + cir::LoadOp load = + builder.createAlignedLoad(loc, elemTy, asc, ptr.getAlignment()); + load.setInvariant(true); + return load.getResult(); +} + /// Emit a CIR LLVMIntrinsicCallOp for a unary NVVM intrinsic. /// The result type is inferred from the single argument. static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf, @@ -206,10 +226,7 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { case NVPTX::BI__nvvm_ldg_f4: case NVPTX::BI__nvvm_ldg_d: case NVPTX::BI__nvvm_ldg_d2: - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented NVPTX builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return makeLdg(*this, expr); case NVPTX::BI__nvvm_ldu_c: case NVPTX::BI__nvvm_ldu_sc: case NVPTX::BI__nvvm_ldu_c2: diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu-ldg.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu-ldg.cu new file mode 100644 index 0000000000000..3b767f24d576e --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu-ldg.cu @@ -0,0 +1,351 @@ +#include "Inputs/cuda.h" + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// 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 floating-point LLVM check lines use {{.*}} to tolerate the +// diff erence between CIR (no flags) and classic codegen ('contract'). + +typedef char char2 __attribute__((ext_vector_type(2))); +typedef unsigned char uchar2 __attribute__((ext_vector_type(2))); +typedef signed char schar2 __attribute__((ext_vector_type(2))); +typedef char char4 __attribute__((ext_vector_type(4))); +typedef unsigned char uchar4 __attribute__((ext_vector_type(4))); +typedef signed char schar4 __attribute__((ext_vector_type(4))); +typedef short short2 __attribute__((ext_vector_type(2))); +typedef unsigned short ushort2 __attribute__((ext_vector_type(2))); +typedef short short4 __attribute__((ext_vector_type(4))); +typedef unsigned short ushort4 __attribute__((ext_vector_type(4))); +typedef int int2 __attribute__((ext_vector_type(2))); +typedef unsigned int uint2 __attribute__((ext_vector_type(2))); +typedef int int4 __attribute__((ext_vector_type(4))); +typedef unsigned int uint4 __attribute__((ext_vector_type(4))); +typedef long long2 __attribute__((ext_vector_type(2))); +typedef unsigned long ulong2 __attribute__((ext_vector_type(2))); +typedef long long longlong2 __attribute__((ext_vector_type(2))); +typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2))); +typedef float float2 __attribute__((ext_vector_type(2))); +typedef float float4 __attribute__((ext_vector_type(4))); +typedef double double2 __attribute__((ext_vector_type(2))); + +// CIR-LABEL: @_Z8nvvm_ldgPKv +// LLVM-LABEL: @_Z8nvvm_ldgPKv +__device__ void nvvm_ldg(const void *p) { + // CIR: %[[C_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!s8i> -> !cir.ptr<!s8i, target_address_space(1)> + // CIR: cir.load invariant align(1) %[[C_CAST]] : !cir.ptr<!s8i, target_address_space(1)>, !s8i + // LLVM: %[[C_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i8, ptr addrspace(1) %[[C_CAST]], align 1, !invariant.load + __nvvm_ldg_c((const char *)p); + + // CIR: %[[UC_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!u8i> -> !cir.ptr<!u8i, target_address_space(1)> + // CIR: cir.load invariant align(1) %[[UC_CAST]] : !cir.ptr<!u8i, target_address_space(1)>, !u8i + // LLVM: %[[UC_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i8, ptr addrspace(1) %[[UC_CAST]], align 1, !invariant.load + __nvvm_ldg_uc((const unsigned char *)p); + + // CIR: %[[SC_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!s8i> -> !cir.ptr<!s8i, target_address_space(1)> + // CIR: cir.load invariant align(1) %[[SC_CAST]] : !cir.ptr<!s8i, target_address_space(1)>, !s8i + // LLVM: %[[SC_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i8, ptr addrspace(1) %[[SC_CAST]], align 1, !invariant.load + __nvvm_ldg_sc((const signed char *)p); + + // CIR: %[[S_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!s16i> -> !cir.ptr<!s16i, target_address_space(1)> + // CIR: cir.load invariant align(2) %[[S_CAST]] : !cir.ptr<!s16i, target_address_space(1)>, !s16i + // LLVM: %[[S_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i16, ptr addrspace(1) %[[S_CAST]], align 2, !invariant.load + __nvvm_ldg_s((const short *)p); + + // CIR: %[[US_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!u16i> -> !cir.ptr<!u16i, target_address_space(1)> + // CIR: cir.load invariant align(2) %[[US_CAST]] : !cir.ptr<!u16i, target_address_space(1)>, !u16i + // LLVM: %[[US_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i16, ptr addrspace(1) %[[US_CAST]], align 2, !invariant.load + __nvvm_ldg_us((const unsigned short *)p); + + // CIR: %[[I_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!s32i> -> !cir.ptr<!s32i, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[I_CAST]] : !cir.ptr<!s32i, target_address_space(1)>, !s32i + // LLVM: %[[I_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i32, ptr addrspace(1) %[[I_CAST]], align 4, !invariant.load + __nvvm_ldg_i((const int *)p); + + // CIR: %[[UI_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!u32i> -> !cir.ptr<!u32i, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[UI_CAST]] : !cir.ptr<!u32i, target_address_space(1)>, !u32i + // LLVM: %[[UI_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i32, ptr addrspace(1) %[[UI_CAST]], align 4, !invariant.load + __nvvm_ldg_ui((const unsigned int *)p); + + // CIR: %[[L_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!s64i> -> !cir.ptr<!s64i, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[L_CAST]] : !cir.ptr<!s64i, target_address_space(1)>, !s64i + // LLVM: %[[L_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i64, ptr addrspace(1) %[[L_CAST]], align 8, !invariant.load + __nvvm_ldg_l((const long *)p); + + // CIR: %[[UL_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!u64i> -> !cir.ptr<!u64i, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[UL_CAST]] : !cir.ptr<!u64i, target_address_space(1)>, !u64i + // LLVM: %[[UL_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load i64, ptr addrspace(1) %[[UL_CAST]], align 8, !invariant.load + __nvvm_ldg_ul((const unsigned long *)p); + + // CIR: %[[F_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.float> -> !cir.ptr<!cir.float, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[F_CAST]] : !cir.ptr<!cir.float, target_address_space(1)>, !cir.float + // LLVM: %[[F_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load float, ptr addrspace(1) %[[F_CAST]], align 4, !invariant.load + __nvvm_ldg_f((const float *)p); + + // CIR: %[[D_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.double> -> !cir.ptr<!cir.double, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[D_CAST]] : !cir.ptr<!cir.double, target_address_space(1)>, !cir.double + // LLVM: %[[D_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load double, ptr addrspace(1) %[[D_CAST]], align 8, !invariant.load + __nvvm_ldg_d((const double *)p); + + // CIR: %[[C2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !s8i>> -> !cir.ptr<!cir.vector<2 x !s8i>, target_address_space(1)> + // CIR: cir.load invariant align(2) %[[C2_CAST]] : !cir.ptr<!cir.vector<2 x !s8i>, target_address_space(1)>, !cir.vector<2 x !s8i> + // LLVM: %[[C2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i8>, ptr addrspace(1) %[[C2_CAST]], align 2, !invariant.load + __nvvm_ldg_c2((const char2 *)p); + + // CIR: %[[UC2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !u8i>> -> !cir.ptr<!cir.vector<2 x !u8i>, target_address_space(1)> + // CIR: cir.load invariant align(2) %[[UC2_CAST]] : !cir.ptr<!cir.vector<2 x !u8i>, target_address_space(1)>, !cir.vector<2 x !u8i> + // LLVM: %[[UC2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i8>, ptr addrspace(1) %[[UC2_CAST]], align 2, !invariant.load + __nvvm_ldg_uc2((const uchar2 *)p); + + // CIR: %[[SC2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !s8i>> -> !cir.ptr<!cir.vector<2 x !s8i>, target_address_space(1)> + // CIR: cir.load invariant align(2) %[[SC2_CAST]] : !cir.ptr<!cir.vector<2 x !s8i>, target_address_space(1)>, !cir.vector<2 x !s8i> + // LLVM: %[[SC2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i8>, ptr addrspace(1) %[[SC2_CAST]], align 2, !invariant.load + __nvvm_ldg_sc2((const schar2 *)p); + + // CIR: %[[C4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !s8i>> -> !cir.ptr<!cir.vector<4 x !s8i>, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[C4_CAST]] : !cir.ptr<!cir.vector<4 x !s8i>, target_address_space(1)>, !cir.vector<4 x !s8i> + // LLVM: %[[C4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i8>, ptr addrspace(1) %[[C4_CAST]], align 4, !invariant.load + __nvvm_ldg_c4((const char4 *)p); + + // CIR: %[[UC4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !u8i>> -> !cir.ptr<!cir.vector<4 x !u8i>, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[UC4_CAST]] : !cir.ptr<!cir.vector<4 x !u8i>, target_address_space(1)>, !cir.vector<4 x !u8i> + // LLVM: %[[UC4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i8>, ptr addrspace(1) %[[UC4_CAST]], align 4, !invariant.load + __nvvm_ldg_uc4((const uchar4 *)p); + + // CIR: %[[SC4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !s8i>> -> !cir.ptr<!cir.vector<4 x !s8i>, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[SC4_CAST]] : !cir.ptr<!cir.vector<4 x !s8i>, target_address_space(1)>, !cir.vector<4 x !s8i> + // LLVM: %[[SC4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i8>, ptr addrspace(1) %[[SC4_CAST]], align 4, !invariant.load + __nvvm_ldg_sc4((const schar4 *)p); + + // CIR: %[[S2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !s16i>> -> !cir.ptr<!cir.vector<2 x !s16i>, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[S2_CAST]] : !cir.ptr<!cir.vector<2 x !s16i>, target_address_space(1)>, !cir.vector<2 x !s16i> + // LLVM: %[[S2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i16>, ptr addrspace(1) %[[S2_CAST]], align 4, !invariant.load + __nvvm_ldg_s2((const short2 *)p); + + // CIR: %[[US2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !u16i>> -> !cir.ptr<!cir.vector<2 x !u16i>, target_address_space(1)> + // CIR: cir.load invariant align(4) %[[US2_CAST]] : !cir.ptr<!cir.vector<2 x !u16i>, target_address_space(1)>, !cir.vector<2 x !u16i> + // LLVM: %[[US2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i16>, ptr addrspace(1) %[[US2_CAST]], align 4, !invariant.load + __nvvm_ldg_us2((const ushort2 *)p); + + // CIR: %[[S4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !s16i>> -> !cir.ptr<!cir.vector<4 x !s16i>, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[S4_CAST]] : !cir.ptr<!cir.vector<4 x !s16i>, target_address_space(1)>, !cir.vector<4 x !s16i> + // LLVM: %[[S4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i16>, ptr addrspace(1) %[[S4_CAST]], align 8, !invariant.load + __nvvm_ldg_s4((const short4 *)p); + + // CIR: %[[US4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !u16i>> -> !cir.ptr<!cir.vector<4 x !u16i>, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[US4_CAST]] : !cir.ptr<!cir.vector<4 x !u16i>, target_address_space(1)>, !cir.vector<4 x !u16i> + // LLVM: %[[US4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i16>, ptr addrspace(1) %[[US4_CAST]], align 8, !invariant.load + __nvvm_ldg_us4((const ushort4 *)p); + + // CIR: %[[I2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !s32i>> -> !cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[I2_CAST]] : !cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)>, !cir.vector<2 x !s32i> + // LLVM: %[[I2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i32>, ptr addrspace(1) %[[I2_CAST]], align 8, !invariant.load + __nvvm_ldg_i2((const int2 *)p); + + // CIR: %[[UI2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !u32i>> -> !cir.ptr<!cir.vector<2 x !u32i>, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[UI2_CAST]] : !cir.ptr<!cir.vector<2 x !u32i>, target_address_space(1)>, !cir.vector<2 x !u32i> + // LLVM: %[[UI2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i32>, ptr addrspace(1) %[[UI2_CAST]], align 8, !invariant.load + __nvvm_ldg_ui2((const uint2 *)p); + + // CIR: %[[I4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !s32i>> -> !cir.ptr<!cir.vector<4 x !s32i>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[I4_CAST]] : !cir.ptr<!cir.vector<4 x !s32i>, target_address_space(1)>, !cir.vector<4 x !s32i> + // LLVM: %[[I4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i32>, ptr addrspace(1) %[[I4_CAST]], align 16, !invariant.load + __nvvm_ldg_i4((const int4 *)p); + + // CIR: %[[UI4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !u32i>> -> !cir.ptr<!cir.vector<4 x !u32i>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[UI4_CAST]] : !cir.ptr<!cir.vector<4 x !u32i>, target_address_space(1)>, !cir.vector<4 x !u32i> + // LLVM: %[[UI4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x i32>, ptr addrspace(1) %[[UI4_CAST]], align 16, !invariant.load + __nvvm_ldg_ui4((const uint4 *)p); + + // CIR: %[[L2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !s64i>> -> !cir.ptr<!cir.vector<2 x !s64i>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[L2_CAST]] : !cir.ptr<!cir.vector<2 x !s64i>, target_address_space(1)>, !cir.vector<2 x !s64i> + // LLVM: %[[L2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i64>, ptr addrspace(1) %[[L2_CAST]], align 16, !invariant.load + __nvvm_ldg_l2((const long2 *)p); + + // CIR: %[[UL2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !u64i>> -> !cir.ptr<!cir.vector<2 x !u64i>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[UL2_CAST]] : !cir.ptr<!cir.vector<2 x !u64i>, target_address_space(1)>, !cir.vector<2 x !u64i> + // LLVM: %[[UL2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i64>, ptr addrspace(1) %[[UL2_CAST]], align 16, !invariant.load + __nvvm_ldg_ul2((const ulong2 *)p); + + // CIR: %[[LL2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !s64i>> -> !cir.ptr<!cir.vector<2 x !s64i>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[LL2_CAST]] : !cir.ptr<!cir.vector<2 x !s64i>, target_address_space(1)>, !cir.vector<2 x !s64i> + // LLVM: %[[LL2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i64>, ptr addrspace(1) %[[LL2_CAST]], align 16, !invariant.load + __nvvm_ldg_ll2((const longlong2 *)p); + + // CIR: %[[ULL2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !u64i>> -> !cir.ptr<!cir.vector<2 x !u64i>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[ULL2_CAST]] : !cir.ptr<!cir.vector<2 x !u64i>, target_address_space(1)>, !cir.vector<2 x !u64i> + // LLVM: %[[ULL2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x i64>, ptr addrspace(1) %[[ULL2_CAST]], align 16, !invariant.load + __nvvm_ldg_ull2((const ulonglong2 *)p); + + // CIR: %[[F2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !cir.float>> -> !cir.ptr<!cir.vector<2 x !cir.float>, target_address_space(1)> + // CIR: cir.load invariant align(8) %[[F2_CAST]] : !cir.ptr<!cir.vector<2 x !cir.float>, target_address_space(1)>, !cir.vector<2 x !cir.float> + // LLVM: %[[F2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x float>, ptr addrspace(1) %[[F2_CAST]], align 8, !invariant.load + __nvvm_ldg_f2((const float2 *)p); + + // CIR: %[[F4_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<4 x !cir.float>> -> !cir.ptr<!cir.vector<4 x !cir.float>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[F4_CAST]] : !cir.ptr<!cir.vector<4 x !cir.float>, target_address_space(1)>, !cir.vector<4 x !cir.float> + // LLVM: %[[F4_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <4 x float>, ptr addrspace(1) %[[F4_CAST]], align 16, !invariant.load + __nvvm_ldg_f4((const float4 *)p); + + // CIR: %[[D2_CAST:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.vector<2 x !cir.double>> -> !cir.ptr<!cir.vector<2 x !cir.double>, target_address_space(1)> + // CIR: cir.load invariant align(16) %[[D2_CAST]] : !cir.ptr<!cir.vector<2 x !cir.double>, target_address_space(1)>, !cir.vector<2 x !cir.double> + // LLVM: %[[D2_CAST:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(1) + // LLVM: load <2 x double>, ptr addrspace(1) %[[D2_CAST]], align 16, !invariant.load + __nvvm_ldg_d2((const double2 *)p); +} + +// CIR-LABEL: @_Z8nvvm_lduPKv +// LLVM-LABEL: @_Z8nvvm_lduPKv +__device__ void nvvm_ldu(const void *p) { + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s8i>, !s32i) -> !s8i + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u8i>, !s32i) -> !u8i + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s8i>, !s32i) -> !s8i + // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1) + // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1) + // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1) + __nvvm_ldu_c((const char *)p); + __nvvm_ldu_uc((const unsigned char *)p); + __nvvm_ldu_sc((const signed char *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s16i>, !s32i) -> !s16i + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u16i>, !s32i) -> !u16i + // LLVM: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2) + // LLVM: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2) + __nvvm_ldu_s((const short *)p); + __nvvm_ldu_us((const unsigned short *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u32i>, !s32i) -> !u32i + // LLVM: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4) + // LLVM: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4) + __nvvm_ldu_i((const int *)p); + __nvvm_ldu_ui((const unsigned int *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s64i>, !s32i) -> !s64i + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u64i>, !s32i) -> !u64i + // LLVM: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8) + // LLVM: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8) + __nvvm_ldu_l((const long *)p); + __nvvm_ldu_ul((const unsigned long *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.float>, !s32i) -> !cir.float + // LLVM: call {{.*}}float @llvm.nvvm.ldu.global.f.f32.p0(ptr {{%[0-9]+}}, i32 4) + __nvvm_ldu_f((const float *)p); + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.double>, !s32i) -> !cir.double + // LLVM: call {{.*}}double @llvm.nvvm.ldu.global.f.f64.p0(ptr {{%[0-9]+}}, i32 8) + __nvvm_ldu_d((const double *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s8i>>, !s32i) -> !cir.vector<2 x !s8i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u8i>>, !s32i) -> !cir.vector<2 x !u8i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s8i>>, !s32i) -> !cir.vector<2 x !s8i> + // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2) + // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2) + // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2) + __nvvm_ldu_c2((const char2 *)p); + __nvvm_ldu_uc2((const uchar2 *)p); + __nvvm_ldu_sc2((const schar2 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s8i>>, !s32i) -> !cir.vector<4 x !s8i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u8i>>, !s32i) -> !cir.vector<4 x !u8i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s8i>>, !s32i) -> !cir.vector<4 x !s8i> + // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4) + // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4) + // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4) + __nvvm_ldu_c4((const char4 *)p); + __nvvm_ldu_uc4((const uchar4 *)p); + __nvvm_ldu_sc4((const schar4 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s16i>>, !s32i) -> !cir.vector<2 x !s16i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u16i>>, !s32i) -> !cir.vector<2 x !u16i> + // LLVM: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4) + // LLVM: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4) + __nvvm_ldu_s2((const short2 *)p); + __nvvm_ldu_us2((const ushort2 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s16i>>, !s32i) -> !cir.vector<4 x !s16i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u16i>>, !s32i) -> !cir.vector<4 x !u16i> + // LLVM: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8) + // LLVM: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8) + __nvvm_ldu_s4((const short4 *)p); + __nvvm_ldu_us4((const ushort4 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>>, !s32i) -> !cir.vector<2 x !s32i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u32i>>, !s32i) -> !cir.vector<2 x !u32i> + // LLVM: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8) + // LLVM: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8) + __nvvm_ldu_i2((const int2 *)p); + __nvvm_ldu_ui2((const uint2 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s32i>>, !s32i) -> !cir.vector<4 x !s32i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u32i>>, !s32i) -> !cir.vector<4 x !u32i> + // LLVM: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16) + // LLVM: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16) + __nvvm_ldu_i4((const int4 *)p); + __nvvm_ldu_ui4((const uint4 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s64i>>, !s32i) -> !cir.vector<2 x !s64i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u64i>>, !s32i) -> !cir.vector<2 x !u64i> + // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) + // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) + __nvvm_ldu_l2((const long2 *)p); + __nvvm_ldu_ul2((const ulong2 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s64i>>, !s32i) -> !cir.vector<2 x !s64i> + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u64i>>, !s32i) -> !cir.vector<2 x !u64i> + // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) + // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) + __nvvm_ldu_ll2((const longlong2 *)p); + __nvvm_ldu_ull2((const ulonglong2 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<2 x !cir.float>>, !s32i) -> !cir.vector<2 x !cir.float> + // LLVM: call {{.*}}<2 x float> @llvm.nvvm.ldu.global.f.v2f32.p0(ptr {{%[0-9]+}}, i32 8) + __nvvm_ldu_f2((const float2 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.float>>, !s32i) -> !cir.vector<4 x !cir.float> + // LLVM: call {{.*}}<4 x float> @llvm.nvvm.ldu.global.f.v4f32.p0(ptr {{%[0-9]+}}, i32 16) + __nvvm_ldu_f4((const float4 *)p); + + // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<2 x !cir.double>>, !s32i) -> !cir.vector<2 x !cir.double> + // LLVM: call {{.*}}<2 x double> @llvm.nvvm.ldu.global.f.v2f64.p0(ptr {{%[0-9]+}}, i32 16) + __nvvm_ldu_d2((const double2 *)p); +} diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu deleted file mode 100644 index cee0038d770da..0000000000000 --- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-ldu.cu +++ /dev/null @@ -1,155 +0,0 @@ -#include "Inputs/cuda.h" - -// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ -// RUN: -fcuda-is-device -fclangir -emit-cir %s -o %t.cir -// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s - -// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ -// RUN: -fcuda-is-device -fclangir -emit-llvm %s -o %t-cir.ll -// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s - -// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ -// 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 floating-point LLVM check lines use {{.*}} to tolerate the -// diff erence between CIR (no flags) and classic codegen ('contract'). - -typedef char char2 __attribute__((ext_vector_type(2))); -typedef unsigned char uchar2 __attribute__((ext_vector_type(2))); -typedef signed char schar2 __attribute__((ext_vector_type(2))); -typedef char char4 __attribute__((ext_vector_type(4))); -typedef unsigned char uchar4 __attribute__((ext_vector_type(4))); -typedef signed char schar4 __attribute__((ext_vector_type(4))); -typedef short short2 __attribute__((ext_vector_type(2))); -typedef unsigned short ushort2 __attribute__((ext_vector_type(2))); -typedef short short4 __attribute__((ext_vector_type(4))); -typedef unsigned short ushort4 __attribute__((ext_vector_type(4))); -typedef int int2 __attribute__((ext_vector_type(2))); -typedef unsigned int uint2 __attribute__((ext_vector_type(2))); -typedef int int4 __attribute__((ext_vector_type(4))); -typedef unsigned int uint4 __attribute__((ext_vector_type(4))); -typedef long long2 __attribute__((ext_vector_type(2))); -typedef unsigned long ulong2 __attribute__((ext_vector_type(2))); -typedef long long longlong2 __attribute__((ext_vector_type(2))); -typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2))); -typedef float float2 __attribute__((ext_vector_type(2))); -typedef float float4 __attribute__((ext_vector_type(4))); -typedef double double2 __attribute__((ext_vector_type(2))); - -// CIR-LABEL: @_Z8nvvm_lduPKv -// LLVM-LABEL: @_Z8nvvm_lduPKv -__device__ void nvvm_ldu(const void *p) { - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s8i>, !s32i) -> !s8i - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u8i>, !s32i) -> !u8i - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s8i>, !s32i) -> !s8i - // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1) - // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1) - // LLVM: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1) - __nvvm_ldu_c((const char *)p); - __nvvm_ldu_uc((const unsigned char *)p); - __nvvm_ldu_sc((const signed char *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s16i>, !s32i) -> !s16i - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u16i>, !s32i) -> !u16i - // LLVM: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2) - // LLVM: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2) - __nvvm_ldu_s((const short *)p); - __nvvm_ldu_us((const unsigned short *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s32i>, !s32i) -> !s32i - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u32i>, !s32i) -> !u32i - // LLVM: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4) - // LLVM: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4) - __nvvm_ldu_i((const int *)p); - __nvvm_ldu_ui((const unsigned int *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!s64i>, !s32i) -> !s64i - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!u64i>, !s32i) -> !u64i - // LLVM: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8) - // LLVM: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8) - __nvvm_ldu_l((const long *)p); - __nvvm_ldu_ul((const unsigned long *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.float>, !s32i) -> !cir.float - // LLVM: call {{.*}}float @llvm.nvvm.ldu.global.f.f32.p0(ptr {{%[0-9]+}}, i32 4) - __nvvm_ldu_f((const float *)p); - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.double>, !s32i) -> !cir.double - // LLVM: call {{.*}}double @llvm.nvvm.ldu.global.f.f64.p0(ptr {{%[0-9]+}}, i32 8) - __nvvm_ldu_d((const double *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s8i>>, !s32i) -> !cir.vector<2 x !s8i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u8i>>, !s32i) -> !cir.vector<2 x !u8i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s8i>>, !s32i) -> !cir.vector<2 x !s8i> - // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2) - // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2) - // LLVM: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2) - __nvvm_ldu_c2((const char2 *)p); - __nvvm_ldu_uc2((const uchar2 *)p); - __nvvm_ldu_sc2((const schar2 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s8i>>, !s32i) -> !cir.vector<4 x !s8i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u8i>>, !s32i) -> !cir.vector<4 x !u8i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s8i>>, !s32i) -> !cir.vector<4 x !s8i> - // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4) - // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4) - // LLVM: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4) - __nvvm_ldu_c4((const char4 *)p); - __nvvm_ldu_uc4((const uchar4 *)p); - __nvvm_ldu_sc4((const schar4 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s16i>>, !s32i) -> !cir.vector<2 x !s16i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u16i>>, !s32i) -> !cir.vector<2 x !u16i> - // LLVM: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4) - // LLVM: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4) - __nvvm_ldu_s2((const short2 *)p); - __nvvm_ldu_us2((const ushort2 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s16i>>, !s32i) -> !cir.vector<4 x !s16i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u16i>>, !s32i) -> !cir.vector<4 x !u16i> - // LLVM: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8) - // LLVM: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8) - __nvvm_ldu_s4((const short4 *)p); - __nvvm_ldu_us4((const ushort4 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>>, !s32i) -> !cir.vector<2 x !s32i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u32i>>, !s32i) -> !cir.vector<2 x !u32i> - // LLVM: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8) - // LLVM: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8) - __nvvm_ldu_i2((const int2 *)p); - __nvvm_ldu_ui2((const uint2 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !s32i>>, !s32i) -> !cir.vector<4 x !s32i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<4 x !u32i>>, !s32i) -> !cir.vector<4 x !u32i> - // LLVM: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16) - // LLVM: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16) - __nvvm_ldu_i4((const int4 *)p); - __nvvm_ldu_ui4((const uint4 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s64i>>, !s32i) -> !cir.vector<2 x !s64i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u64i>>, !s32i) -> !cir.vector<2 x !u64i> - // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) - // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) - __nvvm_ldu_l2((const long2 *)p); - __nvvm_ldu_ul2((const ulong2 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !s64i>>, !s32i) -> !cir.vector<2 x !s64i> - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.i" {{.*}} : (!cir.ptr<!cir.vector<2 x !u64i>>, !s32i) -> !cir.vector<2 x !u64i> - // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) - // LLVM: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16) - __nvvm_ldu_ll2((const longlong2 *)p); - __nvvm_ldu_ull2((const ulonglong2 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<2 x !cir.float>>, !s32i) -> !cir.vector<2 x !cir.float> - // LLVM: call {{.*}}<2 x float> @llvm.nvvm.ldu.global.f.v2f32.p0(ptr {{%[0-9]+}}, i32 8) - __nvvm_ldu_f2((const float2 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<4 x !cir.float>>, !s32i) -> !cir.vector<4 x !cir.float> - // LLVM: call {{.*}}<4 x float> @llvm.nvvm.ldu.global.f.v4f32.p0(ptr {{%[0-9]+}}, i32 16) - __nvvm_ldu_f4((const float4 *)p); - - // CIR: cir.call_llvm_intrinsic "nvvm.ldu.global.f" {{.*}} : (!cir.ptr<!cir.vector<2 x !cir.double>>, !s32i) -> !cir.vector<2 x !cir.double> - // LLVM: call {{.*}}<2 x double> @llvm.nvvm.ldu.global.f.v2f64.p0(ptr {{%[0-9]+}}, i32 16) - __nvvm_ldu_d2((const double2 *)p); -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
