From 277689f8b8f68852740f11649646b8558f11ade3 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sat, 30 May 2026 08:34:19 +0000 Subject: [PATCH] P4: use proper CUDA enum names --- tests/unit/test_p4_tma_descriptor_dump.cu | 117 +++++++--------------- 1 file changed, 36 insertions(+), 81 deletions(-) diff --git a/tests/unit/test_p4_tma_descriptor_dump.cu b/tests/unit/test_p4_tma_descriptor_dump.cu index da4b3377..ce89ee4b 100644 --- a/tests/unit/test_p4_tma_descriptor_dump.cu +++ b/tests/unit/test_p4_tma_descriptor_dump.cu @@ -1,6 +1,5 @@ /** * P4: Dump TMA descriptor bytes for comparison. - * Uses CUDA Driver API directly. */ #include #include @@ -8,14 +7,6 @@ #include #include -/* CUDA 13.2 enum values for cuTensorMapEncodeTiled: - * CUtensorMapDataType: 0=F16, 6=BF16 - * CUtensorMapInterleave: 0=NONE, 1=16B, 2=32B - * CUtensorMapSwizzle: 0=NONE, 1=4B, 2=32B, 3=64B, 4=128B - * CUtensorMapL2Promotion: 0=NONE, 1=64B, 2=128B, 3=256B - * CUtensorMapOOBFill: 0=NONE, 1=ZERO - */ - int main() { const int ROWS = 128; const int COLS = 16; @@ -33,110 +24,74 @@ int main() { CUtensorMap tma_desc; CUresult res; - // Descriptor 1: NO swizzle - res = cuTensorMapEncodeTiled( - &tma_desc, - (CUtensorMapDataType)6, // BF16 - 2, // rank - d_ptr, - tensorDims, - globalStrides, - boxDims, - elementStrides, - (CUtensorMapInterleave)0, // NONE - (CUtensorMapSwizzle)0, // NONE - 0, // NONE - (CUtensorMapOOBFill)0 // NONE - ); - + // Config 1: NO swizzle printf("=== Descriptor 1: NO swizzle ===\n"); + res = cuTensorMapEncodeTiled( + &tma_desc, 2, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, + d_ptr, tensorDims, globalStrides, boxDims, elementStrides, + CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_NONE, + CU_TENSOR_MAP_L2_PROMOTION_NONE, CU_TENSOR_MAP_OOB_FILL_NONE + ); if (res != CUDA_SUCCESS) { printf("FAILED: %d\n", res); } else { - auto* bytes = reinterpret_cast(&tma_desc); + auto* b = reinterpret_cast(&tma_desc); for (int i = 0; i < 128; i += 16) { printf("[%3d-%3d]: ", i, i+15); - for (int j = 0; j < 16; j++) printf("%02x ", bytes[i+j]); + for (int j = 0; j < 16; j++) printf("%02x ", b[i+j]); printf("\n"); } } - // Descriptor 2: SWIZZLE_128B - res = cuTensorMapEncodeTiled( - &tma_desc, - (CUtensorMapDataType)6, - 2, - d_ptr, - tensorDims, - globalStrides, - boxDims, - elementStrides, - (CUtensorMapInterleave)0, - (CUtensorMapSwizzle)4, // 128B - 0, - (CUtensorMapOOBFill)0 - ); - + // Config 2: SWIZZLE_128B printf("\n=== Descriptor 2: SWIZZLE_128B ===\n"); + res = cuTensorMapEncodeTiled( + &tma_desc, 2, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, + d_ptr, tensorDims, globalStrides, boxDims, elementStrides, + CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B, + CU_TENSOR_MAP_L2_PROMOTION_NONE, CU_TENSOR_MAP_OOB_FILL_NONE + ); if (res != CUDA_SUCCESS) { printf("FAILED: %d\n", res); } else { - auto* bytes = reinterpret_cast(&tma_desc); + auto* b = reinterpret_cast(&tma_desc); for (int i = 0; i < 128; i += 16) { printf("[%3d-%3d]: ", i, i+15); - for (int j = 0; j < 16; j++) printf("%02x ", bytes[i+j]); + for (int j = 0; j < 16; j++) printf("%02x ", b[i+j]); printf("\n"); } } - // Descriptor 3: NO swizzle, OOB_FILL_ZERO - res = cuTensorMapEncodeTiled( - &tma_desc, - (CUtensorMapDataType)6, - 2, - d_ptr, - tensorDims, - globalStrides, - boxDims, - elementStrides, - (CUtensorMapInterleave)0, - (CUtensorMapSwizzle)0, - 0, - (CUtensorMapOOBFill)1 // ZERO - ); - + // Config 3: NO swizzle, OOB_FILL_ZERO printf("\n=== Descriptor 3: NO swizzle, OOB_FILL_ZERO ===\n"); + res = cuTensorMapEncodeTiled( + &tma_desc, 2, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, + d_ptr, tensorDims, globalStrides, boxDims, elementStrides, + CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_NONE, + CU_TENSOR_MAP_L2_PROMOTION_NONE, CU_TENSOR_MAP_OOB_FILL_ZERO + ); if (res != CUDA_SUCCESS) { printf("FAILED: %d\n", res); } else { - auto* bytes = reinterpret_cast(&tma_desc); + auto* b = reinterpret_cast(&tma_desc); for (int i = 0; i < 128; i += 16) { printf("[%3d-%3d]: ", i, i+15); - for (int j = 0; j < 16; j++) printf("%02x ", bytes[i+j]); + for (int j = 0; j < 16; j++) printf("%02x ", b[i+j]); printf("\n"); } } - // Descriptor 4: SWIZZLE_128B, OOB_FILL_ZERO - res = cuTensorMapEncodeTiled( - &tma_desc, - (CUtensorMapDataType)6, - 2, - d_ptr, - tensorDims, - globalStrides, - boxDims, - elementStrides, - (CUtensorMapInterleave)0, - (CUtensorMapSwizzle)4, - 0, - (CUtensorMapOOBFill)1 - ); - + // Config 4: SWIZZLE_128B, OOB_FILL_ZERO printf("\n=== Descriptor 4: SWIZZLE_128B, OOB_FILL_ZERO ===\n"); + res = cuTensorMapEncodeTiled( + &tma_desc, 2, CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, + d_ptr, tensorDims, globalStrides, boxDims, elementStrides, + CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B, + CU_TENSOR_MAP_L2_PROMOTION_NONE, CU_TENSOR_MAP_OOB_FILL_ZERO + ); if (res != CUDA_SUCCESS) { printf("FAILED: %d\n", res); } else { - auto* bytes = reinterpret_cast(&tma_desc); + auto* b = reinterpret_cast(&tma_desc); for (int i = 0; i < 128; i += 16) { printf("[%3d-%3d]: ", i, i+15); - for (int j = 0; j < 16; j++) printf("%02x ", bytes[i+j]); + for (int j = 0; j < 16; j++) printf("%02x ", b[i+j]); printf("\n"); } }