Skip to content

Commit f9c8c01

Browse files
AlexMacleanAlexMaclean
authored
[NVPTX] Aggressively try to replace image handles with references (#119730)
Even in cases where handles are supported, references are still preferable for performance. This is because, a ref uses one less register and can avoid the handle creating code associated with taking the address of a tex/surf/sampler.
1 parent bfc2dbe commit f9c8c01

11 files changed

+231
-162
lines changed

‎llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

Lines changed: 15 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) {
209209
TargetMachine &TM = const_cast<TargetMachine &>(MF->getTarget());
210210
NVPTXTargetMachine &nvTM = static_cast<NVPTXTargetMachine &>(TM);
211211
const NVPTXMachineFunctionInfo *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
212-
const char *Sym = MFI->getImageHandleSymbol(Index);
212+
StringRef Sym = MFI->getImageHandleSymbol(Index);
213213
StringRef SymName = nvTM.getStrPool().save(Sym);
214214
MCOp = GetSymbolRef(OutContext.getOrCreateSymbol(SymName));
215215
}
@@ -224,16 +224,13 @@ void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
224224
return;
225225
}
226226

227-
const NVPTXSubtarget &STI = MI->getMF()->getSubtarget<NVPTXSubtarget>();
228227
for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) {
229228
const MachineOperand &MO = MI->getOperand(i);
230229

231230
MCOperand MCOp;
232-
if (!STI.hasImageHandles()) {
233-
if (lowerImageHandleOperand(MI, i, MCOp)) {
234-
OutMI.addOperand(MCOp);
235-
continue;
236-
}
231+
if (lowerImageHandleOperand(MI, i, MCOp)) {
232+
OutMI.addOperand(MCOp);
233+
continue;
237234
}
238235

239236
if (lowerOperand(MO, MCOp))
@@ -1509,13 +1506,14 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
15091506
const AttributeList &PAL = F->getAttributes();
15101507
const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
15111508
const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
1509+
const NVPTXMachineFunctionInfo *MFI =
1510+
MF ? MF->getInfo<NVPTXMachineFunctionInfo>() : nullptr;
15121511

15131512
Function::const_arg_iterator I, E;
15141513
unsigned paramIndex = 0;
15151514
bool first = true;
15161515
bool isKernelFunc = isKernelFunction(*F);
15171516
bool isABI = (STI.getSmVersion() >= 20);
1518-
bool hasImageHandles = STI.hasImageHandles();
15191517

15201518
if (F->arg_empty() && !F->isVarArg()) {
15211519
O << "()";
@@ -1533,25 +1531,30 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
15331531
first = false;
15341532

15351533
// Handle image/sampler parameters
1536-
if (isKernelFunction(*F)) {
1534+
if (isKernelFunc) {
15371535
if (isSampler(*I) || isImage(*I)) {
1536+
std::string ParamSym;
1537+
raw_string_ostream ParamStr(ParamSym);
1538+
ParamStr << F->getName() << "_param_" << paramIndex;
1539+
ParamStr.flush();
1540+
bool EmitImagePtr = !MFI || !MFI->checkImageHandleSymbol(ParamSym);
15381541
if (isImage(*I)) {
15391542
if (isImageWriteOnly(*I) || isImageReadWrite(*I)) {
1540-
if (hasImageHandles)
1543+
if (EmitImagePtr)
15411544
O << "\t.param .u64 .ptr .surfref ";
15421545
else
15431546
O << "\t.param .surfref ";
15441547
O << TLI->getParamName(F, paramIndex);
15451548
}
15461549
else { // Default image is read_only
1547-
if (hasImageHandles)
1550+
if (EmitImagePtr)
15481551
O << "\t.param .u64 .ptr .texref ";
15491552
else
15501553
O << "\t.param .texref ";
15511554
O << TLI->getParamName(F, paramIndex);
15521555
}
15531556
} else {
1554-
if (hasImageHandles)
1557+
if (EmitImagePtr)
15551558
O << "\t.param .u64 .ptr .samplerref ";
15561559
else
15571560
O << "\t.param .samplerref ";

‎llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h

Lines changed: 15 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -14,13 +14,14 @@
1414
#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXMACHINEFUNCTIONINFO_H
1515
#define LLVM_LIB_TARGET_NVPTX_NVPTXMACHINEFUNCTIONINFO_H
1616

17+
#include "llvm/ADT/StringRef.h"
1718
#include "llvm/CodeGen/MachineFunction.h"
1819

1920
namespace llvm {
2021
class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
2122
private:
22-
/// Stores a mapping from index to symbol name for removing image handles
23-
/// on Fermi.
23+
/// Stores a mapping from index to symbol name for image handles that are
24+
/// replaced with image references
2425
SmallVector<std::string, 8> ImageHandleList;
2526

2627
public:
@@ -36,20 +37,27 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
3637
/// Returns the index for the symbol \p Symbol. If the symbol was previously,
3738
/// added, the same index is returned. Otherwise, the symbol is added and the
3839
/// new index is returned.
39-
unsigned getImageHandleSymbolIndex(const char *Symbol) {
40+
unsigned getImageHandleSymbolIndex(StringRef Symbol) {
4041
// Is the symbol already present?
4142
for (unsigned i = 0, e = ImageHandleList.size(); i != e; ++i)
42-
if (ImageHandleList[i] == std::string(Symbol))
43+
if (ImageHandleList[i] == Symbol)
4344
return i;
4445
// Nope, insert it
45-
ImageHandleList.push_back(Symbol);
46+
ImageHandleList.push_back(Symbol.str());
4647
return ImageHandleList.size()-1;
4748
}
4849

4950
/// Returns the symbol name at the given index.
50-
const char *getImageHandleSymbol(unsigned Idx) const {
51+
StringRef getImageHandleSymbol(unsigned Idx) const {
5152
assert(ImageHandleList.size() > Idx && "Bad index");
52-
return ImageHandleList[Idx].c_str();
53+
return ImageHandleList[Idx];
54+
}
55+
56+
/// Check if the symbol has a mapping. Having a mapping means the handle is
57+
/// replaced with a reference
58+
bool checkImageHandleSymbol(StringRef Symbol) const {
59+
return ImageHandleList.end() !=
60+
std::find(ImageHandleList.begin(), ImageHandleList.end(), Symbol);
5361
}
5462
};
5563
}

‎llvm/lib/Target/NVPTX/NVPTXReplaceImageHandles.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1830,7 +1830,7 @@ findIndexForHandle(MachineOperand &Op, MachineFunction &MF, unsigned &Idx) {
18301830
NewSymStr << MF.getName() << "_param_" << Param;
18311831

18321832
InstrsToRemove.insert(&TexHandleDef);
1833-
Idx = MFI->getImageHandleSymbolIndex(NewSymStr.str().c_str());
1833+
Idx = MFI->getImageHandleSymbolIndex(NewSymStr.str());
18341834
return true;
18351835
}
18361836
case NVPTX::texsurf_handles: {
@@ -1839,7 +1839,7 @@ findIndexForHandle(MachineOperand &Op, MachineFunction &MF, unsigned &Idx) {
18391839
const GlobalValue *GV = TexHandleDef.getOperand(1).getGlobal();
18401840
assert(GV->hasName() && "Global sampler must be named!");
18411841
InstrsToRemove.insert(&TexHandleDef);
1842-
Idx = MFI->getImageHandleSymbolIndex(GV->getName().data());
1842+
Idx = MFI->getImageHandleSymbolIndex(GV->getName());
18431843
return true;
18441844
}
18451845
case NVPTX::nvvm_move_i64:

‎llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp

Lines changed: 1 addition & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -56,7 +56,7 @@ NVPTXSubtarget::NVPTXSubtarget(const Triple &TT, const std::string &CPU,
5656
const std::string &FS,
5757
const NVPTXTargetMachine &TM)
5858
: NVPTXGenSubtargetInfo(TT, CPU, /*TuneCPU*/ CPU, FS), PTXVersion(0),
59-
FullSmVersion(200), SmVersion(getSmVersion()), TM(TM),
59+
FullSmVersion(200), SmVersion(getSmVersion()),
6060
TLInfo(TM, initializeSubtargetDependencies(CPU, FS)) {
6161
TSInfo = std::make_unique<NVPTXSelectionDAGInfo>();
6262
}
@@ -67,16 +67,6 @@ const SelectionDAGTargetInfo *NVPTXSubtarget::getSelectionDAGInfo() const {
6767
return TSInfo.get();
6868
}
6969

70-
bool NVPTXSubtarget::hasImageHandles() const {
71-
// Enable handles for Kepler+, where CUDA supports indirect surfaces and
72-
// textures
73-
if (TM.getDrvInterface() == NVPTX::CUDA)
74-
return (SmVersion >= 30);
75-
76-
// Disabled, otherwise
77-
return false;
78-
}
79-
8070
bool NVPTXSubtarget::allowFP16Math() const {
8171
return hasFP16Math() && NoF16Math == false;
8272
}

‎llvm/lib/Target/NVPTX/NVPTXSubtarget.h

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,6 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
4242
// FullSmVersion.
4343
unsigned int SmVersion;
4444

45-
const NVPTXTargetMachine &TM;
4645
NVPTXInstrInfo InstrInfo;
4746
NVPTXTargetLowering TLInfo;
4847
std::unique_ptr<const SelectionDAGTargetInfo> TSInfo;
@@ -81,7 +80,6 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
8180
bool hasClusters() const { return SmVersion >= 90 && PTXVersion >= 78; }
8281
bool hasLDG() const { return SmVersion >= 32; }
8382
bool hasHWROT32() const { return SmVersion >= 32; }
84-
bool hasImageHandles() const;
8583
bool hasFP16Math() const { return SmVersion >= 53; }
8684
bool hasBF16Math() const { return SmVersion >= 80; }
8785
bool allowFP16Math() const;

‎llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -404,14 +404,10 @@ void NVPTXPassConfig::addIRPasses() {
404404
}
405405

406406
bool NVPTXPassConfig::addInstSelector() {
407-
const NVPTXSubtarget &ST = *getTM<NVPTXTargetMachine>().getSubtargetImpl();
408-
409407
addPass(createLowerAggrCopies());
410408
addPass(createAllocaHoisting());
411409
addPass(createNVPTXISelDag(getNVPTXTargetMachine(), getOptLevel()));
412-
413-
if (!ST.hasImageHandles())
414-
addPass(createNVPTXReplaceImageHandlesPass());
410+
addPass(createNVPTXReplaceImageHandlesPass());
415411

416412
return false;
417413
}

‎llvm/test/CodeGen/NVPTX/surf-read-cuda.ll

Lines changed: 32 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,6 @@
1-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
2-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
3+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s
34
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
45
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
56

@@ -9,38 +10,48 @@ declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32)
910
declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
1011

1112

12-
; SM20-LABEL: .entry foo
13-
; SM30-LABEL: .entry foo
1413
define void @foo(i64 %img, ptr %red, i32 %idx) {
15-
; SM20: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0];
16-
; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}]
17-
; SM30: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0];
18-
; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}]
14+
; CHECK-LABEL: foo(
15+
; CHECK: {
16+
; CHECK-NEXT: .reg .b32 %r<3>;
17+
; CHECK-NEXT: .reg .f32 %f<2>;
18+
; CHECK-NEXT: .reg .b64 %rd<4>;
19+
; CHECK-EMPTY:
20+
; CHECK-NEXT: // %bb.0:
21+
; CHECK-NEXT: ld.param.u64 %rd1, [foo_param_0];
22+
; CHECK-NEXT: ld.param.u64 %rd2, [foo_param_1];
23+
; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;
24+
; CHECK-NEXT: ld.param.u32 %r1, [foo_param_2];
25+
; CHECK-NEXT: suld.b.1d.b32.trap {%r2}, [%rd1, {%r1}];
26+
; CHECK-NEXT: cvt.rn.f32.s32 %f1, %r2;
27+
; CHECK-NEXT: st.global.f32 [%rd3], %f1;
28+
; CHECK-NEXT: ret;
1929
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)
20-
; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
21-
; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
2230
%ret = sitofp i32 %val to float
23-
; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
24-
; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
2531
store float %ret, ptr %red
2632
ret void
2733
}
2834

2935
@surf0 = internal addrspace(1) global i64 0, align 8
3036

31-
; SM20-LABEL: .entry bar
32-
; SM30-LABEL: .entry bar
3337
define void @bar(ptr %red, i32 %idx) {
34-
; SM30: mov.u64 %rd[[SURFHANDLE:[0-9]+]], surf0
38+
; CHECK-LABEL: bar(
39+
; CHECK: {
40+
; CHECK-NEXT: .reg .b32 %r<3>;
41+
; CHECK-NEXT: .reg .f32 %f<2>;
42+
; CHECK-NEXT: .reg .b64 %rd<4>;
43+
; CHECK-EMPTY:
44+
; CHECK-NEXT: // %bb.0:
45+
; CHECK-NEXT: ld.param.u64 %rd1, [bar_param_0];
46+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
47+
; CHECK-NEXT: ld.param.u32 %r1, [bar_param_1];
48+
; CHECK-NEXT: suld.b.1d.b32.trap {%r2}, [surf0, {%r1}];
49+
; CHECK-NEXT: cvt.rn.f32.s32 %f1, %r2;
50+
; CHECK-NEXT: st.global.f32 [%rd2], %f1;
51+
; CHECK-NEXT: ret;
3552
%surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)
36-
; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [surf0, {%r{{[0-9]+}}}]
37-
; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFHANDLE]], {%r{{[0-9]+}}}]
3853
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx)
39-
; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
40-
; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
4154
%ret = sitofp i32 %val to float
42-
; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
43-
; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
4455
store float %ret, ptr %red
4556
ret void
4657
}

‎llvm/test/CodeGen/NVPTX/surf-tex.py

Lines changed: 6 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,12 @@
11
# RUN: %python %s --target=cuda --tests=suld,sust,tex,tld4 --gen-list=%t.list > %t-cuda.ll
2-
# RUN: llc -mcpu=sm_60 -mattr=+ptx43 %t-cuda.ll -verify-machineinstrs -o - | FileCheck %t-cuda.ll --check-prefixes=CHECK,CHECK-CUDA
2+
# RUN: llc -mcpu=sm_60 -mattr=+ptx43 %t-cuda.ll -verify-machineinstrs -o - | FileCheck %t-cuda.ll
33
# RUN: %if ptxas %{ llc -mcpu=sm_60 -mattr=+ptx43 %t-cuda.ll -verify-machineinstrs -o - | %ptxas-verify %}
44

55
# We only need to run this second time for texture tests, because
66
# there is a difference between unified and non-unified intrinsics.
77
#
88
# RUN: %python %s --target=nvcl --tests=suld,sust,tex,tld4 --gen-list-append --gen-list=%t.list > %t-nvcl.ll
9-
# RUN: llc %t-nvcl.ll -verify-machineinstrs -o - | FileCheck %t-nvcl.ll --check-prefixes=CHECK,CHECK-NVCL
9+
# RUN: llc %t-nvcl.ll -verify-machineinstrs -o - | FileCheck %t-nvcl.ll
1010
# RUN: %if ptxas %{ llc %t-nvcl.ll -verify-machineinstrs -o - | %ptxas-verify %}
1111

1212
# Verify that all instructions and intrinsics defined in TableGen
@@ -269,9 +269,7 @@ def gen_suld_tests(target, global_surf):
269269
ret void
270270
}
271271
; CHECK-LABEL: .entry ${test_name}_global
272-
; CHECK-CUDA: mov.u64 [[REG${reg_id}:%.*]], ${global_surf}
273-
; CHECK-CUDA: ${instruction} ${reg_ret}, [[[REG${reg_id}]], ${reg_access}]
274-
; CHECK-NVCL: ${instruction} ${reg_ret}, [${global_surf}, ${reg_access}]
272+
; CHECK: ${instruction} ${reg_ret}, [${global_surf}, ${reg_access}]
275273
define void @${test_name}_global(${retty}* %ret, ${access}) {
276274
%gs = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @${global_surf})
277275
%val = tail call ${retty} @${intrinsic}(i64 %gs, ${access})
@@ -314,7 +312,6 @@ def gen_suld_tests(target, global_surf):
314312
"reg_ret": get_ptx_vec_reg(vec, dtype),
315313
"reg_surf": get_ptx_surface(target),
316314
"reg_access": get_ptx_surface_access(geom),
317-
"reg_id": get_table_gen_id(),
318315
}
319316
gen_test(template, params)
320317
generated_items.append((params["intrinsic"], params["instruction"]))
@@ -364,9 +361,7 @@ def gen_sust_tests(target, global_surf):
364361
ret void
365362
}
366363
; CHECK-LABEL: .entry ${test_name}_global
367-
; CHECK-CUDA: mov.u64 [[REG${reg_id}:%.*]], ${global_surf}
368-
; CHECK-CUDA: ${instruction} [[[REG${reg_id}]], ${reg_access}], ${reg_value}
369-
; CHECK-NVCL: ${instruction} [${global_surf}, ${reg_access}], ${reg_value}
364+
; CHECK: ${instruction} [${global_surf}, ${reg_access}], ${reg_value}
370365
define void @${test_name}_global(${value}, ${access}) {
371366
%gs = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @${global_surf})
372367
tail call void @${intrinsic}(i64 %gs, ${access}, ${value})
@@ -420,7 +415,6 @@ def gen_sust_tests(target, global_surf):
420415
"reg_value": get_ptx_vec_reg(vec, ctype),
421416
"reg_surf": get_ptx_surface(target),
422417
"reg_access": get_ptx_surface_access(geom),
423-
"reg_id": get_table_gen_id(),
424418
}
425419
gen_test(template, params)
426420
generated_items.append((params["intrinsic"], params["instruction"]))
@@ -627,9 +621,7 @@ def gen_tex_tests(target, global_tex, global_sampler):
627621
ret void
628622
}
629623
; CHECK-LABEL: .entry ${test_name}_global
630-
; CHECK-CUDA: mov.u64 [[REG${reg_id}:%.*]], ${global_tex}
631-
; CHECK-CUDA: ${instruction} ${ptx_ret}, [[[REG${reg_id}]], ${ptx_global_sampler} ${ptx_access}]
632-
; CHECK-NVCL: ${instruction} ${ptx_ret}, [${global_tex}, ${ptx_global_sampler} ${ptx_access}]
624+
; CHECK: ${instruction} ${ptx_ret}, [${global_tex}, ${ptx_global_sampler} ${ptx_access}]
633625
define void @${test_name}_global(${retty}* %ret, ${access}) {
634626
%gt = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @${global_tex})
635627
${get_sampler_handle}
@@ -713,7 +705,6 @@ def gen_tex_tests(target, global_tex, global_sampler):
713705
"ptx_tex": get_ptx_texture(target),
714706
"ptx_access": get_ptx_texture_access(geom, ctype),
715707
"ptx_global_sampler": get_ptx_global_sampler(target, global_sampler),
716-
"reg_id": get_table_gen_id(),
717708
}
718709
gen_test(template, params)
719710
generated_items.append((params["intrinsic"], params["instruction"]))
@@ -814,9 +805,7 @@ def gen_tld4_tests(target, global_tex, global_sampler):
814805
ret void
815806
}
816807
; CHECK-LABEL: .entry ${test_name}_global
817-
; CHECK-CUDA: mov.u64 [[REG${reg_id}:%.*]], ${global_tex}
818-
; CHECK-CUDA: ${instruction} ${ptx_ret}, [[[REG${reg_id}]], ${ptx_global_sampler} ${ptx_access}]
819-
; CHECK-NVCL: ${instruction} ${ptx_ret}, [${global_tex}, ${ptx_global_sampler} ${ptx_access}]
808+
; CHECK: ${instruction} ${ptx_ret}, [${global_tex}, ${ptx_global_sampler} ${ptx_access}]
820809
define void @${test_name}_global(${retty}* %ret, ${access}) {
821810
%gt = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @${global_tex})
822811
${get_sampler_handle}
@@ -862,7 +851,6 @@ def gen_tld4_tests(target, global_tex, global_sampler):
862851
"ptx_tex": get_ptx_texture(target),
863852
"ptx_access": get_ptx_tld4_access(geom),
864853
"ptx_global_sampler": get_ptx_global_sampler(target, global_sampler),
865-
"reg_id": get_table_gen_id(),
866854
}
867855
gen_test(template, params)
868856
generated_items.append((params["intrinsic"], params["instruction"]))

0 commit comments

Comments
 (0)