csullivan opened a new pull request, #16548:
URL: https://github.com/apache/tvm/pull/16548
```C++
#include <cuda_fp8.h>
using fp8_e4_t = __nv_fp8_e4m3;
using fp8_e4_2_t = __nv_fp8x2_e4m3;
using fp8_e4_4_t = __nv_fp8x4_e4m3;
struct __align__(8) fp8_e4_8_t {
__nv_fp8x2_e4m3 x;
__nv_fp8x2_e4m3 y;
__nv_fp8x2_e4m3 z;
__nv_fp8x2_e4m3 w;
};
using fp8_e5_t = __nv_fp8_e5m2;
using fp8_e5_2_t = __nv_fp8x2_e5m2;
using fp8_e5_4_t = __nv_fp8x4_e5m2;
struct __align__(8) fp8_e5_8_t {
__nv_fp8x2_e5m2 x;
__nv_fp8x2_e5m2 y;
__nv_fp8x2_e5m2 z;
__nv_fp8x2_e5m2 w;
};
#endif
#if (((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 4)) || \
(__CUDACC_VER_MAJOR__ > 11))
#define TVM_ENABLE_L2_PREFETCH 1
#else
#define TVM_ENABLE_L2_PREFETCH 0
#endif
#ifdef _WIN32
using uint = unsigned int;
using uchar = unsigned char;
using ushort = unsigned short;
using int64_t = long long;
using uint64_t = unsigned long long;
#else
#define uint unsigned int
#define uchar unsigned char
#define ushort unsigned short
#define int64_t long long
#define uint64_t unsigned long long
#endif
extern "C" __global__ void __launch_bounds__(32) vector_add_kernel(fp8_e4_t*
__restrict__ A, fp8_e4_t* __restrict__ B, fp8_e4_t* __restrict__ C);
extern "C" __global__ void __launch_bounds__(32) vector_add_kernel(fp8_e4_t*
__restrict__ A, fp8_e4_t* __restrict__ B, fp8_e4_t* __restrict__ C) {
C[((((int)blockIdx.x) * 32) + ((int)threadIdx.x))] =
((fp8_e4_t)(((half)A[((((int)blockIdx.x) * 32) + ((int)threadIdx.x))]) +
((half)B[((((int)blockIdx.x) * 32) + ((int)threadIdx.x))])));
}
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
