[E2E] Add Graphs case

features/config/TEMPLATE_graph.xml

@@ -0,0 +1,13 @@
+<?xml version="1.0" encoding="UTF-8"?>
+
+<test driverID="test_feature" name="TEMPLATE">
+    <description>test</description>
+    <files>
+        <file path="feature_case/graph/${testName}.cu" />
+    </files>
+    <rules>
+        <optlevelRule excludeOptlevelNameString="cpu" />
+	<optlevelRule excludeOptlevelNameString="usmnone" />
+	<optlevelRule excludeOptlevelNameString="syclcompat" />
+    </rules>
+</test>

features/feature_case/graph/graph.cu

@@ -0,0 +1,70 @@
+// ===------- graph.cu ------------------------------------ *- CUDA -* ---===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <stdio.h>
+
+const int blockSize = 256;
+const int numBlocks = (10 + blockSize - 1) / blockSize;
+
+__global__ void init(float *a) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id < 10) {
+    a[id] = 1.0f;
+  }
+}
+
+__global__ void incrementA(float *a) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id < 10) {
+    a[id] += 1.0f;
+  }
+}
+
+int main() {
+
+  cudaGraph_t graph;
+
+  cudaStream_t stream;
+
+  cudaStreamCreate(&stream);
+
+  float *d_a, h_a[10];
+
+  cudaMalloc(&d_a, 10 * sizeof(float));
+
+  init<<<numBlocks, blockSize, 0, stream>>>(d_a);
+
+  cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
+
+  incrementA<<<numBlocks, blockSize, 0, stream>>>(d_a);
+
+  cudaStreamEndCapture(stream, &graph);
+  cudaGraphExec_t execGraph;
+  cudaGraphInstantiate(&execGraph, graph, NULL, NULL, 0);
+
+  cudaGraphLaunch(execGraph, stream);
+
+  cudaStreamSynchronize(stream); // Ensure the graph has completed execution
+
+  cudaMemcpy(h_a, d_a, 10 * sizeof(float), cudaMemcpyDeviceToHost);
+
+  for (int i = 0; i < 10; i++) {
+    if (h_a[i] != 2.0f) {
+      printf("Results do not match\n");
+      return -1;
+    }
+  }
+
+  printf("Passed\n");
+
+  cudaStreamDestroy(stream);
+  cudaFree(d_a);
+  cudaGraphExecDestroy(execGraph);
+
+  return 0;
+}

features/features.xml

@@ -356,6 +356,7 @@
     <test testName="matmul" configFile="config/TEMPLATE_cublasLt.xml" />
     <test testName="transform" configFile="config/TEMPLATE_cublasLt.xml" />
     <test testName="graphics_interop_d3d11" configFile="config/TEMPLATE_cudaGraphics_Interop.xml" />
+    <test testName="graph" configFile="config/TEMPLATE_graph.xml" />
     <test testName="peer_access_using_driver_api" configFile="config/TEMPLATE_peer_access_using_driver_api.xml" />
     <test testName="context_push_n_pop" configFile="config/TEMPLATE_context_push_n_pop.xml" />
   </tests>

features/test_feature.py

@@ -61,8 +61,8 @@
               'thrust_swap_ranges', 'thrust_uninitialized_fill_n', 'thrust_equal', 'system_atomic', 'thrust_detail_types',
               'operator_eq', 'operator_neq', 'operator_lege', 'thrust_system', 'thrust_reverse_copy',
               'thrust_device_new_delete', 'thrust_temporary_buffer', 'thrust_malloc_free', 'codepin', 'thrust_unique_count',
-              'thrust_advance_trans_op_itr', 'cuda_stream_query', "matmul", "transform", "context_push_n_pop",
-              "graphics_interop_d3d11"]
+              'thrust_advance_trans_op_itr', 'cuda_stream_query', "matmul", "transform",  "context_push_n_pop",
+              "graphics_interop_d3d11", 'graph']
 
 occupancy_calculation_exper = ['occupancy_calculation']
 
@@ -126,6 +126,8 @@ def migrate_test():
         src.append(' --use-experimental-features=bindless_images')
     if "codepin" in test_config.current_test:
         src.append(' --enable-codepin ')
+    if test_config.current_test == 'graph':
+        src.append(' --use-experimental-features=graph ')
     return do_migrate(src, in_root, test_config.out_root, extra_args)
 
 def manual_fix_for_cufft_external_workspace(migrated_file):