| /external/tensorflow/tensorflow/stream_executor/ |
| kernel_spec.cc | 50 CudaPtxInMemory::CudaPtxInMemory(port::StringPiece ptx,
57 // the original ptx is compressed.
58 decompressed_ptx_[ptx.data()] = "";
60 ptx_by_compute_capability_[kMinimumCapability] = ptx.data();
70 port::StringPiece ptx; local
71 std::tie(major, minor, ptx) = spec;
74 // that the original ptx is compressed.
75 decompressed_ptx_[ptx.data()] = "";
77 ptx_by_compute_capability_[std::tuple<int, int>{major, minor}] = ptx.data();
81 string CudaPtxInMemory::DecompressPtx(const char *ptx) {
100 auto ptx = ptx_by_compute_capability_.begin()->second; local
[all...] |
| kernel_spec.h | 36 // This lazily instantiates an object that describes how to load CUDA PTX
65 // The filename_or_text field represents the program location (i.e. PTX or
92 // canonical filename suffix is ".ptx".
101 // e.g. PTX files on disk have a canonical suffix of ".ptx".
114 // Kernel loader specification for PTX text that resides on disk.
120 const char *CanonicalSuffix() const override { return ".ptx"; }
142 // Kernel loader specification for PTX text that resides in memory.
146 // number, and PTX source.
149 // Single-PTX constructor. Adds the provided PTX version with an unknow
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
| gpu_compiler.h | 99 // Tries to compile the given ptx string to cubin. Returns a vector with the
101 std::vector<uint8> CompilePtxOrGetCachedResult(const string& ptx,
104 // The compilation_cache_ map is a cache from {ptx string, cc_major, cc_minor}
105 // -> cubin so we don't recompile the same ptx twice. This is important for
107 // we can't realize that two modules are the same until we lower to ptx.)
109 // Compilation of distinct PTX happens in parallel. If more than one thread
110 // attempts to compile the same PTX, the fist thread to obtain
114 // If compiling the ptx fails, we return an empty cubin, cross our fingers,
117 CompilationCacheKey(std::string ptx, int cc_major, int cc_minor)
118 : ptx(std::move(ptx)), cc_major(cc_major), cc_minor(cc_minor) {
119 string ptx; member in struct:xla::gpu::CompilationCacheKey
[all...] |
| kernel_thunk.cc | 46 tensorflow::StringPiece ptx = executable.ptx(); local
50 se::port::StringPiece(ptx.data(), ptx.size()), kernel_name_);
|
| gpu_executable.h | 50 // cubin (i.e. the compiled ptx) may be empty, in which case we leave
52 GpuExecutable(const string& ptx, const std::vector<uint8>& cubin,
68 // Returns the compiled PTX for the computation.
69 tensorflow::StringPiece ptx() const { return ptx_; } function in class:xla::gpu::GpuExecutable
71 // Returns the cubin (compiled PTX) stored in this GpuExecutable. May be
114 // The PTX for the computation.
|
| gpu_compiler.cc | 312 // Already checked this ptx binary, nothing to do.
359 // Prints a warning if the ptx->sass JIT in the driver has known bugs.
361 // Using such a driver only a problem if we fail to use ptxas to compile our ptx
389 << "*** WARNING *** Invoking the PTX->SASS JIT from driver version "
398 // Compiles the given PTX string using ptxas and returns the resulting machine
400 StatusOr<std::vector<uint8>> CompilePtx(const string& ptx, int cc_major,
402 Tracing::TraceMe annotation("Compile PTX", /*is_expensive=*/true);
411 // Write ptx into a temporary file.
414 return InternalError("couldn't get temp PTX file name");
420 TF_RETURN_IF_ERROR(tensorflow::WriteStringToFile(env, ptx_path, ptx));
609 string ptx; local
[all...] |
| gpu_executable.cc | 114 const string& ptx, const std::vector<uint8>& cubin,
123 ptx_(ptx),
|
| /prebuilts/go/darwin-x86/src/crypto/aes/ |
| gcm_amd64.s | 502 #define ptx SI define
553 MOVQ src_base+32(FP), ptx
639 MOVOU (16*0)(ptx), T0
641 MOVOU (16*1)(ptx), T0
643 MOVOU (16*2)(ptx), T0
645 MOVOU (16*3)(ptx), T0
647 MOVOU (16*4)(ptx), T0
649 MOVOU (16*5)(ptx), T0
651 MOVOU (16*6)(ptx), T0
653 MOVOU (16*7)(ptx), T
[all...] |
| /prebuilts/go/linux-x86/src/crypto/aes/ |
| gcm_amd64.s | 502 #define ptx SI define
553 MOVQ src_base+32(FP), ptx
639 MOVOU (16*0)(ptx), T0
641 MOVOU (16*1)(ptx), T0
643 MOVOU (16*2)(ptx), T0
645 MOVOU (16*3)(ptx), T0
647 MOVOU (16*4)(ptx), T0
649 MOVOU (16*5)(ptx), T0
651 MOVOU (16*6)(ptx), T0
653 MOVOU (16*7)(ptx), T
[all...] |
| /external/mesa3d/src/gallium/drivers/nouveau/nvc0/ |
| nvc0_resource.h | 56 struct pipe_transfer *ptx);
|
| /external/libldac/src/ |
| ldacBT_internal.c | 306 LDACBT_TX_INFO *ptx; local
315 ptx = &hLdacBT->tx;
361 ptx->nfrm_in_pkt = ptx->tx_size / hLdacBT->frmlen_tx;
362 if( ptx->nfrm_in_pkt > LDACBT_NFRM_TX_MAX ){
363 ptx->nfrm_in_pkt = LDACBT_NFRM_TX_MAX;
365 else if( ptx->nfrm_in_pkt < 2 ){
367 if( frmlen <= (ptx->tx_size / 2 - LDACBT_FRMHDRBYTES)){
370 frmlen = ptx->tx_size / 2 - LDACBT_FRMHDRBYTES;
375 hLdacBT->eqmid = ldacBT_get_eqmid_from_frmlen( frmlen, ch, hLdacBT->transport, ptx->pkt_type )
[all...] |
| /sdk/eclipse/scripts/ |
| gen_icon.py | 61 ptx = (sz4 - tsx) / 2
64 draw.text((ptx + i, pty), data.letter, font=font, fill=data.letter_color)
|
| /external/tensorflow/tensorflow/compiler/xla/service/gpu/llvm_gpu_backend/ |
| gpu_backend_lib.cc | 259 // Emits the given module to PTX. target_machine is an initialized TargetMachine
262 std::string ptx; // need a std::string instead of a ::string. local
264 llvm::raw_string_ostream stream(ptx);
281 return ptx;
437 // Finally, produce PTX.
492 string ptx; local
500 ptx, CompileModuleToPtx(module, compute_capability, hlo_module_config,
503 return ptx;
|
| /external/tensorflow/tensorflow/stream_executor/cuda/ |
| cuda_gpu_executor.cc | 73 // Hook that can be used to CUBIN-ate PTX before it is loaded into the driver.
74 // It has been observed that loading both PTX and cubins into the driver library
77 // PTX code.
248 const char *ptx = spec.cuda_ptx_in_memory().text(cc_major_, cc_minor_); local
249 if (ptx == nullptr) {
250 ptx = spec.cuda_ptx_in_memory().default_text();
252 if (ptx == nullptr) {
253 LOG(FATAL) << "loader spec has no ptx for kernel " << *kernelname;
259 std::tie(module, module_refcount) = gpu_binary_to_module_[ptx];
262 if (!CUDADriver::LoadPtx(context_, ptx, &module))
[all...] |
| /external/mesa3d/src/gallium/drivers/nouveau/nv50/ |
| nv50_resource.h | 138 struct pipe_transfer *ptx);
|
| /external/wpa_supplicant_8/src/pae/ |
| ieee802_1x_kay_i.h | 255 * @ptx: plain TX, ie protectFrames is False
293 u8 ptx:1; member in struct:ieee802_1x_mka_sak_use_body
295 u8 ptx:1; member in struct:ieee802_1x_mka_sak_use_body
|
| /prebuilts/go/darwin-x86/test/ |
| sinit.go | 205 var ptx *T var
271 // var copy_ptx = ptx
|
| /prebuilts/go/linux-x86/test/ |
| sinit.go | 205 var ptx *T var
271 // var copy_ptx = ptx
|
| /external/mesa3d/src/gallium/drivers/nouveau/nv30/ |
| nv30_miptree.c | 83 nv30_transfer(struct pipe_transfer *ptx)
85 return (struct nv30_transfer *)ptx;
337 struct pipe_transfer *ptx)
340 struct nv30_transfer *tx = nv30_transfer(ptx);
342 if (ptx->usage & PIPE_TRANSFER_WRITE) {
351 pipe_resource_reference(&ptx->resource, NULL);
|
| /external/cmockery/cmockery_0_1_2/ |
| config.sub | 210 -ptx*)
1180 | -ptx* | -coff* | -ecoff* | -winnt* | -domain* | -vsta* \
1444 os=-ptx
1542 -ptx*)
|
| /external/fec/ |
| config.sub | 199 -ptx*)
1019 | -ptx* | -coff* | -ecoff* | -winnt* | -domain* | -vsta* \
1236 os=-ptx
1331 -ptx*)
|
| /external/google-tv-pairing-protocol/cpp/ |
| config.sub | 224 -ptx*)
1281 | -ptx* | -coff* | -ecoff* | -winnt* | -domain* | -vsta* \
1558 os=-ptx
[all...] |
| /external/libexif/ |
| config.sub | 227 -ptx*)
1293 | -ptx* | -coff* | -ecoff* | -winnt* | -domain* | -vsta* \
1570 os=-ptx
[all...] |
| /external/libnetfilter_conntrack/ |
| config.sub | 228 -ptx*)
1300 | -ptx* | -coff* | -ecoff* | -winnt* | -domain* | -vsta* \
[all...] |
| /external/libnfnetlink/ |
| config.sub | 228 -ptx*)
1300 | -ptx* | -coff* | -ecoff* | -winnt* | -domain* | -vsta* \
[all...] |
