diff --git a/benchmarks/hecbench/hip/attention/Makefile b/benchmarks/hecbench/hip/attention/Makefile
new file mode 100644
index 0000000..4afbeac
--- /dev/null
+++ b/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
diff --git a/benchmarks/hecbench/hip/attention/kernels.h b/benchmarks/hecbench/hip/attention/kernels.h
new file mode 100644
index 0000000..76ee2c1
--- /dev/null
+++ b/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;
+  }
+}
diff --git a/benchmarks/hecbench/hip/attention/main.cu b/benchmarks/hecbench/hip/attention/main.cu
new file mode 100644
index 0000000..edcbda5
--- /dev/null
+++ b/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;
+}
diff --git a/benchmarks/hecbench/hip/conv3d/Makefile b/benchmarks/hecbench/hip/conv3d/Makefile
new file mode 100644
index 0000000..2570b18
--- /dev/null
+++ b/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
diff --git a/benchmarks/hecbench/hip/conv3d/main.cu b/benchmarks/hecbench/hip/conv3d/main.cu
new file mode 100644
index 0000000..d486133
--- /dev/null
+++ b/benchmarks/hecbench/hip/conv3d/main.cu
@@ -0,0 +1,303 @@
+/*
+  Reference
+  Chapter 16 in Programming massively parallel processors,
+  A hands-on approach (D. Kirk and W. Hwu)
+ */
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <chrono>
+#include <iostream>
+#include <hip/hip_runtime.h>
+
+
+#define TILE_WIDTH 16
+
+#define II(n,c,h,w) ((n)*C*Hin*Win+(c)*Hin*Win+(h)*Win+w)
+#define WI(n,c,h,w) ((n)*C*K*K+(c)*K*K+(h)*K+w)
+#define OI(n,c,h,w) ((n)*M*Hout*Wout+(c)*Hout*Wout+(h)*Wout+w)
+
+void verify (const float* Y, float* Y_ref, size_t Y_size)
+{
+  bool ok = true;
+  for (size_t i = 0; i < Y_size; i++) {
+    if (fabs(Y[i] - Y_ref[i]) > 1e-3f) {
+      printf("%f (device) != %f (reference)\n", Y[i], Y_ref[i]);
+      ok = false;
+      break;
+    }
+  }
+  printf("%s\n", ok ? "PASS" : "FAIL");
+}
+
+#ifdef ENABLE_PROTEUS
+__global__ __attribute__((annotate("jit", 4, 5, 6, 7, 8, 9, 10, 11)))
+#else
+__global__
+#endif
+void conv3d_s1(const float * __restrict__ X,
+               const float * __restrict__ W,
+                     float * __restrict__ Y,
+               const int C,
+               const int M,
+               const int K,
+               const int Hin,
+               const int Win,
+               const int Hout,
+               const int Wout,
+               const int W_grid)
+{
+  int n = blockIdx.x;
+  int m = blockIdx.y;
+  int h = blockIdx.z / W_grid * TILE_WIDTH + threadIdx.y;
+  int w = blockIdx.z % W_grid * TILE_WIDTH + threadIdx.x;
+  if (h < Hout && w < Wout) {
+    float s = 0;
+    for (int c = 0; c < C; c++) {
+      for (int p = 0; p < K; p++) {
+        for (int q = 0; q < K; q++) {
+          s += X[II(n, c, h+p, w+q)] * W[WI(m, c, p, q)];
+        }
+      }
+    }
+    Y[OI(n, m, h, w)] = s;
+  }
+}
+
+#ifdef ENABLE_PROTEUS
+__global__ __attribute__((annotate("jit", 4, 5, 6, 7, 8, 9, 10, 11)))
+#else
+__global__
+#endif
+void conv3d_s2(const float * __restrict__ X,
+               const float * __restrict__ W,
+                     float * __restrict__ Y,
+               const int C,
+               const int M,
+               const int K,
+               const int Hin,
+               const int Win,
+               const int Hout,
+               const int Wout,
+               const int W_grid)
+{
+  int m = blockIdx.x;
+  int h = blockIdx.y / W_grid * TILE_WIDTH + threadIdx.y;
+  int w = blockIdx.y % W_grid * TILE_WIDTH + threadIdx.x;
+  int n = blockIdx.z;
+  if (h < Hout && w < Wout) {
+    float s = 0;
+    for (int c = 0; c < C; c++) {
+      for (int p = 0; p < K; p++) {
+        for (int q = 0; q < K; q++) {
+          s += X[II(n, c, h+p, w+q)] * W[WI(m, c, p, q)];
+        }
+      }
+    }
+    Y[OI(n, m, h, w)] = s;
+  }
+}
+
+#ifdef ENABLE_PROTEUS
+__global__ __attribute__((annotate("jit", 4, 5, 6, 7, 8, 9, 10, 11)))
+#else
+__global__
+#endif
+void conv3d_s3(const float * __restrict__ X,
+               const float * __restrict__ W,
+                     float * __restrict__ Y,
+               const int C,
+               const int M,
+               const int K,
+               const int Hin,
+               const int Win,
+               const int Hout,
+               const int Wout,
+               const int W_grid)
+{
+  int h = blockIdx.x / W_grid * TILE_WIDTH + threadIdx.y;
+  int w = blockIdx.x % W_grid * TILE_WIDTH + threadIdx.x;
+  int n = blockIdx.y;
+  int m = blockIdx.z;
+  if (h < Hout && w < Wout) {
+    float s = 0;
+    for (int c = 0; c < C; c++) {
+      for (int p = 0; p < K; p++) {
+        for (int q = 0; q < K; q++) {
+          s += X[II(n, c, h+p, w+q)] * W[WI(m, c, p, q)];
+        }
+      }
+    }
+    Y[OI(n, m, h, w)] = s;
+  }
+}
+
+
+// Hin = Hout-1+K; max(h+p) is Hin - 1 as max(h) = Hout-1 and max(p) = K-1
+void reference(const float * __restrict__ X,
+               const float * __restrict__ W,
+                     float * __restrict__ Y,
+               const int N,
+               const int M,
+               const int C,
+               const int K,
+               const int Hin,
+               const int Win,
+               const int Hout,
+               const int Wout)
+{
+  for(int n = 0; n < N; n++)
+    for(int m = 0; m < M; m++)
+      for(int h = 0; h < Hout; h++)
+        for(int w = 0; w < Wout; w++) {
+          Y[OI(n, m, h, w)] = 0;
+          for(int c = 0; c < C; c++)
+            for(int p = 0; p < K; p++)
+              for(int q = 0; q < K; q++)
+                Y[OI(n, m, h, w)] += X[II(n, c, h+p, w+q)] * W[WI(m, c, p, q)];
+        }
+}
+
+void conv3D(const int N, const int C, const int M, const int Win, const int Hin, const int K, const int repeat, const int do_verify)
+{
+  const int Hout = Hin-K+1;
+  const int Wout = Win-K+1;
+
+  size_t X_size = N * C * Hin * Win;
+  size_t W_size = M * C * K * K;
+  size_t Y_size = N * M * Hout * Wout;
+  size_t X_bytes = X_size * sizeof(float);
+  size_t W_bytes = W_size * sizeof(float);
+  size_t Y_bytes = Y_size * sizeof(float);
+
+  float *X, *W, *Y, *Y_ref;
+  X = (float *)malloc(X_bytes); // input
+  W = (float *)malloc(W_bytes); // filter
+  Y = (float *)malloc(Y_bytes); // output
+
+  srand(123);
+
+
+    for (size_t i = 0; i < W_size; i++) W[i] = rand() % 31;
+    for (size_t i = 0; i < X_size; i++) X[i] = rand() % 13;
+
+    for (size_t i = 0; i < Y_size; i++) {
+      Y[i] = -1;
+    }
+
+  if (do_verify) {
+
+    Y_ref = (float *)malloc(Y_bytes);
+    for (size_t i = 0; i < Y_size; i++) {
+      Y_ref[i] = -1;
+    }
+    reference(X, W, Y_ref, N, M, C, K, Hin, Win, Hout, Wout);
+  }
+
+  float *dX, *dW, *dY;
+  hipMalloc((void **)&dX, X_bytes);
+  hipMalloc((void **)&dW, W_bytes);
+  hipMalloc((void **)&dY, Y_bytes);
+
+  hipMemcpy(dX, X, X_bytes, hipMemcpyHostToDevice);
+  hipMemcpy(dW, W, W_bytes, hipMemcpyHostToDevice);
+  hipMemcpy(dY, Y, Y_bytes, hipMemcpyHostToDevice);
+
+  int W_grid = (Wout + TILE_WIDTH - 1) / TILE_WIDTH;
+  int H_grid = (Hout + TILE_WIDTH - 1) / TILE_WIDTH;
+  int Z = H_grid * W_grid;
+
+  printf("input dimensions: C=%d Win=%d Hin=%d\n", C, Win, Hin);
+  printf("output dimensions: M=%d Wout=%d Hout=%d\n", M, Wout, Hout);
+  printf("3D grid dimensions: N=%d M=%d Z=%d\n", N, M, Z);
+
+  // try grid organizations
+  dim3 grids_s1 (N, M, Z);
+  dim3 grids_s2 (M, Z, N);
+  dim3 grids_s3 (Z, N, M);
+  dim3 blocks (TILE_WIDTH, TILE_WIDTH, 1);
+
+  hipDeviceSynchronize();
+
+  auto start = std::chrono::steady_clock::now();
+  for (int i = 0; i < repeat; i++) {
+    conv3d_s1 <<< grids_s1, blocks >>> (dX, dW, dY, C, M, K, Hin, Win, Hout, Wout, W_grid);
+  }
+
+  hipDeviceSynchronize();
+  auto end = std::chrono::steady_clock::now();
+  auto time = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start).count();
+  printf("Average kernel execution time of conv3d_s1 kernel: %f (us)\n",
+         (time * 1e-3f) / repeat);
+  if (do_verify) {
+    hipMemcpy(Y, dY, Y_bytes, hipMemcpyDeviceToHost);
+    verify(Y, Y_ref, Y_size);
+  }
+
+  start = std::chrono::steady_clock::now();
+  for (int i = 0; i < repeat; i++) {
+    conv3d_s2 <<< grids_s2, blocks >>> (dX, dW, dY, C, M, K, Hin, Win, Hout, Wout, W_grid);
+  }
+
+  hipDeviceSynchronize();
+  end = std::chrono::steady_clock::now();
+  time = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start).count();
+  printf("Average kernel execution time of conv3d_s2 kernel: %f (us)\n",
+         (time * 1e-3f) / repeat);
+  if (do_verify) {
+    hipMemcpy(Y, dY, Y_bytes, hipMemcpyDeviceToHost);
+    verify(Y, Y_ref, Y_size);
+  }
+
+  start = std::chrono::steady_clock::now();
+  for (int i = 0; i < repeat; i++) {
+    conv3d_s3 <<< grids_s3, blocks >>> (dX, dW, dY, C, M, K, Hin, Win, Hout, Wout, W_grid);
+  }
+
+  hipDeviceSynchronize();
+  end = std::chrono::steady_clock::now();
+  time = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start).count();
+  printf("Average kernel execution time of conv3d_s3 kernel: %f (us)\n",
+         (time * 1e-3f) / repeat);
+  if (do_verify) {
+    hipMemcpy(Y, dY, Y_bytes, hipMemcpyDeviceToHost);
+    verify(Y, Y_ref, Y_size);
+  }
+
+  free(X);
+  free(W);
+  free(Y);
+  if (do_verify) {
+    free(Y_ref);
+  }
+  hipFree(dX);
+  hipFree(dW);
+  hipFree(dY);
+}
+
+int main(int argc, char* argv[]) {
+  if (argc != 8 && argc != 9) {
+    printf("Usage: %s <batch size:N> <input channels:C> <output feature maps:M>", argv[0]);
+    printf(" <input width:Win> <input height:Hin> <kernel size:K> <repeat> [verify (0 or 1, default 0)]\n");
+    return 1;
+  }
+
+  int N = atoi(argv[1]);
+  int C = atoi(argv[2]);
+  int M = atoi(argv[3]);
+  int W = atoi(argv[4]);
+  int H = atoi(argv[5]);
+  int K = atoi(argv[6]);
+  int repeat = atoi(argv[7]);
+  int verify = (argc == 9) ? atoi(argv[8]) : 0;
+
+  printf("3D convolution (FP32)\n");
+  printf("\n========== Warmup start ==========\n");
+  conv3D(N, C, M, W, H, K, 1000, verify);
+  printf("\n========== Warmup done ==========\n");
+  conv3D(N, C, M, W, H, K, repeat, verify);
+
+  return 0;
+}
diff --git a/hecbench.toml b/hecbench.toml
index 66a7710..7006eab 100644
--- a/hecbench.toml
+++ b/hecbench.toml
@@ -127,3 +127,23 @@ path = "benchmarks/hecbench/hip/wsm5"
 exe = "wsm5-proteus.x"
 [hecbench.wsm5.inputs]
 default = "10"
+
+[hecbench.conv3d]
+[hecbench.conv3d.amd.aot]
+path = "benchmarks/hecbench/hip/conv3d"
+exe = "conv3d-aot.x"
+[hecbench.conv3d.amd.proteus]
+path = "benchmarks/hecbench/hip/conv3d"
+exe = "conv3d-proteus.x"
+[hecbench.conv3d.inputs]
+default = "32 96 256 26 26 5 100"
+
+[hecbench.attention]
+[hecbench.attention.amd.aot]
+path = "benchmarks/hecbench/hip/attention"
+exe = "attention-aot.x"
+[hecbench.attention.amd.proteus]
+path = "benchmarks/hecbench/hip/attention"
+exe = "attention-proteus.x"
+[hecbench.attention.inputs]
+default = "65536 2048 100"