From 7c65b1bfb266a5f6e260b8e4ec685b171db81906 Mon Sep 17 00:00:00 2001 From: Jingyue Wu Date: Thu, 4 Jun 2015 21:07:08 +0000 Subject: [PATCH] Revert r239082 llc crashed for NVPTX backend git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@239094 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/NVPTX/CMakeLists.txt | 2 +- lib/Target/NVPTX/NVPTX.h | 2 +- lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp | 170 ------------------ lib/Target/NVPTX/NVPTXLowerStructArgs.cpp | 136 ++++++++++++++ lib/Target/NVPTX/NVPTXTargetMachine.cpp | 5 +- test/CodeGen/NVPTX/bug21465.ll | 7 +- test/CodeGen/NVPTX/call-with-alloca-buffer.ll | 7 +- test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll | 20 --- test/CodeGen/NVPTX/pr13291-i1-store.ll | 8 +- test/CodeGen/NVPTX/surf-read-cuda.ll | 8 +- test/CodeGen/NVPTX/tex-read-cuda.ll | 8 +- 11 files changed, 158 insertions(+), 215 deletions(-) delete mode 100644 lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp create mode 100644 lib/Target/NVPTX/NVPTXLowerStructArgs.cpp delete mode 100644 test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll diff --git a/lib/Target/NVPTX/CMakeLists.txt b/lib/Target/NVPTX/CMakeLists.txt index d48a7a9b1fc..cdd2f1f5944 100644 --- a/lib/Target/NVPTX/CMakeLists.txt +++ b/lib/Target/NVPTX/CMakeLists.txt @@ -20,7 +20,7 @@ set(NVPTXCodeGen_sources NVPTXImageOptimizer.cpp NVPTXInstrInfo.cpp NVPTXLowerAggrCopies.cpp - NVPTXLowerKernelArgs.cpp + NVPTXLowerStructArgs.cpp NVPTXMCExpr.cpp NVPTXPrologEpilogPass.cpp NVPTXRegisterInfo.cpp diff --git a/lib/Target/NVPTX/NVPTX.h b/lib/Target/NVPTX/NVPTX.h index 477b0bac6ca..382525d27a2 100644 --- a/lib/Target/NVPTX/NVPTX.h +++ b/lib/Target/NVPTX/NVPTX.h @@ -69,7 +69,7 @@ ModulePass *createNVVMReflectPass(const StringMap& Mapping); MachineFunctionPass *createNVPTXPrologEpilogPass(); MachineFunctionPass *createNVPTXReplaceImageHandlesPass(); FunctionPass *createNVPTXImageOptimizerPass(); -FunctionPass *createNVPTXLowerKernelArgsPass(const NVPTXTargetMachine *TM); +FunctionPass *createNVPTXLowerStructArgsPass(); bool isImageOrSamplerVal(const Value *, const Module *); diff --git a/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp b/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp deleted file mode 100644 index 24dcb122b94..00000000000 --- a/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp +++ /dev/null @@ -1,170 +0,0 @@ -//===-- NVPTXLowerKernelArgs.cpp - Lower kernel arguments -----------------===// -// -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. -// -//===----------------------------------------------------------------------===// -// -// Pointer arguments to kernel functions need to be lowered specially. -// -// 1. Copy byval struct args to local memory. This is a preparation for handling -// cases like -// -// kernel void foo(struct A arg, ...) -// { -// struct A *p = &arg; -// ... -// ... = p->filed1 ... (this is no generic address for .param) -// p->filed2 = ... (this is no write access to .param) -// } -// -// 2. Convert non-byval pointer arguments of CUDA kernels to pointers in the -// global address space. This allows later optimizations to emit -// ld.global.*/st.global.* for accessing these pointer arguments. For -// example, -// -// define void @foo(float* %input) { -// %v = load float, float* %input, align 4 -// ... -// } -// -// becomes -// -// define void @foo(float* %input) { -// %input2 = addrspacecast float* %input to float addrspace(1)* -// %input3 = addrspacecast float addrspace(1)* %input2 to float* -// %v = load float, float* %input3, align 4 -// ... -// } -// -// Later, NVPTXFavorNonGenericAddrSpaces will optimize it to -// -// define void @foo(float* %input) { -// %input2 = addrspacecast float* %input to float addrspace(1)* -// %v = load float, float addrspace(1)* %input2, align 4 -// ... -// } -// -// TODO: merge this pass with NVPTXFavorNonGenericAddrSpace so that other passes -// don't cancel the addrspacecast pair this pass emits. -//===----------------------------------------------------------------------===// - -#include "NVPTX.h" -#include "NVPTXUtilities.h" -#include "NVPTXTargetMachine.h" -#include "llvm/IR/Function.h" -#include "llvm/IR/Instructions.h" -#include "llvm/IR/Module.h" -#include "llvm/IR/Type.h" -#include "llvm/Pass.h" - -using namespace llvm; - -namespace llvm { -void initializeNVPTXLowerKernelArgsPass(PassRegistry &); -} - -namespace { -class NVPTXLowerKernelArgs : public FunctionPass { - bool runOnFunction(Function &F) override; - - // handle byval parameters - void handleByValParam(Argument *); - // handle non-byval pointer parameters - void handlePointerParam(Argument *); - -public: - static char ID; // Pass identification, replacement for typeid - NVPTXLowerKernelArgs(const NVPTXTargetMachine *TM = nullptr) - : FunctionPass(ID), TM(TM) {} - const char *getPassName() const override { - return "Lower pointer arguments of CUDA kernels"; - } - -private: - const NVPTXTargetMachine *TM; -}; -} // namespace - -char NVPTXLowerKernelArgs::ID = 1; - -INITIALIZE_PASS(NVPTXLowerKernelArgs, "nvptx-lower-kernel-args", - "Lower kernel arguments (NVPTX)", false, false) - -// ============================================================================= -// If the function had a byval struct ptr arg, say foo(%struct.x *byval %d), -// then add the following instructions to the first basic block: -// -// %temp = alloca %struct.x, align 8 -// %tempd = addrspacecast %struct.x* %d to %struct.x addrspace(101)* -// %tv = load %struct.x addrspace(101)* %tempd -// store %struct.x %tv, %struct.x* %temp, align 8 -// -// The above code allocates some space in the stack and copies the incoming -// struct from param space to local space. -// Then replace all occurences of %d by %temp. -// ============================================================================= -void NVPTXLowerKernelArgs::handleByValParam(Argument *Arg) { - Function *Func = Arg->getParent(); - Instruction *FirstInst = &(Func->getEntryBlock().front()); - PointerType *PType = dyn_cast(Arg->getType()); - - assert(PType && "Expecting pointer type in handleByValParam"); - - Type *StructType = PType->getElementType(); - AllocaInst *AllocA = new AllocaInst(StructType, Arg->getName(), FirstInst); - // Set the alignment to alignment of the byval parameter. This is because, - // later load/stores assume that alignment, and we are going to replace - // the use of the byval parameter with this alloca instruction. - AllocA->setAlignment(Func->getParamAlignment(Arg->getArgNo() + 1)); - Arg->replaceAllUsesWith(AllocA); - - Value *ArgInParam = new AddrSpaceCastInst( - Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(), - FirstInst); - LoadInst *LI = new LoadInst(ArgInParam, Arg->getName(), FirstInst); - new StoreInst(LI, AllocA, FirstInst); -} - -void NVPTXLowerKernelArgs::handlePointerParam(Argument *Arg) { - assert(!Arg->hasByValAttr() && - "byval params should be handled by handleByValParam"); - - Instruction *FirstInst = Arg->getParent()->getEntryBlock().begin(); - Instruction *ArgInGlobal = new AddrSpaceCastInst( - Arg, PointerType::get(Arg->getType()->getPointerElementType(), - ADDRESS_SPACE_GLOBAL), - Arg->getName(), FirstInst); - Value *ArgInGeneric = new AddrSpaceCastInst(ArgInGlobal, Arg->getType(), - Arg->getName(), FirstInst); - // Replace with ArgInGeneric all uses of Args except ArgInGlobal. - Arg->replaceAllUsesWith(ArgInGeneric); - ArgInGlobal->setOperand(0, Arg); -} - - -// ============================================================================= -// Main function for this pass. -// ============================================================================= -bool NVPTXLowerKernelArgs::runOnFunction(Function &F) { - // Skip non-kernels. See the comments at the top of this file. - if (!isKernelFunction(F)) - return false; - - for (Argument &Arg : F.args()) { - if (Arg.getType()->isPointerTy()) { - if (Arg.hasByValAttr()) - handleByValParam(&Arg); - else if (TM && TM->getDrvInterface() == NVPTX::CUDA) - handlePointerParam(&Arg); - } - } - return true; -} - -FunctionPass * -llvm::createNVPTXLowerKernelArgsPass(const NVPTXTargetMachine *TM) { - return new NVPTXLowerKernelArgs(TM); -} diff --git a/lib/Target/NVPTX/NVPTXLowerStructArgs.cpp b/lib/Target/NVPTX/NVPTXLowerStructArgs.cpp new file mode 100644 index 00000000000..68dfbb71613 --- /dev/null +++ b/lib/Target/NVPTX/NVPTXLowerStructArgs.cpp @@ -0,0 +1,136 @@ +//===-- NVPTXLowerStructArgs.cpp - Copy struct args to local memory =====--===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +// +// Copy struct args to local memory. This is needed for kernel functions only. +// This is a preparation for handling cases like +// +// kernel void foo(struct A arg, ...) +// { +// struct A *p = &arg; +// ... +// ... = p->filed1 ... (this is no generic address for .param) +// p->filed2 = ... (this is no write access to .param) +// } +// +//===----------------------------------------------------------------------===// + +#include "NVPTX.h" +#include "NVPTXUtilities.h" +#include "llvm/IR/Function.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/IntrinsicInst.h" +#include "llvm/IR/Module.h" +#include "llvm/IR/Type.h" +#include "llvm/Pass.h" + +using namespace llvm; + +namespace llvm { +void initializeNVPTXLowerStructArgsPass(PassRegistry &); +} + +namespace { +class NVPTXLowerStructArgs : public FunctionPass { + bool runOnFunction(Function &F) override; + + void handleStructPtrArgs(Function &); + void handleParam(Argument *); + +public: + static char ID; // Pass identification, replacement for typeid + NVPTXLowerStructArgs() : FunctionPass(ID) {} + const char *getPassName() const override { + return "Copy structure (byval *) arguments to stack"; + } +}; +} // namespace + +char NVPTXLowerStructArgs::ID = 1; + +INITIALIZE_PASS(NVPTXLowerStructArgs, "nvptx-lower-struct-args", + "Lower structure arguments (NVPTX)", false, false) + +void NVPTXLowerStructArgs::handleParam(Argument *Arg) { + Function *Func = Arg->getParent(); + Instruction *FirstInst = &(Func->getEntryBlock().front()); + PointerType *PType = dyn_cast(Arg->getType()); + + assert(PType && "Expecting pointer type in handleParam"); + + Type *StructType = PType->getElementType(); + AllocaInst *AllocA = new AllocaInst(StructType, Arg->getName(), FirstInst); + + /* Set the alignment to alignment of the byval parameter. This is because, + * later load/stores assume that alignment, and we are going to replace + * the use of the byval parameter with this alloca instruction. + */ + AllocA->setAlignment(Func->getParamAlignment(Arg->getArgNo() + 1)); + + Arg->replaceAllUsesWith(AllocA); + + // Get the cvt.gen.to.param intrinsic + Type *CvtTypes[] = { + Type::getInt8PtrTy(Func->getParent()->getContext(), ADDRESS_SPACE_PARAM), + Type::getInt8PtrTy(Func->getParent()->getContext(), + ADDRESS_SPACE_GENERIC)}; + Function *CvtFunc = Intrinsic::getDeclaration( + Func->getParent(), Intrinsic::nvvm_ptr_gen_to_param, CvtTypes); + + Value *BitcastArgs[] = { + new BitCastInst(Arg, Type::getInt8PtrTy(Func->getParent()->getContext(), + ADDRESS_SPACE_GENERIC), + Arg->getName(), FirstInst)}; + CallInst *CallCVT = + CallInst::Create(CvtFunc, BitcastArgs, "cvt_to_param", FirstInst); + + BitCastInst *BitCast = new BitCastInst( + CallCVT, PointerType::get(StructType, ADDRESS_SPACE_PARAM), + Arg->getName(), FirstInst); + LoadInst *LI = new LoadInst(BitCast, Arg->getName(), FirstInst); + new StoreInst(LI, AllocA, FirstInst); +} + +// ============================================================================= +// If the function had a struct ptr arg, say foo(%struct.x *byval %d), then +// add the following instructions to the first basic block : +// +// %temp = alloca %struct.x, align 8 +// %tt1 = bitcast %struct.x * %d to i8 * +// %tt2 = llvm.nvvm.cvt.gen.to.param %tt2 +// %tempd = bitcast i8 addrspace(101) * to %struct.x addrspace(101) * +// %tv = load %struct.x addrspace(101) * %tempd +// store %struct.x %tv, %struct.x * %temp, align 8 +// +// The above code allocates some space in the stack and copies the incoming +// struct from param space to local space. +// Then replace all occurences of %d by %temp. +// ============================================================================= +void NVPTXLowerStructArgs::handleStructPtrArgs(Function &F) { + for (Argument &Arg : F.args()) { + if (Arg.getType()->isPointerTy() && Arg.hasByValAttr()) { + handleParam(&Arg); + } + } +} + +// ============================================================================= +// Main function for this pass. +// ============================================================================= +bool NVPTXLowerStructArgs::runOnFunction(Function &F) { + // Skip non-kernels. See the comments at the top of this file. + if (!isKernelFunction(F)) + return false; + + handleStructPtrArgs(F); + return true; +} + +FunctionPass *llvm::createNVPTXLowerStructArgsPass() { + return new NVPTXLowerStructArgs(); +} diff --git a/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/lib/Target/NVPTX/NVPTXTargetMachine.cpp index 7fdfdbfb83f..c839ef0289b 100644 --- a/lib/Target/NVPTX/NVPTXTargetMachine.cpp +++ b/lib/Target/NVPTX/NVPTXTargetMachine.cpp @@ -53,7 +53,7 @@ void initializeGenericToNVVMPass(PassRegistry&); void initializeNVPTXAllocaHoistingPass(PassRegistry &); void initializeNVPTXAssignValidGlobalNamesPass(PassRegistry&); void initializeNVPTXFavorNonGenericAddrSpacesPass(PassRegistry &); -void initializeNVPTXLowerKernelArgsPass(PassRegistry &); +void initializeNVPTXLowerStructArgsPass(PassRegistry &); } extern "C" void LLVMInitializeNVPTXTarget() { @@ -69,7 +69,7 @@ extern "C" void LLVMInitializeNVPTXTarget() { initializeNVPTXAssignValidGlobalNamesPass(*PassRegistry::getPassRegistry()); initializeNVPTXFavorNonGenericAddrSpacesPass( *PassRegistry::getPassRegistry()); - initializeNVPTXLowerKernelArgsPass(*PassRegistry::getPassRegistry()); + initializeNVPTXLowerStructArgsPass(*PassRegistry::getPassRegistry()); } static std::string computeDataLayout(bool is64Bit) { @@ -163,7 +163,6 @@ void NVPTXPassConfig::addIRPasses() { TargetPassConfig::addIRPasses(); addPass(createNVPTXAssignValidGlobalNamesPass()); addPass(createGenericToNVVMPass()); - addPass(createNVPTXLowerKernelArgsPass(&getNVPTXTargetMachine())); addPass(createNVPTXFavorNonGenericAddrSpacesPass()); // FavorNonGenericAddrSpaces shortcuts unnecessary addrspacecasts, and leave // them unused. We could remove dead code in an ad-hoc manner, but that diff --git a/test/CodeGen/NVPTX/bug21465.ll b/test/CodeGen/NVPTX/bug21465.ll index e998e97b8ee..76af386c651 100644 --- a/test/CodeGen/NVPTX/bug21465.ll +++ b/test/CodeGen/NVPTX/bug21465.ll @@ -1,4 +1,4 @@ -; RUN: opt < %s -nvptx-lower-kernel-args -S | FileCheck %s +; RUN: opt < %s -nvptx-lower-struct-args -S | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" target triple = "nvptx64-unknown-unknown" @@ -8,8 +8,9 @@ target triple = "nvptx64-unknown-unknown" ; Function Attrs: nounwind define void @_Z11TakesStruct1SPi(%struct.S* byval nocapture readonly %input, i32* nocapture %output) #0 { entry: -; CHECK-LABEL: @_Z11TakesStruct1SPi -; CHECK: addrspacecast %struct.S* %input to %struct.S addrspace(101)* +; CHECK-LABEL @_Z22TakesStruct1SPi +; CHECK: bitcast %struct.S* %input to i8* +; CHECK: call i8 addrspace(101)* @llvm.nvvm.ptr.gen.to.param.p101i8.p0i8 %b = getelementptr inbounds %struct.S, %struct.S* %input, i64 0, i32 1 %0 = load i32, i32* %b, align 4 store i32 %0, i32* %output, align 4 diff --git a/test/CodeGen/NVPTX/call-with-alloca-buffer.ll b/test/CodeGen/NVPTX/call-with-alloca-buffer.ll index c70670da13d..58b19112991 100644 --- a/test/CodeGen/NVPTX/call-with-alloca-buffer.ll +++ b/test/CodeGen/NVPTX/call-with-alloca-buffer.ll @@ -24,10 +24,7 @@ entry: ; CHECK: cvta.local.u64 %SP, %rd[[BUF_REG]] ; CHECK: ld.param.u64 %rd[[A_REG:[0-9]+]], [kernel_func_param_0] -; CHECK: cvta.to.global.u64 %rd[[A1_REG:[0-9]+]], %rd[[A_REG]] -; FIXME: casting A1_REG to A2_REG is unnecessary; A2_REG is essentially A_REG -; CHECK: cvta.global.u64 %rd[[A2_REG:[0-9]+]], %rd[[A1_REG]] -; CHECK: ld.global.f32 %f[[A0_REG:[0-9]+]], [%rd[[A1_REG]]] +; CHECK: ld.f32 %f[[A0_REG:[0-9]+]], [%rd[[A_REG]]] ; CHECK: st.f32 [%SP+0], %f[[A0_REG]] %0 = load float, float* %a, align 4 @@ -51,7 +48,7 @@ entry: ; CHECK: add.u64 %rd[[SP_REG:[0-9]+]], %SP, 0 ; CHECK: .param .b64 param0; -; CHECK-NEXT: st.param.b64 [param0+0], %rd[[A2_REG]] +; CHECK-NEXT: st.param.b64 [param0+0], %rd[[A_REG]] ; CHECK-NEXT: .param .b64 param1; ; CHECK-NEXT: st.param.b64 [param1+0], %rd[[SP_REG]] ; CHECK-NEXT: call.uni diff --git a/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll b/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll deleted file mode 100644 index 53220bd905b..00000000000 --- a/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll +++ /dev/null @@ -1,20 +0,0 @@ -; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s - -target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" -target triple = "nvptx64-unknown-unknown" - -; Verify that both %input and %output are converted to global pointers and then -; addrspacecast'ed back to the original type. -define void @kernel(float* %input, float* %output) { -; CHECK-LABEL: .visible .entry kernel( -; CHECK: cvta.to.global.u64 -; CHECK: cvta.to.global.u64 - %1 = load float, float* %input, align 4 -; CHECK: ld.global.f32 - store float %1, float* %output, align 4 -; CHECK: st.global.f32 - ret void -} - -!nvvm.annotations = !{!0} -!0 = !{void (float*, float*)* @kernel, !"kernel", i32 1} diff --git a/test/CodeGen/NVPTX/pr13291-i1-store.ll b/test/CodeGen/NVPTX/pr13291-i1-store.ll index 934df30a3a7..d4f7c3bd210 100644 --- a/test/CodeGen/NVPTX/pr13291-i1-store.ll +++ b/test/CodeGen/NVPTX/pr13291-i1-store.ll @@ -3,19 +3,19 @@ define ptx_kernel void @t1(i1* %a) { ; PTX32: mov.u16 %rs{{[0-9]+}}, 0; -; PTX32-NEXT: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}; +; PTX32-NEXT: st.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}; ; PTX64: mov.u16 %rs{{[0-9]+}}, 0; -; PTX64-NEXT: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}; +; PTX64-NEXT: st.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}; store i1 false, i1* %a ret void } define ptx_kernel void @t2(i1* %a, i8* %b) { -; PTX32: ld.global.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}] +; PTX32: ld.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: and.b16 %rs{{[0-9]+}}, %rs{{[0-9]+}}, 1; ; PTX32: setp.eq.b16 %p{{[0-9]+}}, %rs{{[0-9]+}}, 1; -; PTX64: ld.global.u8 %rs{{[0-9]+}}, [%rd{{[0-9]+}}] +; PTX64: ld.u8 %rs{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: and.b16 %rs{{[0-9]+}}, %rs{{[0-9]+}}, 1; ; PTX64: setp.eq.b16 %p{{[0-9]+}}, %rs{{[0-9]+}}, 1; diff --git a/test/CodeGen/NVPTX/surf-read-cuda.ll b/test/CodeGen/NVPTX/surf-read-cuda.ll index c17c71e01d3..ed021346c0f 100644 --- a/test/CodeGen/NVPTX/surf-read-cuda.ll +++ b/test/CodeGen/NVPTX/surf-read-cuda.ll @@ -18,8 +18,8 @@ define void @foo(i64 %img, float* %red, i32 %idx) { ; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] %ret = sitofp i32 %val to float -; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] -; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] +; SM20: st.f32 [%r{{[0-9]+}}], %f[[REDF]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[REDF]] store float %ret, float* %red ret void } @@ -37,8 +37,8 @@ define void @bar(float* %red, i32 %idx) { ; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] %ret = sitofp i32 %val to float -; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] -; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] +; SM20: st.f32 [%r{{[0-9]+}}], %f[[REDF]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[REDF]] store float %ret, float* %red ret void } diff --git a/test/CodeGen/NVPTX/tex-read-cuda.ll b/test/CodeGen/NVPTX/tex-read-cuda.ll index d5f7c1667f1..c5b5600de87 100644 --- a/test/CodeGen/NVPTX/tex-read-cuda.ll +++ b/test/CodeGen/NVPTX/tex-read-cuda.ll @@ -16,8 +16,8 @@ define void @foo(i64 %img, float* %red, i32 %idx) { ; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx) %ret = extractvalue { float, float, float, float } %val, 0 -; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] -; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] +; SM20: st.f32 [%r{{[0-9]+}}], %f[[RED]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[RED]] store float %ret, float* %red ret void } @@ -34,8 +34,8 @@ define void @bar(float* %red, i32 %idx) { ; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}] %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx) %ret = extractvalue { float, float, float, float } %val, 0 -; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] -; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] +; SM20: st.f32 [%r{{[0-9]+}}], %f[[RED]] +; SM30: st.f32 [%r{{[0-9]+}}], %f[[RED]] store float %ret, float* %red ret void }