From d855d0461b3b9590d6030635f3fb2273b48735f9 Mon Sep 17 00:00:00 2001 From: castano Date: Tue, 6 May 2008 19:52:27 +0000 Subject: [PATCH] Add single color checks to CUDA compressors. Use optimized bitmap table for CTX compressor. --- src/nvtt/cuda/Bitmaps.h | 767 ++++++++++++++++++++++++++++++ src/nvtt/cuda/CompressKernel.cu | 74 ++- src/nvtt/cuda/CudaCompressDXT.cpp | 10 +- src/nvtt/cuda/CudaCompressDXT.h | 1 + src/nvtt/cuda/CudaMath.h | 5 + src/nvtt/cuda/CudaUtils.cpp | 2 +- src/nvtt/tests/stress.cpp | 24 +- 7 files changed, 863 insertions(+), 20 deletions(-) diff --git a/src/nvtt/cuda/Bitmaps.h b/src/nvtt/cuda/Bitmaps.h index 1494092..d05e9f3 100644 --- a/src/nvtt/cuda/Bitmaps.h +++ b/src/nvtt/cuda/Bitmaps.h @@ -1117,3 +1117,770 @@ const static uint s_bitmapTable[992] = 0x55555557, 0x55555557, }; + + +/* +void precomp() +{ + unsigned int bitmaps[1024]; + + int num = 0; + + printf("const static uint s_bitmapTableCTX[704] =\n{\n"); + + for (int a = 1; a <= 15; a++) + { + for (int b = a; b <= 15; b++) + { + for (int c = b; c <= 15; c++) + { + int indices[16]; + + int i = 0; + for(; i < a; i++) { + indices[i] = 0; + } + for(; i < a+b; i++) { + indices[i] = 2; + } + for(; i < a+b+c; i++) { + indices[i] = 3; + } + for(; i < 16; i++) { + indices[i] = 1; + } + + unsigned int bm = 0; + for(i = 0; i < 16; i++) { + bm |= indices[i] << (i * 2); + } + + printf("\t0x%8X, // %d %d %d %d\n", bm, a-0, b-a, c-b, 16-c); + + bitmaps[num] = bm; + num++; + } + } + } + + // Align to 32: 680 -> 704 + while (num < 704) + { + printf("\t0x80000000,\n"); + + bitmaps[num] = 0x80000000; // 15 0 0 1; + num++; + } + + printf("}; // num = %d\n", num); +} +*/ + +const static uint s_bitmapTableCTX[704] = +{ + 0x55555578, // 1 0 0 15 + 0x555555F8, // 1 0 1 14 + 0x555557F8, // 1 0 2 13 + 0x55555FF8, // 1 0 3 12 + 0x55557FF8, // 1 0 4 11 + 0x5555FFF8, // 1 0 5 10 + 0x5557FFF8, // 1 0 6 9 + 0x555FFFF8, // 1 0 7 8 + 0x557FFFF8, // 1 0 8 7 + 0x55FFFFF8, // 1 0 9 6 + 0x57FFFFF8, // 1 0 10 5 + 0x5FFFFFF8, // 1 0 11 4 + 0x7FFFFFF8, // 1 0 12 3 + 0xFFFFFFF8, // 1 0 13 2 + 0xFFFFFFF8, // 1 0 14 1 + 0x555557E8, // 1 1 0 14 + 0x55555FE8, // 1 1 1 13 + 0x55557FE8, // 1 1 2 12 + 0x5555FFE8, // 1 1 3 11 + 0x5557FFE8, // 1 1 4 10 + 0x555FFFE8, // 1 1 5 9 + 0x557FFFE8, // 1 1 6 8 + 0x55FFFFE8, // 1 1 7 7 + 0x57FFFFE8, // 1 1 8 6 + 0x5FFFFFE8, // 1 1 9 5 + 0x7FFFFFE8, // 1 1 10 4 + 0xFFFFFFE8, // 1 1 11 3 + 0xFFFFFFE8, // 1 1 12 2 + 0xFFFFFFE8, // 1 1 13 1 + 0x55557FA8, // 1 2 0 13 + 0x5555FFA8, // 1 2 1 12 + 0x5557FFA8, // 1 2 2 11 + 0x555FFFA8, // 1 2 3 10 + 0x557FFFA8, // 1 2 4 9 + 0x55FFFFA8, // 1 2 5 8 + 0x57FFFFA8, // 1 2 6 7 + 0x5FFFFFA8, // 1 2 7 6 + 0x7FFFFFA8, // 1 2 8 5 + 0xFFFFFFA8, // 1 2 9 4 + 0xFFFFFFA8, // 1 2 10 3 + 0xFFFFFFA8, // 1 2 11 2 + 0xFFFFFFA8, // 1 2 12 1 + 0x5557FEA8, // 1 3 0 12 + 0x555FFEA8, // 1 3 1 11 + 0x557FFEA8, // 1 3 2 10 + 0x55FFFEA8, // 1 3 3 9 + 0x57FFFEA8, // 1 3 4 8 + 0x5FFFFEA8, // 1 3 5 7 + 0x7FFFFEA8, // 1 3 6 6 + 0xFFFFFEA8, // 1 3 7 5 + 0xFFFFFEA8, // 1 3 8 4 + 0xFFFFFEA8, // 1 3 9 3 + 0xFFFFFEA8, // 1 3 10 2 + 0xFFFFFEA8, // 1 3 11 1 + 0x557FFAA8, // 1 4 0 11 + 0x55FFFAA8, // 1 4 1 10 + 0x57FFFAA8, // 1 4 2 9 + 0x5FFFFAA8, // 1 4 3 8 + 0x7FFFFAA8, // 1 4 4 7 + 0xFFFFFAA8, // 1 4 5 6 + 0xFFFFFAA8, // 1 4 6 5 + 0xFFFFFAA8, // 1 4 7 4 + 0xFFFFFAA8, // 1 4 8 3 + 0xFFFFFAA8, // 1 4 9 2 + 0xFFFFFAA8, // 1 4 10 1 + 0x57FFEAA8, // 1 5 0 10 + 0x5FFFEAA8, // 1 5 1 9 + 0x7FFFEAA8, // 1 5 2 8 + 0xFFFFEAA8, // 1 5 3 7 + 0xFFFFEAA8, // 1 5 4 6 + 0xFFFFEAA8, // 1 5 5 5 + 0xFFFFEAA8, // 1 5 6 4 + 0xFFFFEAA8, // 1 5 7 3 + 0xFFFFEAA8, // 1 5 8 2 + 0xFFFFEAA8, // 1 5 9 1 + 0x7FFFAAA8, // 1 6 0 9 + 0xFFFFAAA8, // 1 6 1 8 + 0xFFFFAAA8, // 1 6 2 7 + 0xFFFFAAA8, // 1 6 3 6 + 0xFFFFAAA8, // 1 6 4 5 + 0xFFFFAAA8, // 1 6 5 4 + 0xFFFFAAA8, // 1 6 6 3 + 0xFFFFAAA8, // 1 6 7 2 + 0xFFFFAAA8, // 1 6 8 1 + 0xFFFEAAA8, // 1 7 0 8 + 0xFFFEAAA8, // 1 7 1 7 + 0xFFFEAAA8, // 1 7 2 6 + 0xFFFEAAA8, // 1 7 3 5 + 0xFFFEAAA8, // 1 7 4 4 + 0xFFFEAAA8, // 1 7 5 3 + 0xFFFEAAA8, // 1 7 6 2 + 0xFFFEAAA8, // 1 7 7 1 + 0xFFFAAAA8, // 1 8 0 7 + 0xFFFAAAA8, // 1 8 1 6 + 0xFFFAAAA8, // 1 8 2 5 + 0xFFFAAAA8, // 1 8 3 4 + 0xFFFAAAA8, // 1 8 4 3 + 0xFFFAAAA8, // 1 8 5 2 + 0xFFFAAAA8, // 1 8 6 1 + 0xFFEAAAA8, // 1 9 0 6 + 0xFFEAAAA8, // 1 9 1 5 + 0xFFEAAAA8, // 1 9 2 4 + 0xFFEAAAA8, // 1 9 3 3 + 0xFFEAAAA8, // 1 9 4 2 + 0xFFEAAAA8, // 1 9 5 1 + 0xFFAAAAA8, // 1 10 0 5 + 0xFFAAAAA8, // 1 10 1 4 + 0xFFAAAAA8, // 1 10 2 3 + 0xFFAAAAA8, // 1 10 3 2 + 0xFFAAAAA8, // 1 10 4 1 + 0xFEAAAAA8, // 1 11 0 4 + 0xFEAAAAA8, // 1 11 1 3 + 0xFEAAAAA8, // 1 11 2 2 + 0xFEAAAAA8, // 1 11 3 1 + 0xFAAAAAA8, // 1 12 0 3 + 0xFAAAAAA8, // 1 12 1 2 + 0xFAAAAAA8, // 1 12 2 1 + 0xEAAAAAA8, // 1 13 0 2 + 0xEAAAAAA8, // 1 13 1 1 + 0xAAAAAAA8, // 1 14 0 1 + 0x55555FA0, // 2 0 0 14 + 0x55557FA0, // 2 0 1 13 + 0x5555FFA0, // 2 0 2 12 + 0x5557FFA0, // 2 0 3 11 + 0x555FFFA0, // 2 0 4 10 + 0x557FFFA0, // 2 0 5 9 + 0x55FFFFA0, // 2 0 6 8 + 0x57FFFFA0, // 2 0 7 7 + 0x5FFFFFA0, // 2 0 8 6 + 0x7FFFFFA0, // 2 0 9 5 + 0xFFFFFFA0, // 2 0 10 4 + 0xFFFFFFA0, // 2 0 11 3 + 0xFFFFFFA0, // 2 0 12 2 + 0xFFFFFFA0, // 2 0 13 1 + 0x5555FEA0, // 2 1 0 13 + 0x5557FEA0, // 2 1 1 12 + 0x555FFEA0, // 2 1 2 11 + 0x557FFEA0, // 2 1 3 10 + 0x55FFFEA0, // 2 1 4 9 + 0x57FFFEA0, // 2 1 5 8 + 0x5FFFFEA0, // 2 1 6 7 + 0x7FFFFEA0, // 2 1 7 6 + 0xFFFFFEA0, // 2 1 8 5 + 0xFFFFFEA0, // 2 1 9 4 + 0xFFFFFEA0, // 2 1 10 3 + 0xFFFFFEA0, // 2 1 11 2 + 0xFFFFFEA0, // 2 1 12 1 + 0x555FFAA0, // 2 2 0 12 + 0x557FFAA0, // 2 2 1 11 + 0x55FFFAA0, // 2 2 2 10 + 0x57FFFAA0, // 2 2 3 9 + 0x5FFFFAA0, // 2 2 4 8 + 0x7FFFFAA0, // 2 2 5 7 + 0xFFFFFAA0, // 2 2 6 6 + 0xFFFFFAA0, // 2 2 7 5 + 0xFFFFFAA0, // 2 2 8 4 + 0xFFFFFAA0, // 2 2 9 3 + 0xFFFFFAA0, // 2 2 10 2 + 0xFFFFFAA0, // 2 2 11 1 + 0x55FFEAA0, // 2 3 0 11 + 0x57FFEAA0, // 2 3 1 10 + 0x5FFFEAA0, // 2 3 2 9 + 0x7FFFEAA0, // 2 3 3 8 + 0xFFFFEAA0, // 2 3 4 7 + 0xFFFFEAA0, // 2 3 5 6 + 0xFFFFEAA0, // 2 3 6 5 + 0xFFFFEAA0, // 2 3 7 4 + 0xFFFFEAA0, // 2 3 8 3 + 0xFFFFEAA0, // 2 3 9 2 + 0xFFFFEAA0, // 2 3 10 1 + 0x5FFFAAA0, // 2 4 0 10 + 0x7FFFAAA0, // 2 4 1 9 + 0xFFFFAAA0, // 2 4 2 8 + 0xFFFFAAA0, // 2 4 3 7 + 0xFFFFAAA0, // 2 4 4 6 + 0xFFFFAAA0, // 2 4 5 5 + 0xFFFFAAA0, // 2 4 6 4 + 0xFFFFAAA0, // 2 4 7 3 + 0xFFFFAAA0, // 2 4 8 2 + 0xFFFFAAA0, // 2 4 9 1 + 0xFFFEAAA0, // 2 5 0 9 + 0xFFFEAAA0, // 2 5 1 8 + 0xFFFEAAA0, // 2 5 2 7 + 0xFFFEAAA0, // 2 5 3 6 + 0xFFFEAAA0, // 2 5 4 5 + 0xFFFEAAA0, // 2 5 5 4 + 0xFFFEAAA0, // 2 5 6 3 + 0xFFFEAAA0, // 2 5 7 2 + 0xFFFEAAA0, // 2 5 8 1 + 0xFFFAAAA0, // 2 6 0 8 + 0xFFFAAAA0, // 2 6 1 7 + 0xFFFAAAA0, // 2 6 2 6 + 0xFFFAAAA0, // 2 6 3 5 + 0xFFFAAAA0, // 2 6 4 4 + 0xFFFAAAA0, // 2 6 5 3 + 0xFFFAAAA0, // 2 6 6 2 + 0xFFFAAAA0, // 2 6 7 1 + 0xFFEAAAA0, // 2 7 0 7 + 0xFFEAAAA0, // 2 7 1 6 + 0xFFEAAAA0, // 2 7 2 5 + 0xFFEAAAA0, // 2 7 3 4 + 0xFFEAAAA0, // 2 7 4 3 + 0xFFEAAAA0, // 2 7 5 2 + 0xFFEAAAA0, // 2 7 6 1 + 0xFFAAAAA0, // 2 8 0 6 + 0xFFAAAAA0, // 2 8 1 5 + 0xFFAAAAA0, // 2 8 2 4 + 0xFFAAAAA0, // 2 8 3 3 + 0xFFAAAAA0, // 2 8 4 2 + 0xFFAAAAA0, // 2 8 5 1 + 0xFEAAAAA0, // 2 9 0 5 + 0xFEAAAAA0, // 2 9 1 4 + 0xFEAAAAA0, // 2 9 2 3 + 0xFEAAAAA0, // 2 9 3 2 + 0xFEAAAAA0, // 2 9 4 1 + 0xFAAAAAA0, // 2 10 0 4 + 0xFAAAAAA0, // 2 10 1 3 + 0xFAAAAAA0, // 2 10 2 2 + 0xFAAAAAA0, // 2 10 3 1 + 0xEAAAAAA0, // 2 11 0 3 + 0xEAAAAAA0, // 2 11 1 2 + 0xEAAAAAA0, // 2 11 2 1 + 0xAAAAAAA0, // 2 12 0 2 + 0xAAAAAAA0, // 2 12 1 1 + 0xAAAAAAA0, // 2 13 0 1 + 0x5557FA80, // 3 0 0 13 + 0x555FFA80, // 3 0 1 12 + 0x557FFA80, // 3 0 2 11 + 0x55FFFA80, // 3 0 3 10 + 0x57FFFA80, // 3 0 4 9 + 0x5FFFFA80, // 3 0 5 8 + 0x7FFFFA80, // 3 0 6 7 + 0xFFFFFA80, // 3 0 7 6 + 0xFFFFFA80, // 3 0 8 5 + 0xFFFFFA80, // 3 0 9 4 + 0xFFFFFA80, // 3 0 10 3 + 0xFFFFFA80, // 3 0 11 2 + 0xFFFFFA80, // 3 0 12 1 + 0x557FEA80, // 3 1 0 12 + 0x55FFEA80, // 3 1 1 11 + 0x57FFEA80, // 3 1 2 10 + 0x5FFFEA80, // 3 1 3 9 + 0x7FFFEA80, // 3 1 4 8 + 0xFFFFEA80, // 3 1 5 7 + 0xFFFFEA80, // 3 1 6 6 + 0xFFFFEA80, // 3 1 7 5 + 0xFFFFEA80, // 3 1 8 4 + 0xFFFFEA80, // 3 1 9 3 + 0xFFFFEA80, // 3 1 10 2 + 0xFFFFEA80, // 3 1 11 1 + 0x57FFAA80, // 3 2 0 11 + 0x5FFFAA80, // 3 2 1 10 + 0x7FFFAA80, // 3 2 2 9 + 0xFFFFAA80, // 3 2 3 8 + 0xFFFFAA80, // 3 2 4 7 + 0xFFFFAA80, // 3 2 5 6 + 0xFFFFAA80, // 3 2 6 5 + 0xFFFFAA80, // 3 2 7 4 + 0xFFFFAA80, // 3 2 8 3 + 0xFFFFAA80, // 3 2 9 2 + 0xFFFFAA80, // 3 2 10 1 + 0x7FFEAA80, // 3 3 0 10 + 0xFFFEAA80, // 3 3 1 9 + 0xFFFEAA80, // 3 3 2 8 + 0xFFFEAA80, // 3 3 3 7 + 0xFFFEAA80, // 3 3 4 6 + 0xFFFEAA80, // 3 3 5 5 + 0xFFFEAA80, // 3 3 6 4 + 0xFFFEAA80, // 3 3 7 3 + 0xFFFEAA80, // 3 3 8 2 + 0xFFFEAA80, // 3 3 9 1 + 0xFFFAAA80, // 3 4 0 9 + 0xFFFAAA80, // 3 4 1 8 + 0xFFFAAA80, // 3 4 2 7 + 0xFFFAAA80, // 3 4 3 6 + 0xFFFAAA80, // 3 4 4 5 + 0xFFFAAA80, // 3 4 5 4 + 0xFFFAAA80, // 3 4 6 3 + 0xFFFAAA80, // 3 4 7 2 + 0xFFFAAA80, // 3 4 8 1 + 0xFFEAAA80, // 3 5 0 8 + 0xFFEAAA80, // 3 5 1 7 + 0xFFEAAA80, // 3 5 2 6 + 0xFFEAAA80, // 3 5 3 5 + 0xFFEAAA80, // 3 5 4 4 + 0xFFEAAA80, // 3 5 5 3 + 0xFFEAAA80, // 3 5 6 2 + 0xFFEAAA80, // 3 5 7 1 + 0xFFAAAA80, // 3 6 0 7 + 0xFFAAAA80, // 3 6 1 6 + 0xFFAAAA80, // 3 6 2 5 + 0xFFAAAA80, // 3 6 3 4 + 0xFFAAAA80, // 3 6 4 3 + 0xFFAAAA80, // 3 6 5 2 + 0xFFAAAA80, // 3 6 6 1 + 0xFEAAAA80, // 3 7 0 6 + 0xFEAAAA80, // 3 7 1 5 + 0xFEAAAA80, // 3 7 2 4 + 0xFEAAAA80, // 3 7 3 3 + 0xFEAAAA80, // 3 7 4 2 + 0xFEAAAA80, // 3 7 5 1 + 0xFAAAAA80, // 3 8 0 5 + 0xFAAAAA80, // 3 8 1 4 + 0xFAAAAA80, // 3 8 2 3 + 0xFAAAAA80, // 3 8 3 2 + 0xFAAAAA80, // 3 8 4 1 + 0xEAAAAA80, // 3 9 0 4 + 0xEAAAAA80, // 3 9 1 3 + 0xEAAAAA80, // 3 9 2 2 + 0xEAAAAA80, // 3 9 3 1 + 0xAAAAAA80, // 3 10 0 3 + 0xAAAAAA80, // 3 10 1 2 + 0xAAAAAA80, // 3 10 2 1 + 0xAAAAAA80, // 3 11 0 2 + 0xAAAAAA80, // 3 11 1 1 + 0xAAAAAA80, // 3 12 0 1 + 0x55FFAA00, // 4 0 0 12 + 0x57FFAA00, // 4 0 1 11 + 0x5FFFAA00, // 4 0 2 10 + 0x7FFFAA00, // 4 0 3 9 + 0xFFFFAA00, // 4 0 4 8 + 0xFFFFAA00, // 4 0 5 7 + 0xFFFFAA00, // 4 0 6 6 + 0xFFFFAA00, // 4 0 7 5 + 0xFFFFAA00, // 4 0 8 4 + 0xFFFFAA00, // 4 0 9 3 + 0xFFFFAA00, // 4 0 10 2 + 0xFFFFAA00, // 4 0 11 1 + 0x5FFEAA00, // 4 1 0 11 + 0x7FFEAA00, // 4 1 1 10 + 0xFFFEAA00, // 4 1 2 9 + 0xFFFEAA00, // 4 1 3 8 + 0xFFFEAA00, // 4 1 4 7 + 0xFFFEAA00, // 4 1 5 6 + 0xFFFEAA00, // 4 1 6 5 + 0xFFFEAA00, // 4 1 7 4 + 0xFFFEAA00, // 4 1 8 3 + 0xFFFEAA00, // 4 1 9 2 + 0xFFFEAA00, // 4 1 10 1 + 0xFFFAAA00, // 4 2 0 10 + 0xFFFAAA00, // 4 2 1 9 + 0xFFFAAA00, // 4 2 2 8 + 0xFFFAAA00, // 4 2 3 7 + 0xFFFAAA00, // 4 2 4 6 + 0xFFFAAA00, // 4 2 5 5 + 0xFFFAAA00, // 4 2 6 4 + 0xFFFAAA00, // 4 2 7 3 + 0xFFFAAA00, // 4 2 8 2 + 0xFFFAAA00, // 4 2 9 1 + 0xFFEAAA00, // 4 3 0 9 + 0xFFEAAA00, // 4 3 1 8 + 0xFFEAAA00, // 4 3 2 7 + 0xFFEAAA00, // 4 3 3 6 + 0xFFEAAA00, // 4 3 4 5 + 0xFFEAAA00, // 4 3 5 4 + 0xFFEAAA00, // 4 3 6 3 + 0xFFEAAA00, // 4 3 7 2 + 0xFFEAAA00, // 4 3 8 1 + 0xFFAAAA00, // 4 4 0 8 + 0xFFAAAA00, // 4 4 1 7 + 0xFFAAAA00, // 4 4 2 6 + 0xFFAAAA00, // 4 4 3 5 + 0xFFAAAA00, // 4 4 4 4 + 0xFFAAAA00, // 4 4 5 3 + 0xFFAAAA00, // 4 4 6 2 + 0xFFAAAA00, // 4 4 7 1 + 0xFEAAAA00, // 4 5 0 7 + 0xFEAAAA00, // 4 5 1 6 + 0xFEAAAA00, // 4 5 2 5 + 0xFEAAAA00, // 4 5 3 4 + 0xFEAAAA00, // 4 5 4 3 + 0xFEAAAA00, // 4 5 5 2 + 0xFEAAAA00, // 4 5 6 1 + 0xFAAAAA00, // 4 6 0 6 + 0xFAAAAA00, // 4 6 1 5 + 0xFAAAAA00, // 4 6 2 4 + 0xFAAAAA00, // 4 6 3 3 + 0xFAAAAA00, // 4 6 4 2 + 0xFAAAAA00, // 4 6 5 1 + 0xEAAAAA00, // 4 7 0 5 + 0xEAAAAA00, // 4 7 1 4 + 0xEAAAAA00, // 4 7 2 3 + 0xEAAAAA00, // 4 7 3 2 + 0xEAAAAA00, // 4 7 4 1 + 0xAAAAAA00, // 4 8 0 4 + 0xAAAAAA00, // 4 8 1 3 + 0xAAAAAA00, // 4 8 2 2 + 0xAAAAAA00, // 4 8 3 1 + 0xAAAAAA00, // 4 9 0 3 + 0xAAAAAA00, // 4 9 1 2 + 0xAAAAAA00, // 4 9 2 1 + 0xAAAAAA00, // 4 10 0 2 + 0xAAAAAA00, // 4 10 1 1 + 0xAAAAAA00, // 4 11 0 1 + 0x7FFAA800, // 5 0 0 11 + 0xFFFAA800, // 5 0 1 10 + 0xFFFAA800, // 5 0 2 9 + 0xFFFAA800, // 5 0 3 8 + 0xFFFAA800, // 5 0 4 7 + 0xFFFAA800, // 5 0 5 6 + 0xFFFAA800, // 5 0 6 5 + 0xFFFAA800, // 5 0 7 4 + 0xFFFAA800, // 5 0 8 3 + 0xFFFAA800, // 5 0 9 2 + 0xFFFAA800, // 5 0 10 1 + 0xFFEAA800, // 5 1 0 10 + 0xFFEAA800, // 5 1 1 9 + 0xFFEAA800, // 5 1 2 8 + 0xFFEAA800, // 5 1 3 7 + 0xFFEAA800, // 5 1 4 6 + 0xFFEAA800, // 5 1 5 5 + 0xFFEAA800, // 5 1 6 4 + 0xFFEAA800, // 5 1 7 3 + 0xFFEAA800, // 5 1 8 2 + 0xFFEAA800, // 5 1 9 1 + 0xFFAAA800, // 5 2 0 9 + 0xFFAAA800, // 5 2 1 8 + 0xFFAAA800, // 5 2 2 7 + 0xFFAAA800, // 5 2 3 6 + 0xFFAAA800, // 5 2 4 5 + 0xFFAAA800, // 5 2 5 4 + 0xFFAAA800, // 5 2 6 3 + 0xFFAAA800, // 5 2 7 2 + 0xFFAAA800, // 5 2 8 1 + 0xFEAAA800, // 5 3 0 8 + 0xFEAAA800, // 5 3 1 7 + 0xFEAAA800, // 5 3 2 6 + 0xFEAAA800, // 5 3 3 5 + 0xFEAAA800, // 5 3 4 4 + 0xFEAAA800, // 5 3 5 3 + 0xFEAAA800, // 5 3 6 2 + 0xFEAAA800, // 5 3 7 1 + 0xFAAAA800, // 5 4 0 7 + 0xFAAAA800, // 5 4 1 6 + 0xFAAAA800, // 5 4 2 5 + 0xFAAAA800, // 5 4 3 4 + 0xFAAAA800, // 5 4 4 3 + 0xFAAAA800, // 5 4 5 2 + 0xFAAAA800, // 5 4 6 1 + 0xEAAAA800, // 5 5 0 6 + 0xEAAAA800, // 5 5 1 5 + 0xEAAAA800, // 5 5 2 4 + 0xEAAAA800, // 5 5 3 3 + 0xEAAAA800, // 5 5 4 2 + 0xEAAAA800, // 5 5 5 1 + 0xAAAAA800, // 5 6 0 5 + 0xAAAAA800, // 5 6 1 4 + 0xAAAAA800, // 5 6 2 3 + 0xAAAAA800, // 5 6 3 2 + 0xAAAAA800, // 5 6 4 1 + 0xAAAAA800, // 5 7 0 4 + 0xAAAAA800, // 5 7 1 3 + 0xAAAAA800, // 5 7 2 2 + 0xAAAAA800, // 5 7 3 1 + 0xAAAAA800, // 5 8 0 3 + 0xAAAAA800, // 5 8 1 2 + 0xAAAAA800, // 5 8 2 1 + 0xAAAAA800, // 5 9 0 2 + 0xAAAAA800, // 5 9 1 1 + 0xAAAAA800, // 5 10 0 1 + 0xFFAAA000, // 6 0 0 10 + 0xFFAAA000, // 6 0 1 9 + 0xFFAAA000, // 6 0 2 8 + 0xFFAAA000, // 6 0 3 7 + 0xFFAAA000, // 6 0 4 6 + 0xFFAAA000, // 6 0 5 5 + 0xFFAAA000, // 6 0 6 4 + 0xFFAAA000, // 6 0 7 3 + 0xFFAAA000, // 6 0 8 2 + 0xFFAAA000, // 6 0 9 1 + 0xFEAAA000, // 6 1 0 9 + 0xFEAAA000, // 6 1 1 8 + 0xFEAAA000, // 6 1 2 7 + 0xFEAAA000, // 6 1 3 6 + 0xFEAAA000, // 6 1 4 5 + 0xFEAAA000, // 6 1 5 4 + 0xFEAAA000, // 6 1 6 3 + 0xFEAAA000, // 6 1 7 2 + 0xFEAAA000, // 6 1 8 1 + 0xFAAAA000, // 6 2 0 8 + 0xFAAAA000, // 6 2 1 7 + 0xFAAAA000, // 6 2 2 6 + 0xFAAAA000, // 6 2 3 5 + 0xFAAAA000, // 6 2 4 4 + 0xFAAAA000, // 6 2 5 3 + 0xFAAAA000, // 6 2 6 2 + 0xFAAAA000, // 6 2 7 1 + 0xEAAAA000, // 6 3 0 7 + 0xEAAAA000, // 6 3 1 6 + 0xEAAAA000, // 6 3 2 5 + 0xEAAAA000, // 6 3 3 4 + 0xEAAAA000, // 6 3 4 3 + 0xEAAAA000, // 6 3 5 2 + 0xEAAAA000, // 6 3 6 1 + 0xAAAAA000, // 6 4 0 6 + 0xAAAAA000, // 6 4 1 5 + 0xAAAAA000, // 6 4 2 4 + 0xAAAAA000, // 6 4 3 3 + 0xAAAAA000, // 6 4 4 2 + 0xAAAAA000, // 6 4 5 1 + 0xAAAAA000, // 6 5 0 5 + 0xAAAAA000, // 6 5 1 4 + 0xAAAAA000, // 6 5 2 3 + 0xAAAAA000, // 6 5 3 2 + 0xAAAAA000, // 6 5 4 1 + 0xAAAAA000, // 6 6 0 4 + 0xAAAAA000, // 6 6 1 3 + 0xAAAAA000, // 6 6 2 2 + 0xAAAAA000, // 6 6 3 1 + 0xAAAAA000, // 6 7 0 3 + 0xAAAAA000, // 6 7 1 2 + 0xAAAAA000, // 6 7 2 1 + 0xAAAAA000, // 6 8 0 2 + 0xAAAAA000, // 6 8 1 1 + 0xAAAAA000, // 6 9 0 1 + 0xFAAA8000, // 7 0 0 9 + 0xFAAA8000, // 7 0 1 8 + 0xFAAA8000, // 7 0 2 7 + 0xFAAA8000, // 7 0 3 6 + 0xFAAA8000, // 7 0 4 5 + 0xFAAA8000, // 7 0 5 4 + 0xFAAA8000, // 7 0 6 3 + 0xFAAA8000, // 7 0 7 2 + 0xFAAA8000, // 7 0 8 1 + 0xEAAA8000, // 7 1 0 8 + 0xEAAA8000, // 7 1 1 7 + 0xEAAA8000, // 7 1 2 6 + 0xEAAA8000, // 7 1 3 5 + 0xEAAA8000, // 7 1 4 4 + 0xEAAA8000, // 7 1 5 3 + 0xEAAA8000, // 7 1 6 2 + 0xEAAA8000, // 7 1 7 1 + 0xAAAA8000, // 7 2 0 7 + 0xAAAA8000, // 7 2 1 6 + 0xAAAA8000, // 7 2 2 5 + 0xAAAA8000, // 7 2 3 4 + 0xAAAA8000, // 7 2 4 3 + 0xAAAA8000, // 7 2 5 2 + 0xAAAA8000, // 7 2 6 1 + 0xAAAA8000, // 7 3 0 6 + 0xAAAA8000, // 7 3 1 5 + 0xAAAA8000, // 7 3 2 4 + 0xAAAA8000, // 7 3 3 3 + 0xAAAA8000, // 7 3 4 2 + 0xAAAA8000, // 7 3 5 1 + 0xAAAA8000, // 7 4 0 5 + 0xAAAA8000, // 7 4 1 4 + 0xAAAA8000, // 7 4 2 3 + 0xAAAA8000, // 7 4 3 2 + 0xAAAA8000, // 7 4 4 1 + 0xAAAA8000, // 7 5 0 4 + 0xAAAA8000, // 7 5 1 3 + 0xAAAA8000, // 7 5 2 2 + 0xAAAA8000, // 7 5 3 1 + 0xAAAA8000, // 7 6 0 3 + 0xAAAA8000, // 7 6 1 2 + 0xAAAA8000, // 7 6 2 1 + 0xAAAA8000, // 7 7 0 2 + 0xAAAA8000, // 7 7 1 1 + 0xAAAA8000, // 7 8 0 1 + 0xAAAA0000, // 8 0 0 8 + 0xAAAA0000, // 8 0 1 7 + 0xAAAA0000, // 8 0 2 6 + 0xAAAA0000, // 8 0 3 5 + 0xAAAA0000, // 8 0 4 4 + 0xAAAA0000, // 8 0 5 3 + 0xAAAA0000, // 8 0 6 2 + 0xAAAA0000, // 8 0 7 1 + 0xAAAA0000, // 8 1 0 7 + 0xAAAA0000, // 8 1 1 6 + 0xAAAA0000, // 8 1 2 5 + 0xAAAA0000, // 8 1 3 4 + 0xAAAA0000, // 8 1 4 3 + 0xAAAA0000, // 8 1 5 2 + 0xAAAA0000, // 8 1 6 1 + 0xAAAA0000, // 8 2 0 6 + 0xAAAA0000, // 8 2 1 5 + 0xAAAA0000, // 8 2 2 4 + 0xAAAA0000, // 8 2 3 3 + 0xAAAA0000, // 8 2 4 2 + 0xAAAA0000, // 8 2 5 1 + 0xAAAA0000, // 8 3 0 5 + 0xAAAA0000, // 8 3 1 4 + 0xAAAA0000, // 8 3 2 3 + 0xAAAA0000, // 8 3 3 2 + 0xAAAA0000, // 8 3 4 1 + 0xAAAA0000, // 8 4 0 4 + 0xAAAA0000, // 8 4 1 3 + 0xAAAA0000, // 8 4 2 2 + 0xAAAA0000, // 8 4 3 1 + 0xAAAA0000, // 8 5 0 3 + 0xAAAA0000, // 8 5 1 2 + 0xAAAA0000, // 8 5 2 1 + 0xAAAA0000, // 8 6 0 2 + 0xAAAA0000, // 8 6 1 1 + 0xAAAA0000, // 8 7 0 1 + 0xAAA80000, // 9 0 0 7 + 0xAAA80000, // 9 0 1 6 + 0xAAA80000, // 9 0 2 5 + 0xAAA80000, // 9 0 3 4 + 0xAAA80000, // 9 0 4 3 + 0xAAA80000, // 9 0 5 2 + 0xAAA80000, // 9 0 6 1 + 0xAAA80000, // 9 1 0 6 + 0xAAA80000, // 9 1 1 5 + 0xAAA80000, // 9 1 2 4 + 0xAAA80000, // 9 1 3 3 + 0xAAA80000, // 9 1 4 2 + 0xAAA80000, // 9 1 5 1 + 0xAAA80000, // 9 2 0 5 + 0xAAA80000, // 9 2 1 4 + 0xAAA80000, // 9 2 2 3 + 0xAAA80000, // 9 2 3 2 + 0xAAA80000, // 9 2 4 1 + 0xAAA80000, // 9 3 0 4 + 0xAAA80000, // 9 3 1 3 + 0xAAA80000, // 9 3 2 2 + 0xAAA80000, // 9 3 3 1 + 0xAAA80000, // 9 4 0 3 + 0xAAA80000, // 9 4 1 2 + 0xAAA80000, // 9 4 2 1 + 0xAAA80000, // 9 5 0 2 + 0xAAA80000, // 9 5 1 1 + 0xAAA80000, // 9 6 0 1 + 0xAAA00000, // 10 0 0 6 + 0xAAA00000, // 10 0 1 5 + 0xAAA00000, // 10 0 2 4 + 0xAAA00000, // 10 0 3 3 + 0xAAA00000, // 10 0 4 2 + 0xAAA00000, // 10 0 5 1 + 0xAAA00000, // 10 1 0 5 + 0xAAA00000, // 10 1 1 4 + 0xAAA00000, // 10 1 2 3 + 0xAAA00000, // 10 1 3 2 + 0xAAA00000, // 10 1 4 1 + 0xAAA00000, // 10 2 0 4 + 0xAAA00000, // 10 2 1 3 + 0xAAA00000, // 10 2 2 2 + 0xAAA00000, // 10 2 3 1 + 0xAAA00000, // 10 3 0 3 + 0xAAA00000, // 10 3 1 2 + 0xAAA00000, // 10 3 2 1 + 0xAAA00000, // 10 4 0 2 + 0xAAA00000, // 10 4 1 1 + 0xAAA00000, // 10 5 0 1 + 0xAA800000, // 11 0 0 5 + 0xAA800000, // 11 0 1 4 + 0xAA800000, // 11 0 2 3 + 0xAA800000, // 11 0 3 2 + 0xAA800000, // 11 0 4 1 + 0xAA800000, // 11 1 0 4 + 0xAA800000, // 11 1 1 3 + 0xAA800000, // 11 1 2 2 + 0xAA800000, // 11 1 3 1 + 0xAA800000, // 11 2 0 3 + 0xAA800000, // 11 2 1 2 + 0xAA800000, // 11 2 2 1 + 0xAA800000, // 11 3 0 2 + 0xAA800000, // 11 3 1 1 + 0xAA800000, // 11 4 0 1 + 0xAA000000, // 12 0 0 4 + 0xAA000000, // 12 0 1 3 + 0xAA000000, // 12 0 2 2 + 0xAA000000, // 12 0 3 1 + 0xAA000000, // 12 1 0 3 + 0xAA000000, // 12 1 1 2 + 0xAA000000, // 12 1 2 1 + 0xAA000000, // 12 2 0 2 + 0xAA000000, // 12 2 1 1 + 0xAA000000, // 12 3 0 1 + 0xA8000000, // 13 0 0 3 + 0xA8000000, // 13 0 1 2 + 0xA8000000, // 13 0 2 1 + 0xA8000000, // 13 1 0 2 + 0xA8000000, // 13 1 1 1 + 0xA8000000, // 13 2 0 1 + 0xA0000000, // 14 0 0 2 + 0xA0000000, // 14 0 1 1 + 0xA0000000, // 14 1 0 1 + 0x80000000, // 15 0 0 1 + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, + 0x80000000, +}; + diff --git a/src/nvtt/cuda/CompressKernel.cu b/src/nvtt/cuda/CompressKernel.cu index 43fc3f7..7df67b4 100644 --- a/src/nvtt/cuda/CompressKernel.cu +++ b/src/nvtt/cuda/CompressKernel.cu @@ -159,7 +159,7 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum } } -__device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sums[16], float weights[16], int xrefs[16]) +__device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sums[16], float weights[16], int xrefs[16], int * sameColor) { const int bid = blockIdx.x; const int idx = threadIdx.x; @@ -189,6 +189,8 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum colorSums(colors, sums); float3 axis = bestFitLine(colors, sums[0], kColorMetric); + *sameColor = (axis == make_float3(0, 0, 0)); + dps[idx] = dot(rawColors[idx], axis); #if __DEVICE_EMULATION__ @@ -205,7 +207,7 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum } } -__device__ void loadColorBlock(const uint * image, float2 colors[16], float2 sums[16], int xrefs[16]) +__device__ void loadColorBlock(const uint * image, float2 colors[16], float2 sums[16], int xrefs[16], int * sameColor) { const int bid = blockIdx.x; const int idx = threadIdx.x; @@ -229,6 +231,8 @@ __device__ void loadColorBlock(const uint * image, float2 colors[16], float2 sum colorSums(colors, sums); float2 axis = bestFitLine(colors, sums[0]); + *sameColor = (axis == make_float2(0, 0)); + dps[idx] = dot(colors[idx], axis); #if __DEVICE_EMULATION__ @@ -861,7 +865,7 @@ __device__ void evalAllPermutationsCTX(const float2 * colors, float2 colorSum, c for(int i = 0; i < 16; i++) { int pidx = idx + NUM_THREADS * i; - if (pidx >= 992) break; + if (pidx >= 704) break; ushort start, end; uint permutation = permutations[pidx]; @@ -1024,6 +1028,41 @@ __device__ void saveSingleColorBlockDXT1(float3 color, uint2 * result) } } +__device__ void saveSingleColorBlockDXT1(float2 color, uint2 * result) +{ + const int bid = blockIdx.x; + + int r = color.x * 255; + int g = color.y * 255; + + ushort color0 = (OMatch5[r][0] << 11) | (OMatch6[g][0] << 5); + ushort color1 = (OMatch5[r][1] << 11) | (OMatch6[g][1] << 5); + + if (color0 < color1) + { + result[bid].x = (color0 << 16) | color1; + result[bid].y = 0xffffffff; + } + else + { + result[bid].x = (color1 << 16) | color0; + result[bid].y = 0xaaaaaaaa; + } +} + +__device__ void saveSingleColorBlockCTX1(float2 color, uint2 * result) +{ + const int bid = blockIdx.x; + + int r = color.x * 255; + int g = color.y * 255; + + ushort color0 = (r << 8) | (g); + + result[bid].x = (color0 << 16) | color0; + result[bid].y = 0x00000000; +} + //////////////////////////////////////////////////////////////////////////////// // Compress color block @@ -1069,11 +1108,18 @@ __global__ void compressWeightedDXT1(const uint * permutations, const uint * ima __shared__ float3 sums[16]; __shared__ float weights[16]; __shared__ int xrefs[16]; + __shared__ int sameColor; - loadColorBlock(image, colors, sums, weights, xrefs); + loadColorBlock(image, colors, sums, weights, xrefs, &sameColor); __syncthreads(); + if (sameColor) + { + if (threadIdx.x == 0) saveSingleColorBlockDXT1(colors[0], result); + return; + } + ushort bestStart, bestEnd; uint bestPermutation; @@ -1097,11 +1143,18 @@ __global__ void compressNormalDXT1(const uint * permutations, const uint * image __shared__ float2 colors[16]; __shared__ float2 sums[16]; __shared__ int xrefs[16]; - - loadColorBlock(image, colors, sums, xrefs); + __shared__ int sameColor; + + loadColorBlock(image, colors, sums, xrefs, &sameColor); __syncthreads(); + if (sameColor) + { + if (threadIdx.x == 0) saveSingleColorBlockDXT1(colors[0], result); + return; + } + ushort bestStart, bestEnd; uint bestPermutation; @@ -1124,11 +1177,18 @@ __global__ void compressCTX1(const uint * permutations, const uint * image, uint __shared__ float2 colors[16]; __shared__ float2 sums[16]; __shared__ int xrefs[16]; + __shared__ int sameColor; - loadColorBlock(image, colors, sums, xrefs); + loadColorBlock(image, colors, sums, xrefs, &sameColor); __syncthreads(); + if (sameColor) + { + if (threadIdx.x == 0) saveSingleColorBlockCTX1(colors[0], result); + return; + } + ushort bestStart, bestEnd; uint bestPermutation; diff --git a/src/nvtt/cuda/CudaCompressDXT.cpp b/src/nvtt/cuda/CudaCompressDXT.cpp index b18adf9..fd89c1c 100644 --- a/src/nvtt/cuda/CudaCompressDXT.cpp +++ b/src/nvtt/cuda/CudaCompressDXT.cpp @@ -83,7 +83,7 @@ static void convertToBlockLinear(const Image * image, uint * blockLinearImage) #endif -CudaCompressor::CudaCompressor() : m_bitmapTable(NULL), m_data(NULL), m_result(NULL) +CudaCompressor::CudaCompressor() : m_bitmapTable(NULL), m_bitmapTableCTX(NULL), m_data(NULL), m_result(NULL) { #if defined HAVE_CUDA // Allocate and upload bitmaps. @@ -93,6 +93,12 @@ CudaCompressor::CudaCompressor() : m_bitmapTable(NULL), m_data(NULL), m_result(N cudaMemcpy(m_bitmapTable, s_bitmapTable, 992 * sizeof(uint), cudaMemcpyHostToDevice); } + cudaMalloc((void**) &m_bitmapTableCTX, 704 * sizeof(uint)); + if (m_bitmapTableCTX != NULL) + { + cudaMemcpy(m_bitmapTableCTX, s_bitmapTableCTX, 704 * sizeof(uint), cudaMemcpyHostToDevice); + } + // Allocate scratch buffers. cudaMalloc((void**) &m_data, MAX_BLOCKS * 64U); cudaMalloc((void**) &m_result, MAX_BLOCKS * 8U); @@ -558,7 +564,7 @@ void CudaCompressor::compressCTX1(const nvtt::CompressionOptions::Private & comp cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice); // Launch kernel. - compressKernelCTX1(count, m_data, m_result, m_bitmapTable); + compressKernelCTX1(count, m_data, m_result, m_bitmapTableCTX); // Check for errors. cudaError_t err = cudaGetLastError(); diff --git a/src/nvtt/cuda/CudaCompressDXT.h b/src/nvtt/cuda/CudaCompressDXT.h index 5fbccde..7567943 100644 --- a/src/nvtt/cuda/CudaCompressDXT.h +++ b/src/nvtt/cuda/CudaCompressDXT.h @@ -50,6 +50,7 @@ namespace nv private: uint * m_bitmapTable; + uint * m_bitmapTableCTX; uint * m_data; uint * m_result; diff --git a/src/nvtt/cuda/CudaMath.h b/src/nvtt/cuda/CudaMath.h index 13e27df..88cd233 100644 --- a/src/nvtt/cuda/CudaMath.h +++ b/src/nvtt/cuda/CudaMath.h @@ -139,6 +139,11 @@ inline __device__ __host__ void operator /=(float2 & b, float f) b.y *= inv; } +inline __device__ __host__ bool operator ==(float2 a, float2 b) +{ + return a.x == b.x && a.y == b.y; +} + inline __device__ __host__ float dot(float2 a, float2 b) { diff --git a/src/nvtt/cuda/CudaUtils.cpp b/src/nvtt/cuda/CudaUtils.cpp index 3c5ca7c..cd84cf1 100644 --- a/src/nvtt/cuda/CudaUtils.cpp +++ b/src/nvtt/cuda/CudaUtils.cpp @@ -74,7 +74,7 @@ bool nv::cuda::isHardwarePresent() { #if defined HAVE_CUDA #if NV_OS_WIN32 - if (isWindowsVista()) return false; + //if (isWindowsVista()) return false; //if (isWindowsVista() || !isWow32()) return false; #endif int count = deviceCount(); diff --git a/src/nvtt/tests/stress.cpp b/src/nvtt/tests/stress.cpp index ea86d1b..cdfd6a6 100644 --- a/src/nvtt/tests/stress.cpp +++ b/src/nvtt/tests/stress.cpp @@ -34,7 +34,7 @@ #define WIDTH 2048 #define HEIGHT 2048 #define INPUT_SIZE (WIDTH*HEIGHT) -#define OUTPUT_SIZE (WIDTH*HEIGHT/16*2) +#define OUTPUT_SIZE (WIDTH*HEIGHT/16*4) static int s_input[INPUT_SIZE]; static int s_reference[OUTPUT_SIZE]; @@ -89,12 +89,8 @@ void precomp() int num = 0; - printf("{\n"); - printf("\t%8X,\n", 0); + printf("const static uint s_bitmapTableCTX[704] =\n{\n"); - bitmaps[0] = 0; - - num = 1; for (int a = 1; a <= 15; a++) { for (int b = a; b <= 15; b++) @@ -130,9 +126,16 @@ void precomp() } } - printf("}\n"); + // Align to 32: 680 -> 704 + while (num < 704) + { + printf("\t0x80000000,\n"); - printf("// num = %d\n", num); + bitmaps[num] = 0x80000000; // 15 0 0 1; + num++; + } + + printf("}; // num = %d\n", num); /* for( int i = imax; i >= 0; --i ) @@ -177,7 +180,7 @@ void precomp() int main(int argc, char *argv[]) { - //precomp(); +// precomp(); nvtt::InputOptions inputOptions; inputOptions.setTextureLayout(nvtt::TextureType_2D, WIDTH, HEIGHT); @@ -191,7 +194,7 @@ int main(int argc, char *argv[]) inputOptions.setMipmapGeneration(false); nvtt::CompressionOptions compressionOptions; - compressionOptions.setFormat(nvtt::Format_DXT1); +// compressionOptions.setFormat(nvtt::Format_DXT3); // compressionOptions.setFormat(nvtt::Format_DXT1n); // compressionOptions.setFormat(nvtt::Format_CTX1); @@ -203,6 +206,7 @@ int main(int argc, char *argv[]) nvtt::Compressor compressor; +// compressor.enableCudaAcceleration(false); for (s_frame = 0; s_frame < FRAME_COUNT; s_frame++) {