Add 3D Convolution and Attention Benchmarks

benchmarks/hecbench/hip/attention/Makefile

@@ -0,0 +1,75 @@
+#===============================================================================
+# User Options
+#===============================================================================
+
+# Compiler can be set below, or via environment variable
+CC        = ${PROTEUS_CC}
+OPTIMIZE  = yes
+DEBUG     = no
+PROTEUS_PATH  ?= /path/to/proteus/install
+ENABLE_PROTEUS ?= yes
+LAUNCHER  =
+
+#===============================================================================
+# Program name & source code list
+#===============================================================================
+
+ifeq ($(ENABLE_PROTEUS),yes)
+  SUFFIX = -proteus
+else
+  SUFFIX = -aot
+endif
+
+program = attention$(SUFFIX).x
+
+source = main.cu
+obj = $(source:.cu=$(SUFFIX).o)
+
+#===============================================================================
+# Sets Flags
+#===============================================================================
+
+# Standard Flags
+CFLAGS := $(EXTRA_CFLAGS) -std=c++14 -Wall
+
+# Linker Flags
+LDFLAGS =
+
+# Debug Flags
+ifeq ($(DEBUG),yes)
+	CFLAGS += -g
+	LDFLAGS  += -g
+endif
+
+# Optimization Flags
+ifeq ($(OPTIMIZE),yes)
+	CFLAGS += -O3
+endif
+
+ifeq ($(ENABLE_PROTEUS),yes)
+	CFLAGS += -fpass-plugin=${PROTEUS_PATH}/lib64/libProteusPass.so -DENABLE_PROTEUS
+	LDFLAGS += -Wl,-rpath,${PROTEUS_PATH}/lib64 -L${PROTEUS_PATH}/lib64/ -lproteus \
+	          -L${ROCM_PATH}/lib -L${ROCM_PATH}/llvm/lib \
+	          -Wl,--start-group \
+	          $(shell ls ${ROCM_PATH}/llvm/lib/libclang*.a) \
+	          $(shell ${ROCM_PATH}/llvm/bin/llvm-config --libs) \
+	          -Wl,--end-group \
+	          $(shell ${ROCM_PATH}/llvm/bin/llvm-config --system-libs) \
+	          -llldCommon -llldELF
+endif
+
+#===============================================================================
+# Targets to Build
+#===============================================================================
+
+$(program): $(obj) Makefile
+	$(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS)
+
+%$(SUFFIX).o: %.cu kernels.h Makefile
+	$(CC) $(CFLAGS) -x hip -c $< -o $@
+
+clean:
+	rm -rf *.x *.o *.ll *.bc .proteus
+
+run: $(program)
+	$(LAUNCHER) ./$(program) 65536 2048 100

benchmarks/hecbench/hip/attention/kernels.h

@@ -0,0 +1,56 @@
+#ifdef ENABLE_PROTEUS
+__attribute__((annotate("jit", 5, 6)))
+#endif
+__global__
+void attention_kernel1 (
+    const float*__restrict__ key,
+    const float*__restrict__ query,
+    float*__restrict__ dot_product,
+    float*__restrict__ exp_sum,
+    const int n,
+    const int d)
+{
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n) {
+    float sum = 0;
+    for (int j = 0; j < d; j++)
+      sum += key[i * d + j] * query[j];
+    dot_product[i] = sum;
+    atomicAdd(exp_sum, __expf(sum));
+  }
+}
+
+#ifdef ENABLE_PROTEUS
+__attribute__((annotate("jit", 4)))
+#endif
+__global__
+void attention_kernel2 (
+    const float*__restrict__ exp_sum,
+    const float*__restrict__ dot_product,
+    float*__restrict__ score,
+    const int n)
+{
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n)
+    score[i] = __expf(dot_product[i]) / exp_sum[0];
+}
+
+#ifdef ENABLE_PROTEUS
+__attribute__((annotate("jit", 4, 5)))
+#endif
+__global__
+void attention_kernel3 (
+    const float*__restrict__ score,
+    const float*__restrict__ value,
+    float*__restrict__ output,
+    const int n,
+    const int d)
+{
+  int j = blockIdx.x * blockDim.x + threadIdx.x;
+  if (j < d) {
+    float sum = 0;
+    for (int i = 0; i < n; i++)
+      sum += score[i] * value[i * d + j];
+    output[j] = sum;
+  }
+}

benchmarks/hecbench/hip/attention/main.cu

@@ -0,0 +1,152 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include <math.h>
+#include <chrono>
+#include <random>
+#include <hip/hip_runtime.h>
+#include "kernels.h"
+
+float* attention_host(const float* key, const float* value, const float* query,
+  const int n, const int d) 
+{
+// intermediate
+float* dot_product = (float*) malloc (n * sizeof(float));
+float* score = (float*) malloc (n * sizeof(float));
+// result
+float* output = (float*) malloc (d * sizeof(float));
+
+for (int i = 0; i < n; i++) {
+float sum = 0;
+for (int j = 0; j < d; j++)
+sum += key[i * d + j] * query[j];
+dot_product[i] = sum;
+}
+
+float sum = 0;
+for (int i = 0; i < n; i++)
+sum += expf(dot_product[i]);
+
+for (int i = 0; i < n; i++)
+score[i] = expf(dot_product[i]) / sum;
+
+for (int j = 0; j < d; j++) {
+float sum = 0;
+for (int i = 0; i < n; i++)
+sum += score[i] * value[i * d + j];
+output[j] = sum;
+}
+
+free(dot_product);
+free(score);
+return output;
+}
+
+
+
+
+float* attention_device(const float* key, const float* value, const float* query,
+                        const int n, const int d, const int repeat, const int verify)
+{
+  // input
+  float *d_key;
+  hipMalloc((void**)&d_key, n * d * sizeof(float));
+  hipMemcpy(d_key, key, n * d * sizeof(float), hipMemcpyHostToDevice);
+
+  float *d_value;
+  hipMalloc((void**)&d_value, n * d * sizeof(float));
+  hipMemcpy(d_value, value, n * d * sizeof(float), hipMemcpyHostToDevice);
+
+  float *d_query;
+  hipMalloc((void**)&d_query, d * sizeof(float));
+  hipMemcpy(d_query, query, d * sizeof(float), hipMemcpyHostToDevice);
+
+  // intermediate
+  float *d_dot_product;
+  hipMalloc((void**)&d_dot_product, n * sizeof(float));
+
+  float *d_exp_sum;
+  hipMalloc((void**)&d_exp_sum, sizeof(float));
+
+  // result
+  float *output = (float*) malloc (d * sizeof(float));
+  float *d_output;
+  hipMalloc((void**)&d_output, d * sizeof(float));
+
+  float *d_score;
+  hipMalloc((void**)&d_score, n * sizeof(float));
+
+  hipDeviceSynchronize();
+
+  auto start = std::chrono::steady_clock::now();
+
+  for (int k = 0; k < repeat; k++) {
+    if(verify) {
+      hipMemset(d_exp_sum, 0, 4);
+    }
+    attention_kernel1<<<(n+255)/256, 256>>>(d_key, d_query, d_dot_product, d_exp_sum, n, d);
+    attention_kernel2<<<(n+255)/256, 256>>>(d_exp_sum, d_dot_product, d_score, n);
+    attention_kernel3<<<(d+255)/256, 256>>>(d_score, d_value, d_output, n, d);
+  }
+
+  hipDeviceSynchronize();
+  auto end = std::chrono::steady_clock::now();
+  auto time = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start).count();
+  printf("Average execution time of kernels %f (ms)\n", time * 1e-6f / repeat);
+
+  hipMemcpy(output, d_output, d * sizeof(float), hipMemcpyDeviceToHost);
+  hipFree(d_score);
+  hipFree(d_value);
+  hipFree(d_output);
+  hipFree(d_key);
+  hipFree(d_dot_product);
+  hipFree(d_exp_sum);
+  hipFree(d_query);
+  return output;
+}
+
+int main(int argc, char* argv[]) {
+  if (argc != 4 && argc != 5) {
+    printf("Usage: %s <rows> <columns> <repeat> [verify]\n", argv[0]);
+    return 1;
+  }
+  const int n = atoi(argv[1]);
+  const int d = atoi(argv[2]);
+  const int r = atoi(argv[3]);
+  const int verify = (argc == 5) ? atoi(argv[4]) : 0;
+
+  // input
+  float* key = (float*) malloc (n * d * sizeof(float));
+  float* value = (float*) malloc (n * d * sizeof(float));
+  float* query = (float*) malloc (d * sizeof(float));
+
+  std::mt19937 gen(19937);
+  std::uniform_real_distribution<float> dist(-0.01f, 0.01f);
+
+  if (verify) {
+  for (int i = 0; i < n * d; i++) {
+    key[i] = dist(gen);
+    value[i] = dist(gen);
+      query[i % d] = dist(gen);
+    }
+  }
+
+  float* dout = attention_device(key, value, query, n, d, r, verify);
+
+  if (verify) {
+    float* hout = attention_host(key, value, query, n, d);
+
+    float rmse = 0;
+    for (int i = 0; i < d; i++) {
+      rmse += (hout[i] - dout[i]) * (hout[i] - dout[i]);
+    }
+    printf("RMSE = %f\n", sqrtf(rmse / d));
+
+    free(hout);
+  }
+
+  free(key);
+  free(value);
+  free(query);
+  free(dout);
+  return 0;
+}

benchmarks/hecbench/hip/conv3d/Makefile

@@ -0,0 +1,76 @@
+#===============================================================================
+# User Options
+#===============================================================================
+
+# Compiler can be set below, or via environment variable
+CC        = ${PROTEUS_CC}
+OPTIMIZE  = yes
+DEBUG     = no
+LAUNCHER  =
+PROTEUS_PATH  ?=/path/to/proteus/install
+ENABLE_PROTEUS ?= no
+
+#===============================================================================
+# Program name & source code list
+#===============================================================================
+
+ifeq ($(ENABLE_PROTEUS),yes)
+  SUFFIX = "-proteus"
+else
+  SUFFIX = "-aot"
+endif
+
+program = conv3d$(SUFFIX).x
+
+source = main.cu
+obj = $(source:.cu=$(SUFFIX).o)
+
+#===============================================================================
+# Sets Flags
+#===============================================================================
+
+# Standard Flags
+CFLAGS := $(EXTRA_CFLAGS) -std=c++14 -Wall
+
+# Linker Flags
+LDFLAGS =
+
+# Debug Flags
+ifeq ($(DEBUG),yes)
+	CFLAGS += -g
+	LDFLAGS  += -g
+endif
+
+# Optimization Flags
+ifeq ($(OPTIMIZE),yes)
+	CFLAGS += -O3
+endif
+
+# Proteus
+ifeq ($(ENABLE_PROTEUS),yes)
+  CFLAGS += -fpass-plugin=${PROTEUS_PATH}/lib64/libProteusPass.so -DENABLE_PROTEUS
+  LDFLAGS += -Wl,-rpath,${PROTEUS_PATH}/lib64 -L${PROTEUS_PATH}/lib64/ -lproteus \
+          -L${ROCM_PATH}/lib -L${ROCM_PATH}/llvm/lib \
+          -Wl,--start-group \
+          $(shell ls ${ROCM_PATH}/llvm/lib/libclang*.a) \
+          $(shell ${ROCM_PATH}/llvm/bin/llvm-config --libs) \
+          -Wl,--end-group \
+          $(shell ${ROCM_PATH}/llvm/bin/llvm-config --system-libs) \
+          -llldCommon -llldELF
+endif
+
+#===============================================================================
+# Targets to Build
+#===============================================================================
+
+$(program): $(obj) Makefile
+	$(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS)
+
+%$(SUFFIX).o: %.cu Makefile
+	$(CC) $(CFLAGS) -x hip -c $< -o $@
+
+clean:
+	rm -rf *.x *.o *.ll *.bc .proteus
+
+run: $(program)
+	$(LAUNCHER) ./$(program) 32 96 256 26 26 5 100