Start to experiment with bindings for cuda_kernel

Author: caugonnet

c/experimental/stf/include/cccl/c/experimental/stf/stf.h

@@ -44,17 +44,14 @@ cudaStream_t stf_task_get_stream(stf_task_handle t);
 void* stf_task_get(stf_task_handle t, size_t submitted_index);
 void stf_task_destroy(stf_task_handle t);
 
-typedef struct stf_kernel_desc_handle_t* stf_kernel_desc_handle;
-
-void stf_kernel_create(stf_kernel_desc_handle* d);
-void stf_kernel_destroy(stf_kernel_desc_handle d);
-// TODO stf_cuda_kernel_desc : symbol, deps, args... ?
-// void stf_kernel_set_symbol((stf_kernel_handle k, const char* symbol)
-// void stf_kernel_add_dep(stf_kernel_handle k, stf_logical_data_handle ld, stf_access_mode m);
-// void stf_kernel_start(stf_kernel_handle k);
-// void stf_kernel_set_args(stf_kernel_handle k, size_t cnt, void **args);
-// void stf_kernel_end(stf_kernel_handle k);
-// void stf_kernel_destroy(stf_kernel_handle k);
+typedef struct stf_cuda_kernel_handle_t* stf_cuda_kernel_handle;
+
+void stf_cuda_kernel_create(stf_ctx_handle ctx, stf_cuda_kernel_handle* k);
+void stf_cuda_kernel_set_symbol(stf_cuda_kernel_handle k, const char* symbol);
+void stf_cuda_kernel_add_dep(stf_cuda_kernel_handle k, stf_logical_data_handle ld, stf_access_mode m);
+void stf_cuda_kernel_start(stf_cuda_kernel_handle k);
+void stf_cuda_kernel_end(stf_cuda_kernel_handle k);
+void stf_cuda_kernel_destroy(stf_cuda_kernel_handle t);
 
 #ifdef __cplusplus
 }

c/experimental/stf/src/stf.cu

@@ -153,20 +153,20 @@ void stf_cuda_kernel_add_dep(stf_cuda_kernel_handle k, stf_logical_data_handle l
   assert(k);
   assert(ld);
 
-  k->k.add_deps(cuda_kernel_dep_untyped(ld->ld, access_mode(m)));
-}
-
-// void stf_cuda_kernel_start(stf_cuda_kernel_handle k)
-// {
-//   assert(k);
-//   k->k.start();
-// }
-//
-// void stf_cuda_kernel_end(stf_cuda_kernel_handle k)
-// {
-//   assert(k);
-//   k->k.end();
-// }
+  k->k.add_deps(task_dep_untyped(ld->ld, access_mode(m)));
+}
+
+void stf_cuda_kernel_start(stf_cuda_kernel_handle k)
+{
+  assert(k);
+  k->k.start();
+}
+
+void stf_cuda_kernel_end(stf_cuda_kernel_handle k)
+{
+  assert(k);
+  k->k.end();
+}
 
 void stf_cuda_kernel_destroy(stf_cuda_kernel_handle t)
 {

c/experimental/stf/test/test_cuda_kernel.cpp

@@ -0,0 +1,65 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// 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
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda_runtime.h>
+
+#include <c2h/catch2_test_helper.h>
+#include <cccl/c/experimental/stf/stf.h>
+
+using namespace cuda::experimental::stf;
+
+__global__ void axpy(int cnt, double a, const double *x, double *y)
+{
+  int tid      = blockIdx.x * blockDim.x + threadIdx.x;
+  int nthreads = gridDim.x * blockDim.x;
+
+  for (int i = tid; i < cnt; i += nthreads)
+  {
+    y[i] += a * x[i];
+  }
+}
+
+C2H_TEST("axpy with stf cuda_kernel", "[cuda_kernel]")
+{
+  size_t N = 1000000;
+
+  stf_ctx_handle ctx;
+  stf_ctx_create(&ctx);
+
+  stf_logical_data_handle lX, lY;
+
+  float *X, *Y;
+  X = (float*) malloc(N * sizeof(float));
+  Y = (float*) malloc(N * sizeof(float));
+
+  stf_logical_data(ctx, &lX, X, N * sizeof(float));
+  stf_logical_data(ctx, &lY, Y, N * sizeof(float));
+
+  stf_logical_data_set_symbol(lX, "X");
+  stf_logical_data_set_symbol(lY, "Y");
+
+  stf_cuda_kernel_handle k;
+  stf_cuda_kernel_create(ctx, &k);
+  stf_cuda_kernel_set_symbol(k, "axpy");
+  stf_cuda_kernel_add_dep(k, lX, STF_READ);
+  stf_cuda_kernel_add_dep(k, lY, STF_RW);
+  stf_cuda_kernel_start(k);
+  // TODO add descs
+  stf_cuda_kernel_end(k);
+  stf_cuda_kernel_destroy(k);
+
+  stf_logical_data_destroy(lX);
+  stf_logical_data_destroy(lY);
+
+  stf_ctx_finalize(ctx);
+
+  free(X);
+  free(Y);
+}