[WIP] Refactor spec

cava/nightwatch/parser/c/nightwatch.h

@@ -1,3 +1,5 @@
+#ifndef _AVA_NIGHTWATCH_PARSER_C_NIGHTWATCH_H_
+#define _AVA_NIGHTWATCH_PARSER_C_NIGHTWATCH_H_
 #include <assert.h>
 #include <stdlib.h>
 
@@ -452,3 +454,5 @@ struct __ava_unknown;
 /// API. The actual return type is the return type of the API call. If
 /// the return value is captured, the variable name must be "ret".
 struct __ava_unknown *ava_execute();
+
+#endif  // _AVA_NIGHTWATCH_PARSER_C_NIGHTWATCH_H_

cava/samples/cuda_common_spec/cublas/blas3.h

@@ -0,0 +1,55 @@
+#ifndef _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_BLAS3_H_
+#define _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_BLAS3_H_
+#include "cava/nightwatch/parser/c/nightwatch.h"
+
+/* --------------- CUBLAS BLAS3 functions  ---------------- */
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasGemmBatchedEx(cublasHandle_t handle, cublasOperation_t transa,
+                                                          cublasOperation_t transb, int m, int n, int k,
+                                                          const void *alpha, /* host or device pointer */
+                                                          const void *const Aarray[], cudaDataType Atype, int lda,
+                                                          const void *const Barray[], cudaDataType Btype, int ldb,
+                                                          const void *beta, /* host or device pointer */
+                                                          void *const Carray[], cudaDataType Ctype, int ldc,
+                                                          int batchCount, cudaDataType computeType,
+                                                          cublasGemmAlgo_t algo) {
+  ava_argument(handle) ava_handle;
+  // In tensorflow, Aarray, Barray and Carray are device memory
+  ava_argument(Aarray) ava_opaque;
+  ava_argument(Barray) ava_opaque;
+  ava_argument(Carray) ava_opaque;
+  // If they are host memory, use the following code:
+  /*
+  ava_argument(Aarray) {
+      ava_in; ava_buffer(batchCount);
+      ava_element {
+          ava_in; ava_buffer(lda * __helper_a_last_dim_size(transa, k, m) * __helper_type_size(Atype));
+      }
+  }
+
+  ava_argument(Barray) {
+      ava_in; ava_buffer(batchCount);
+      ava_element {
+          ava_in; ava_buffer(ldb * __helper_b_last_dim_size(transb, k, n) * __helper_type_size(Btype));
+      }
+  }
+
+  ava_argument(Carray) {
+      ava_type_cast(void**);
+      ava_out; ava_buffer(batchCount);
+      ava_element {
+          ava_out; ava_buffer(ldc * n * __helper_type_size(Ctype));
+      }
+  }
+  */
+  // TODO: figure out alpha and beta
+  ava_argument(alpha) {
+    ava_in;
+    ava_buffer(1);
+  }
+  ava_argument(beta) {
+    ava_in;
+    ava_buffer(1);
+  }
+}
+
+#endif  // _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_BLAS3_H_

cava/samples/cuda_common_spec/cublas/blas_like_ext.h

@@ -0,0 +1,43 @@
+#ifndef _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_BLAS_LIKE_EXT_H_
+#define _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_BLAS_LIKE_EXT_H_
+#include "cava/nightwatch/parser/c/nightwatch.h"
+/* ---------------- CUBLAS BLAS-like extension ---------------- */
+/* GEAM */
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSgeam(cublasHandle_t handle, cublasOperation_t transa,
+                                                  cublasOperation_t transb, int m, int n,
+                                                  const float *alpha, /* host or device pointer */
+                                                  const float *A, int lda,
+                                                  const float *beta, /* host or device pointer */
+                                                  const float *B, int ldb, float *C, int ldc) {
+  ava_argument(handle) ava_handle;
+  ava_argument(alpha) {
+    ava_in;
+    ava_buffer(1);
+  }
+  ava_argument(A) ava_opaque;
+  ava_argument(beta) {
+    ava_in;
+    ava_buffer(1);
+  }
+  ava_argument(B) ava_opaque;
+  ava_argument(C) ava_opaque;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSetStream(cublasHandle_t handle, cudaStream_t streamId) {
+  ava_async;
+  ava_argument(handle) ava_handle;
+  ava_argument(streamId) ava_handle;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSscal(cublasHandle_t handle, int n,
+                                                  const float *alpha, /* host or device pointer */
+                                                  float *x, int incx) {
+  ava_argument(handle) ava_handle;
+  ava_argument(alpha) {
+    ava_in;
+    ava_buffer(1);
+  }
+  ava_argument(x) ava_opaque;
+}
+
+#endif  // _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_BLAS_LIKE_EXT_H_

cava/samples/cuda_common_spec/cublas/cublas.h

@@ -0,0 +1,55 @@
+#ifndef _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_H_
+#define _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_H_
+
+#include <cublas_api.h>
+#include <cublas_v2.h>
+
+#include "cava/nightwatch/parser/c/nightwatch.h"
+#include "common/extensions/cudnn_optimization.h"
+
+/* CUDABLAS API */
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasCreate(cublasHandle_t *handle) {
+  ava_disable_native_call;
+  ava_argument(handle) {
+    ava_out;
+    ava_buffer(1);
+    ava_element { ava_handle; }
+  }
+
+  cublasStatus_t ret;
+  if (ava_is_worker) {
+    ret = __cublasCreate(handle);
+    return ret;
+  }
+}
+
+cublasStatus_t CUBLASWINAPI cublasSetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B,
+                                            int ldb) {
+  ava_argument(A) {
+    ava_in;
+    ava_buffer(rows * cols * elemSize);
+  }
+
+  ava_argument(B) { ava_opaque; }
+}
+
+cublasStatus_t CUBLASWINAPI cublasGetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B,
+                                            int ldb) {
+  ava_argument(A) { ava_opaque; }
+
+  ava_argument(B) {
+    ava_out;
+    ava_buffer(rows * cols * elemSize);
+  }
+}
+
+#include "blas1_unimplemented.h"
+#include "blas2_unimplemented.h"
+#include "blas3.h"
+#include "blas3_unimplemented.h"
+#include "blas_like_ext.h"
+#include "blas_like_ext_unimplemented.h"
+#include "cublas_unimplemented.h"
+
+#endif  // _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_H_

cava/samples/cuda_common_spec/cublas/cublas_unimplemented.h

@@ -0,0 +1,64 @@
+#ifndef _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_UNIMPLEMENTED_H_
+#define _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_UNIMPLEMENTED_H_
+#include "cava/nightwatch/parser/c/nightwatch.h"
+/* CUDABLAS API */
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasGetAtomicsMode(cublasHandle_t handle, cublasAtomicsMode_t *mode) {
+  ava_unsupported;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSetAtomicsMode(cublasHandle_t handle, cublasAtomicsMode_t mode) {
+  ava_unsupported;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasLoggerConfigure(int logIsOn, int logToStdOut, int logToStdErr,
+                                                            const char *logFileName) {
+  ava_unsupported;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSetLoggerCallback(cublasLogCallback userCallback) { ava_unsupported; }
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasGetLoggerCallback(cublasLogCallback *userCallback) { ava_unsupported; }
+
+cublasStatus_t CUBLASWINAPI cublasSetVector(int n, int elemSize, const void *x, int incx, void *devicePtr, int incy) {
+  ava_unsupported;
+}
+
+cublasStatus_t CUBLASWINAPI cublasGetVector(int n, int elemSize, const void *x, int incx, void *y, int incy) {
+  ava_unsupported;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasGetMathMode(cublasHandle_t handle, cublasMath_t *mode) { ava_unsupported; }
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSetMathMode(cublasHandle_t handle, cublasMath_t mode) { ava_unsupported; }
+
+cublasStatus_t CUBLASWINAPI cublasSetMatrixAsync(int rows, int cols, int elemSize, const void *A, int lda, void *B,
+                                                 int ldb, cudaStream_t stream) {
+  ava_unsupported;
+}
+
+cublasStatus_t CUBLASWINAPI cublasGetMatrixAsync(int rows, int cols, int elemSize, const void *A, int lda, void *B,
+                                                 int ldb, cudaStream_t stream) {
+  ava_unsupported;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasHgemmBatched(cublasHandle_t handle, cublasOperation_t transa,
+                                                         cublasOperation_t transb, int m, int n, int k,
+                                                         const __half *alpha, /* host or device pointer */
+                                                         const __half *const Aarray[], int lda,
+                                                         const __half *const Barray[], int ldb,
+                                                         const __half *beta, /* host or device pointer */
+                                                         __half *const Carray[], int ldc, int batchCount) {
+  ava_unsupported;
+}
+
+CUBLASAPI cublasStatus_t CUBLASWINAPI cublasHgemmStridedBatched(
+    cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb, int m, int n, int k,
+    const __half *alpha,                                                 /* host or device pointer */
+    const __half *A, int lda, long long int strideA,                     /* purposely signed */
+    const __half *B, int ldb, long long int strideB, const __half *beta, /* host or device pointer */
+    __half *C, int ldc, long long int strideC, int batchCount) {
+  ava_unsupported;
+}
+
+#endif  // _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUBLAS_UNIMPLEMENTED_H_

cava/samples/cuda_common_spec/cudadrv_unimplemented.h

@@ -0,0 +1,43 @@
+#ifndef _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUDADRV_UNIMPLEMENTED_H_
+#define _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUDADRV_UNIMPLEMENTED_H_
+#include <cuda.h>
+
+#include "cava/nightwatch/parser/c/nightwatch.h"
+
+CUresult CUDAAPI cuFuncSetCacheConfig(CUfunction hfunc, CUfunc_cache config) { ava_unsupported; }
+
+CUresult CUDAAPI cuCtxGetSharedMemConfig(CUsharedconfig *pConfig) { ava_unsupported; }
+
+CUresult CUDAAPI cuFuncGetAttribute(int *pi, CUfunction_attribute attrib, CUfunction hfunc) { ava_unsupported; }
+
+CUresult CUDAAPI cuModuleLoadFatBinary(CUmodule *module, const void *fatCubin) { ava_unsupported; }
+
+CUresult CUDAAPI cuStreamAddCallback(CUstream hStream, CUstreamCallback callback, void *userData, unsigned int flags) {
+  ava_unsupported;
+}
+
+CUresult CUDAAPI cuDeviceGetProperties(CUdevprop *prop, CUdevice dev) { ava_unsupported; }
+
+CUresult cuGetExportTable(const void **ppExportTable, const CUuuid *pExportTableId) { ava_unsupported; }
+
+CUresult CUDAAPI cuOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, CUfunction func, int blockSize,
+                                                             size_t dynamicSMemSize) {
+  ava_argument(numBlocks) {
+    ava_out;
+    ava_buffer(1);
+  }
+
+  ava_unsupported;
+}
+
+CUresult CUDAAPI cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, CUfunction func, int blockSize,
+                                                                      size_t dynamicSMemSize, unsigned int flags) {
+  ava_argument(numBlocks) {
+    ava_out;
+    ava_buffer(1);
+  }
+
+  ava_unsupported;
+}
+
+#endif  // _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUDADRV_UNIMPLEMENTED_H_

cava/samples/cuda_common_spec/cudart.h

@@ -0,0 +1,63 @@
+#ifndef _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUDART_H_
+#define _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUDART_H_
+#include <cuda_runtime_api.h>
+
+#include "cava/nightwatch/parser/c/nightwatch.h"
+
+__host__ __cudart_builtin__ const char *CUDARTAPI cudaGetErrorName(cudaError_t error) {
+  const char *ret = reinterpret_cast<const char *>(ava_execute());
+  ava_return_value {
+    ava_out;
+    ava_buffer(strlen(ret) + 1);
+    ava_lifetime_static;
+  }
+}
+
+__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) {
+  ava_argument(pStream) {
+    ava_out;
+    ava_buffer(1);
+    ava_element ava_handle;
+  }
+}
+
+__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) {
+  ava_argument(stream) ava_handle;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { ava_argument(stream) ava_handle; }
+
+__host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) { ava_argument(stream) ava_handle; }
+
+__host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, unsigned int flags) {
+  ava_argument(event) {
+    ava_out;
+    ava_buffer(1);
+    ava_element ava_handle;
+  }
+}
+
+__host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event) { ava_argument(event) ava_handle; }
+
+__host__ cudaError_t CUDARTAPI cudaMemGetInfo(size_t *_free, size_t *total) {
+  ava_argument(_free) {
+    ava_out;
+    ava_buffer(1);
+  }
+  ava_argument(total) {
+    ava_out;
+    ava_buffer(1);
+  }
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset __dv(0),
+                                                    enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) {
+  /* kind is always cudaMemcpyDeviceToHost */
+  ava_argument(dst) {
+    ava_out;
+    ava_buffer(count);
+  }
+  ava_argument(symbol) ava_opaque;
+}
+
+#endif  // _AVA_CAVA_SAMPLES_CUDA_COMMON_SPEC_CUDART_H_