| 48 | } |
| 49 | |
| 50 | inline float dDequantizeNF4(unsigned char val) { |
| 51 | |
| 52 | // the values for this tree was generated by test_normal_map_tree |
| 53 | // in the file tests/test_functional.py |
| 54 | if ((val & 0b1000) == 8) |
| 55 | if ((val & 0b0100) == 4) // 1 |
| 56 | if ((val & 0b0010) == 2) // 11 |
| 57 | if ((val & 0b0001) == 1) // 111 |
| 58 | return 1.0f; //*1111 |
| 59 | else |
| 60 | return 0.7229568362236023f; //*1110 |
| 61 | else if ((val & 0b0001) == 1) // 110 |
| 62 | return 0.5626170039176941f; //*1101 |
| 63 | else |
| 64 | return 0.44070982933044434f; //*1100 |
| 65 | else if ((val & 0b0010) == 2) // 10 |
| 66 | if ((val & 0b0001) == 1) // 101 |
| 67 | return 0.33791524171829224f; //*1011 |
| 68 | else |
| 69 | return 0.24611230194568634f; //*1010 |
| 70 | else if ((val & 0b0001) == 1) // 100 |
| 71 | return 0.16093020141124725f; //*1001 |
| 72 | else |
| 73 | return 0.07958029955625534f; //*1000 |
| 74 | |
| 75 | else if ((val & 0b0100) == 4) // 0 |
| 76 | if ((val & 0b0010) == 2) // 01 |
| 77 | if ((val & 0b0001) == 1) // 011 |
| 78 | return 0.0f; //*0111 |
| 79 | else |
| 80 | return -0.09105003625154495f; //*0110 |
| 81 | else if ((val & 0b0001) == 1) // 010 |
| 82 | return -0.18477343022823334f; //*0101 |
| 83 | else |
| 84 | return -0.28444138169288635f; //*0100 |
| 85 | else if ((val & 0b0010) == 2) // 00 |
| 86 | if ((val & 0b0001) == 1) // 001 |
| 87 | return -0.39491748809814453f; //*0011 |
| 88 | else |
| 89 | return -0.5250730514526367f; //*0010 |
| 90 | else if ((val & 0b0001) == 1) // 000 |
| 91 | return -0.6961928009986877f; //*0001 |
| 92 | else |
| 93 | return -1.0f; //*0000 |
| 94 | } |
| 95 | |
| 96 | template <typename T, int TILE_SIZE, int NUM_PER_TH, int DATA_TYPE> |
| 97 | SYCL_EXTERNAL void kDequantizeBlockwise<T, TILE_SIZE, NUM_PER_TH, DATA_TYPE>::operator()(sycl::nd_item<1> item) const { |