From e920848d1a39ec3b401126fdf957be51757e6e4a Mon Sep 17 00:00:00 2001 From: morluto <76467478+morluto@users.noreply.github.com> Date: Thu, 10 Sep 2026 17:02:40 +0800 Subject: [PATCH] perf: reduce redundant Top-K host setup --- .../kerutils/include/kerutils/host/host.h | 144 ++++++++------ csrc/cuda_kernels/common_parts.cuh | 4 +- csrc/cuda_kernels/v3/topk_select.cuh | 1 - csrc/cuda_kernels/v3_cluster/topk_select.cuh | 4 - csrc/cuda_kernels/v3_fp32/topk_select.cuh | 1 - tests/host/README.md | 17 ++ tests/host/allocations.cpp | 13 ++ tests/host/tensor_map.cu | 176 ++++++++++++++++++ tests/host/tensor_map_peer.cu | 9 + 9 files changed, 308 insertions(+), 61 deletions(-) create mode 100644 tests/host/README.md create mode 100644 tests/host/allocations.cpp create mode 100644 tests/host/tensor_map.cu create mode 100644 tests/host/tensor_map_peer.cu diff --git a/csrc/3rdparty/kerutils/include/kerutils/host/host.h b/csrc/3rdparty/kerutils/include/kerutils/host/host.h index e4be82f..d1b0150 100644 --- a/csrc/3rdparty/kerutils/include/kerutils/host/host.h +++ b/csrc/3rdparty/kerutils/include/kerutils/host/host.h @@ -1,6 +1,7 @@ #pragma once #include +#include #include #include #include @@ -97,81 +98,70 @@ inline __host__ __device__ constexpr T find_next_power_of_2(const T& x) { return find_next_power_of_2(x); } -// A wrapper for make_tensor_map -static inline CUtensorMap make_tensor_map( - const std::vector &size, - const std::vector &strides, // PAY ATTENTION: In BYTES - const std::vector &box_size, - void* global_ptr, - CUtensorMapDataType data_type, - CUtensorMapSwizzle swizzle_mode, - CUtensorMapL2promotion l2_promotion, - CUtensorMapInterleave interleave_mode = CUtensorMapInterleave::CU_TENSOR_MAP_INTERLEAVE_NONE, - CUtensorMapFloatOOBfill oob_fill = CUtensorMapFloatOOBfill::CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE, - const std::vector &element_strides_ = {} -) { - int dim = size.size(); - KU_ASSERT(dim >= 1); - - std::vector element_strides; - if (element_strides_.empty()) { - for (int i = 0; i < dim; ++i) - element_strides.push_back(1); - } else { - element_strides = element_strides_; - } - KU_ASSERT(strides.size() == (uint32_t)dim-1 && box_size.size() == (uint32_t)dim && element_strides.size() == (uint32_t)dim); - - auto call_cuTensorMapEncodeTiled = [&](Args... args) { +// Driver symbols are independent of tensor storage and the current device. +// Inline linkage shares this cache across translation units. A failed lookup +// throws before initialization completes, so a later call retries it. +inline PFN_cuTensorMapEncodeTiled_v12000 get_tensor_map_encoder() { + static const auto encoder = [] { cudaDriverEntryPointQueryResult cuda_status; void* pfn = nullptr; -#if (__CUDACC_VER_MAJOR__ > 12) +#if CUDA_VERSION >= 13000 KU_CUDA_CHECK(cudaGetDriverEntryPointByVersion( - "cuTensorMapEncodeTiled", - &pfn, 12000, - cudaEnableDefault, - &cuda_status)); + "cuTensorMapEncodeTiled", &pfn, 12000, cudaEnableDefault, &cuda_status)); #else KU_CUDA_CHECK(cudaGetDriverEntryPoint( - "cuTensorMapEncodeTiled", - &pfn, - cudaEnableDefault, - &cuda_status)); + "cuTensorMapEncodeTiled", &pfn, cudaEnableDefault, &cuda_status)); #endif - if (cuda_status != cudaDriverEntryPointSuccess) { - KU_ASSERT(false, "Failed to load `cuTensorMapEncodeTiled`. cuda_status = %d", cuda_status); - } - return reinterpret_cast(pfn)(args...); \ - }; + KU_ASSERT(cuda_status == cudaDriverEntryPointSuccess && pfn != nullptr, + "Failed to load `cuTensorMapEncodeTiled`. cuda_status = %d", cuda_status); + return reinterpret_cast(pfn); + }(); + return encoder; +} + +namespace detail { +inline CUtensorMap encode_tensor_map( + int dim, + const uint64_t* size, + const uint64_t* strides, + const uint32_t* box_size, + void* global_ptr, + CUtensorMapDataType data_type, + CUtensorMapSwizzle swizzle_mode, + CUtensorMapL2promotion l2_promotion, + CUtensorMapInterleave interleave_mode, + CUtensorMapFloatOOBfill oob_fill, + const uint32_t* element_strides +) { CUtensorMap result; - CUresult ret_code = call_cuTensorMapEncodeTiled( + CUresult ret_code = get_tensor_map_encoder()( &result, data_type, dim, global_ptr, - size.data(), - strides.data(), - box_size.data(), - element_strides.data(), + size, + strides, + box_size, + element_strides, interleave_mode, swizzle_mode, l2_promotion, oob_fill ); if (ret_code != CUresult::CUDA_SUCCESS) { - auto print_vector = [&](auto t, const char* fmt, const char end='\n') { - for (auto elem : t) { - printf(fmt, elem); + auto print_vector = [&](auto t, int count, const char* fmt, const char end='\n') { + for (int i = 0; i < count; ++i) { + printf(fmt, t[i]); } printf("%c", end); }; fprintf(stderr, "Failed to create tensormap\n"); fprintf(stderr, "Dim: %d\n", dim); - printf("size: "); print_vector(size, "%lu "); - printf("strides: "); print_vector(strides, "%lu "); - printf("box_size: "); print_vector(box_size, "%u "); - printf("element_strides: "); print_vector(element_strides, "%u "); + printf("size: "); print_vector(size, dim, "%lu "); + printf("strides: "); print_vector(strides, dim-1, "%lu "); + printf("box_size: "); print_vector(box_size, dim, "%u "); + printf("element_strides: "); print_vector(element_strides, dim, "%u "); printf("global ptr: 0x%lx\n", (int64_t)global_ptr); printf("data_type: %d\n", (int)data_type); printf("swizzle_mode: %d\n", (int)swizzle_mode); @@ -183,6 +173,54 @@ static inline CUtensorMap make_tensor_map( return result; } +} // namespace detail + +// Preserve the dynamically ranked interface for existing callers. +static inline CUtensorMap make_tensor_map( + const std::vector& size, + const std::vector& strides, // In bytes + const std::vector& box_size, + void* global_ptr, + CUtensorMapDataType data_type, + CUtensorMapSwizzle swizzle_mode, + CUtensorMapL2promotion l2_promotion, + CUtensorMapInterleave interleave_mode = CU_TENSOR_MAP_INTERLEAVE_NONE, + CUtensorMapFloatOOBfill oob_fill = CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE, + const std::vector& element_strides_ = {} +) { + int dim = size.size(); + KU_ASSERT(dim >= 1); + std::vector element_strides = element_strides_.empty() + ? std::vector(dim, 1) : element_strides_; + KU_ASSERT(strides.size() == (uint32_t)dim-1 && box_size.size() == (uint32_t)dim && element_strides.size() == (uint32_t)dim); + return detail::encode_tensor_map(dim, size.data(), strides.data(), box_size.data(), + global_ptr, data_type, swizzle_mode, l2_promotion, interleave_mode, oob_fill, element_strides.data()); +} + +// Fixed rank keeps metadata on the stack; the descriptor is still rebuilt for +// every call's pointer, dimensions, and strides. +template +inline CUtensorMap make_tensor_map( + const std::array& size, + const std::array& strides, // In bytes + const std::array& box_size, + void* global_ptr, + CUtensorMapDataType data_type, + CUtensorMapSwizzle swizzle_mode, + CUtensorMapL2promotion l2_promotion, + CUtensorMapInterleave interleave_mode = CU_TENSOR_MAP_INTERLEAVE_NONE, + CUtensorMapFloatOOBfill oob_fill = CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE, + const std::array& element_strides = [] { + std::array ones; + ones.fill(1); + return ones; + }() +) { + static_assert(Rank >= 1 && Rank <= 5); + return detail::encode_tensor_map(Rank, size.data(), strides.data(), box_size.data(), + global_ptr, data_type, swizzle_mode, l2_promotion, interleave_mode, oob_fill, element_strides.data()); +} + // Given strides (in number of elements), this function converts their datatype in uint64_t and then multiplies by elem_size template static inline std::vector make_stride_helper(const std::vector &strides_in_elems, size_t elem_size) { @@ -279,4 +317,4 @@ void launch_kernel(const KernelLaunchConfig &cfg, KernelFunc kernel, Args&&... a #endif // KERUTILS_IS_BUILD_ON_CUDA -} // namespace kerutils \ No newline at end of file +} // namespace kerutils diff --git a/csrc/cuda_kernels/common_parts.cuh b/csrc/cuda_kernels/common_parts.cuh index aa347eb..77a85e3 100644 --- a/csrc/cuda_kernels/common_parts.cuh +++ b/csrc/cuda_kernels/common_parts.cuh @@ -411,7 +411,7 @@ public: constexpr CUtensorMapDataType dtype = std::is_same_v ? CU_TENSOR_MAP_DATA_TYPE_BFLOAT16 : CU_TENSOR_MAP_DATA_TYPE_FLOAT32; - return ku::make_tensor_map( + return ku::make_tensor_map<3>( // Split args.vocab_size into two dims because // - TMA requires innermost box dim <= swizzle size // - To avoid OOB as we have `INPUT_STRIDE_ALIGNMENT_REQUIREMENT` @@ -420,7 +420,7 @@ public: ku::ceil_div((uint64_t)args.vocab_size, (uint64_t)NUM_ELEMS_PER_TMA_ROW), args.batch_size }, - ku::make_stride_helper({NUM_ELEMS_PER_TMA_ROW, args.stride_input_batch}, sizeof(ValueT)), + {uint64_t(NUM_ELEMS_PER_TMA_ROW) * sizeof(ValueT), uint64_t(args.stride_input_batch) * sizeof(ValueT)}, { NUM_ELEMS_PER_TMA_ROW, NUM_TMA_ROWS_PER_SEG, diff --git a/csrc/cuda_kernels/v3/topk_select.cuh b/csrc/cuda_kernels/v3/topk_select.cuh index fd30af5..d0b2d6b 100644 --- a/csrc/cuda_kernels/v3/topk_select.cuh +++ b/csrc/cuda_kernels/v3/topk_select.cuh @@ -130,7 +130,6 @@ void run_topk_select_kernel(const TopkSelectArgs &args) { auto kernel = topk_kernel; constexpr size_t smem_size = sizeof(typename Kernel::SharedMemoryPlan); KU_ASSERT(smem_size * Kernel::TARGET_OCCUPANCY <= args.shared_memory_size_per_sm); - KU_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size)); KU_ASSERT(args.stride_input_batch % 8 == 0, "stride_input_batch must be 16B-aligned"); typename Kernel::TmaParams tma_params = {Kernel::make_topk_tensor_map(args)}; diff --git a/csrc/cuda_kernels/v3_cluster/topk_select.cuh b/csrc/cuda_kernels/v3_cluster/topk_select.cuh index 3d880f8..fa82726 100644 --- a/csrc/cuda_kernels/v3_cluster/topk_select.cuh +++ b/csrc/cuda_kernels/v3_cluster/topk_select.cuh @@ -286,10 +286,6 @@ void run_topk_select_kernel(const TopkSelectArgs &args) { auto kernel = topk_kernel; constexpr size_t smem_size = sizeof(typename Kernel::SharedMemoryPlanBF16Cluster); KU_ASSERT(smem_size * Kernel::TARGET_OCCUPANCY <= args.shared_memory_size_per_sm); - KU_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size)); - if constexpr (Kernel::CLUSTER_SIZE > 8) { - KU_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeNonPortableClusterSizeAllowed, 1)); - } KU_ASSERT(args.stride_input_batch % 8 == 0, "stride_input_batch must be 16B-aligned"); typename Kernel::TmaParams tma_params = {Kernel::make_topk_tensor_map(args)}; diff --git a/csrc/cuda_kernels/v3_fp32/topk_select.cuh b/csrc/cuda_kernels/v3_fp32/topk_select.cuh index 897b9e1..9345175 100644 --- a/csrc/cuda_kernels/v3_fp32/topk_select.cuh +++ b/csrc/cuda_kernels/v3_fp32/topk_select.cuh @@ -589,7 +589,6 @@ void run_topk_select_kernel(const TopkSelectArgs &args) { auto kernel = topk_kernel; constexpr size_t smem_size = sizeof(typename Kernel::SharedMemoryPlanFP32); KU_ASSERT(smem_size * Kernel::TARGET_OCCUPANCY <= args.shared_memory_size_per_sm); - KU_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size)); KU_ASSERT(args.stride_input_batch % 4 == 0, "stride_input_batch must be 16B-aligned"); typename Kernel::TmaParams tma_params = {Kernel::make_topk_tensor_map(args)}; diff --git a/tests/host/README.md b/tests/host/README.md new file mode 100644 index 0000000..ee9ef8a --- /dev/null +++ b/tests/host/README.md @@ -0,0 +1,17 @@ +# Host setup regression tests + +These tests use real CUDA headers and a mocked driver resolver and encoder. They require CUDA 13's compiler, but do not initialize CUDA or require a GPU. The mock records encoder arguments rather than validating device-specific tensor-map constraints. + +From the repository root, build into a directory outside the checkout: + +```sh +nvcc -std=c++20 -O2 --cudart shared -I csrc/3rdparty/kerutils/include -Xcompiler -pthread tests/host/tensor_map.cu tests/host/tensor_map_peer.cu tests/host/allocations.cpp -o /path/to/artifacts/tensor-map-test +/path/to/artifacts/tensor-map-test +/path/to/artifacts/tensor-map-test retry +``` + +The normal process tests concurrent cold initialization, one lookup shared across two translation units, current per-call tensor arguments, retry after an encoder failure, and vector/array equivalence for ranks one through five with default and explicit element strides. The separate `retry` process additionally tests runtime lookup failure, missing symbol, and a null function pointer before successful initialization. Assertion diagnostics from these intentional failures are expected. + +For the baseline regression, compile the same sources with `-DTEST_BASELINE` and point the include path at baseline kerutils headers. Run without `retry`: the single-lookup assertion must fail after sixteen concurrent calls. This mode excludes the new array overload and does not attempt the baseline's unsafe null-pointer call. + +Run either binary with `benchmark` to measure rank-three metadata preparation through its actual headers. It warms up 1,000 calls, then reports seven rounds of 200,000 calls, allocations, and resolver calls. The baseline uses the original vector/stride-helper path; the candidate uses fixed arrays. Both use the same recording mock encoder. This isolates host preparation and does not measure driver encoding, kernel launch, or end-to-end Top-K latency. Alternate baseline/candidate runs to check timing noise. The candidate additionally asserts zero steady-state allocations and lookups. diff --git a/tests/host/allocations.cpp b/tests/host/allocations.cpp new file mode 100644 index 0000000..62326a0 --- /dev/null +++ b/tests/host/allocations.cpp @@ -0,0 +1,13 @@ +#include +#include + +thread_local bool count_allocations = false; +thread_local size_t allocations = 0; + +void* operator new(size_t size) { + if (count_allocations) ++allocations; + if (auto* ptr = std::malloc(size ? size : 1)) return ptr; + throw std::bad_alloc(); +} +void operator delete(void* ptr) noexcept { std::free(ptr); } +void operator delete(void* ptr, size_t) noexcept { std::free(ptr); } diff --git a/tests/host/tensor_map.cu b/tests/host/tensor_map.cu new file mode 100644 index 0000000..ff715ae --- /dev/null +++ b/tests/host/tensor_map.cu @@ -0,0 +1,176 @@ +// Real CUDA types and ABI, mocked driver lookup/encoder: no GPU is required. +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "kerutils/host/host.h" + +extern thread_local bool count_allocations; +extern thread_local size_t allocations; + +void encode_from_peer(); +static std::atomic lookups{0}, encodes{0}; +static int lookup_failure = 0; +static bool encode_failure = false; +struct Arguments { + unsigned rank; + void* pointer; + CUtensorMapDataType dtype; + CUtensorMapInterleave interleave; + CUtensorMapSwizzle swizzle; + CUtensorMapL2promotion promotion; + CUtensorMapFloatOOBfill fill; + std::array sizes{}, strides{}; + std::array boxes{}, elements{}; + bool operator==(const Arguments&) const = default; +}; +static thread_local Arguments recorded; + +static CUresult CUDAAPI encode(CUtensorMap* result, CUtensorMapDataType dtype, + cuuint32_t rank, void* pointer, const cuuint64_t* sizes, + const cuuint64_t* strides, const cuuint32_t* boxes, const cuuint32_t* elements, + CUtensorMapInterleave interleave, CUtensorMapSwizzle swizzle, + CUtensorMapL2promotion promotion, CUtensorMapFloatOOBfill fill) { + ++encodes; + assert(rank >= 1 && rank <= 5); + recorded = {}; + recorded.rank = rank; recorded.pointer = pointer; recorded.dtype = dtype; + recorded.interleave = interleave; recorded.swizzle = swizzle; + recorded.promotion = promotion; recorded.fill = fill; + std::copy_n(sizes, rank, recorded.sizes.begin()); + std::copy_n(strides, rank - 1, recorded.strides.begin()); + std::copy_n(boxes, rank, recorded.boxes.begin()); + std::copy_n(elements, rank, recorded.elements.begin()); + std::memset(result, 0, sizeof(*result)); + return encode_failure ? CUDA_ERROR_INVALID_VALUE : CUDA_SUCCESS; +} + +extern "C" cudaError_t CUDARTAPI cudaGetDriverEntryPointByVersion( + const char* name, void** pointer, unsigned version, unsigned long long flags, + cudaDriverEntryPointQueryResult* status) { + ++lookups; + assert(std::strcmp(name, "cuTensorMapEncodeTiled") == 0); + assert(version == 12000 && flags == cudaEnableDefault); + *status = lookup_failure == 2 ? cudaDriverEntryPointSymbolNotFound : cudaDriverEntryPointSuccess; + *pointer = lookup_failure == 3 ? nullptr : reinterpret_cast(&encode); + return lookup_failure == 1 ? cudaErrorUnknown : cudaSuccess; +} + +template void throws(F f) { + bool caught = false; + try { f(); } catch (const ku::KUException&) { caught = true; } + assert(caught); +} + +#ifndef TEST_BASELINE +template void compare_storage() { + std::array sizes; + std::array strides; + std::array boxes, elements; + sizes.fill(32); strides.fill(128); boxes.fill(1); elements.fill(2); + for (bool explicit_elements : {false, true}) { + const auto ptr = reinterpret_cast(0x2000 + Rank * 128); + const auto dtype = CU_TENSOR_MAP_DATA_TYPE_FLOAT32; + const auto swizzle = CU_TENSOR_MAP_SWIZZLE_NONE; + const auto promotion = CU_TENSOR_MAP_L2_PROMOTION_L2_256B; + const auto interleave = CU_TENSOR_MAP_INTERLEAVE_NONE; + const auto fill = CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA; + std::vector e; + if (explicit_elements) e.assign(elements.begin(), elements.end()); + ku::make_tensor_map({sizes.begin(), sizes.end()}, {strides.begin(), strides.end()}, + {boxes.begin(), boxes.end()}, ptr, dtype, swizzle, promotion, interleave, fill, e); + auto expected = recorded; + if (explicit_elements) + ku::make_tensor_map(sizes, strides, boxes, ptr, dtype, swizzle, promotion, interleave, fill, elements); + else + ku::make_tensor_map(sizes, strides, boxes, ptr, dtype, swizzle, promotion, interleave, fill); + assert(recorded == expected); + } +} +#endif + +// Measures actual metadata wrappers with a mock encoder, not CUDA execution. +void benchmark() { + constexpr int count = 200000; + auto prepare = [](int i) { + const uint64_t vocab = 8192 + (i & 31); +#ifdef TEST_BASELINE + ku::make_tensor_map({32, (vocab + 31) / 32, 6}, + ku::make_stride_helper({32, 8448}, 4), {32, 8, 1}, +#else + ku::make_tensor_map<3>({32, (vocab + 31) / 32, 6}, + {128, 33792}, {32, 8, 1}, +#endif + reinterpret_cast(0x1000), CU_TENSOR_MAP_DATA_TYPE_FLOAT32, + CU_TENSOR_MAP_SWIZZLE_128B, CU_TENSOR_MAP_L2_PROMOTION_L2_256B); + assert(recorded.sizes[1] == (vocab + 31) / 32); + }; + for (int i = 0; i < 1000; ++i) prepare(i); + for (int round = 0; round < 7; ++round) { + allocations = 0; + const int before = lookups; + count_allocations = true; + const auto start = std::chrono::steady_clock::now(); + for (int i = 0; i < count; ++i) prepare(i); + const auto elapsed = std::chrono::steady_clock::now() - start; + count_allocations = false; + std::printf("round=%d calls=%d allocations=%zu lookups=%d ns_per_call=%.3f\n", + round, count, allocations, int(lookups) - before, + std::chrono::duration(elapsed).count() / count); +#ifndef TEST_BASELINE + assert(allocations == 0 && int(lookups) == before); +#endif + } +} + +int main(int argc, char** argv) { + if (argc > 1 && std::strcmp(argv[1], "benchmark") == 0) { + benchmark(); + return 0; + } + if (argc > 1 && std::strcmp(argv[1], "retry") == 0) { + for (int failure : {1, 2, 3}) { + lookup_failure = failure; + throws(encode_from_peer); + assert(lookups == failure && encodes == 0); + } + lookup_failure = 0; + } + const int failed_lookups = lookups; + std::atomic ready{0}; std::atomic start{false}; + std::vector threads; + for (int i = 0; i < 16; ++i) threads.emplace_back([&] { + ++ready; while (!start.load()) std::this_thread::yield(); encode_from_peer(); + }); + while (ready != 16) std::this_thread::yield(); + start = true; + for (auto& t : threads) t.join(); + assert(encodes == 16); + assert(lookups == failed_lookups + 1); + ku::make_tensor_map({64, 8, 3}, {256, 2048}, {64, 2, 1}, + reinterpret_cast(0x3000), CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, + CU_TENSOR_MAP_SWIZZLE_128B, CU_TENSOR_MAP_L2_PROMOTION_NONE); + assert(recorded.pointer == reinterpret_cast(0x3000)); + assert(recorded.sizes[0] == 64 && recorded.sizes[1] == 8 && recorded.sizes[2] == 3); + assert(recorded.strides[0] == 256 && recorded.strides[1] == 2048); + assert(recorded.boxes[1] == 2 && recorded.dtype == CU_TENSOR_MAP_DATA_TYPE_BFLOAT16); + assert(lookups == failed_lookups + 1); // Shared with the other translation unit. + encode_failure = true; + throws(encode_from_peer); + encode_failure = false; + encode_from_peer(); + assert(encodes == 19 && lookups == failed_lookups + 1); +#ifndef TEST_BASELINE + compare_storage<1>(); compare_storage<2>(); compare_storage<3>(); + compare_storage<4>(); compare_storage<5>(); +#endif + std::puts("tensor-map host tests passed"); +} diff --git a/tests/host/tensor_map_peer.cu b/tests/host/tensor_map_peer.cu new file mode 100644 index 0000000..26d8be7 --- /dev/null +++ b/tests/host/tensor_map_peer.cu @@ -0,0 +1,9 @@ +#include +#include +#include "kerutils/host/host.h" + +void encode_from_peer() { + ku::make_tensor_map({32, 4, 2}, {128, 512}, {32, 1, 1}, + reinterpret_cast(0x1000), CU_TENSOR_MAP_DATA_TYPE_FLOAT32, + CU_TENSOR_MAP_SWIZZLE_128B, CU_TENSOR_MAP_L2_PROMOTION_NONE); +}