ci: fix env var conflict, use clEnqueueWriteBuffer

libpocl is broken on github action cpus, use clEnqueueWriteBuffer
instead of compiling a kernel

.github/workflows/presubmit.yml

@@ -112,7 +112,7 @@ jobs:
     env:
       THAPI_TEST_BIN: clinfo
       THAPI_BIN_DIR: ${{ github.workspace }}/build/ici/bin
-      THAPI_HOME: ${{ github.workspace }}/build/ici
+      THAPI_INSTALL_DIR: ${{ github.workspace }}/build/ici
       THAPI_SYNC_DAEMON: ${{ matrix.thapi_sync_daemon }}
     steps:
       - uses: actions/checkout@v4

integration_tests/libthapi-ctl.bats

@@ -1,7 +1,7 @@
 #!/usr/bin/env bats
 
 setup_file() {
-   export IPROF=$THAPI_HOME/bin/iprof
+   export IPROF=$THAPI_INSTALL_DIR/bin/iprof
 
    (cd integration_tests/libthapi-ctl;
     ld --verbose | grep SEARCH_DIR | tr -s ' ;' \\012;
@@ -11,7 +11,7 @@ setup_file() {
     objdump -p hello-cl-startstop)
 
    export THAPI_TEST_BIN=$(pwd)/integration_tests/libthapi-ctl/hello-cl-startstop
-   export LD_LIBRARY_PATH=$THAPI_HOME/lib:$LD_LIBRARY_PATH
+   export LD_LIBRARY_PATH=$THAPI_INSTALL_DIR/lib:$LD_LIBRARY_PATH
 }
 
 teardown_file() {

integration_tests/libthapi-ctl/Makefile

@@ -1,5 +1,5 @@
-CFLAGS = -Wall -Wextra $(WERROR) -DCL_TARGET_OPENCL_VERSION=220 -DCL_USE_DEPRECATED_OPENCL_1_2_APIS -I$(THAPI_HOME)/include
-LDFLAGS = -lOpenCL -L$(THAPI_HOME)/lib -lthapi-ctl -llttng-ctl
+CFLAGS = -Wall -Wextra $(WERROR) -DCL_TARGET_OPENCL_VERSION=220 -DCL_USE_DEPRECATED_OPENCL_1_2_APIS -I$(THAPI_INSTALL_DIR)/include
+LDFLAGS = /usr/lib/x86_64-linux-gnu/libOpenCL.so -L$(THAPI_INSTALL_DIR)/lib -lthapi-ctl -llttng-ctl
 
 
 .PHONY: all

integration_tests/libthapi-ctl/hello-cl-startstop.c

@@ -6,7 +6,7 @@
 
 #include "thapi-ctl.h"
 
-#define N 11
+#define NLOOP 11
 
 int main(int argc, char* argv[]) {
 
@@ -86,85 +86,29 @@ int main(int argc, char* argv[]) {
     cl_command_queue queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &err);
     check_error(err,"clCreateCommandQueue");
 
-    // |/  _  ._ ._   _  |
-    // |\ (/_ |  | | (/_ |
-    //
-    printf(">>> Kernel configuration...\n");
-
-    // String kernel.
-    // Use MPI terminology (we are HPC!)
-    const char *kernelstring =
-        "__kernel void hello_world(const int i) {"
-        "   const int world_rank = get_global_id(0);" 
-        "   if (world_rank == i) printf(\"Hello world from rank %d \\n\", world_rank);" 
-        "}";
-
-    // Create the program
-    cl_program program = clCreateProgramWithSource(context, 1, &kernelstring, NULL, &err);
-    check_error(err,"clCreateProgramWithSource");
-
-    //Build / Compile the program executable
-    err = clBuildProgram(program, device_count, devices, "-cl-unsafe-math-optimizations", NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        printf("Error: Failed to build program executable!\n");
-
-        size_t logSize;
-        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
-
-        // there's no information in the reference whether the string is 0 terminated or not.
-        char* messages = (char*)malloc((1+logSize)*sizeof(char));
-        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, logSize, messages, NULL);
-        messages[logSize] = '\0';
-
-        printf("%s", messages);
-        free(messages);
-        return EXIT_FAILURE;
+    // create a device buffer
+    size_t a_size = 100;
+    size_t a_bytes = a_size * sizeof(float);
+    float *h_a = (float*)malloc(a_bytes);
+    for (unsigned int i = 0; i < a_size; i++) {
+        h_a[i] = (float)i;
     }
-
-   /* - - - -
-    Create
-    - - - - */
-    cl_kernel kernel = clCreateKernel(program, "hello_world", &err);
-    check_error(err,"clCreateKernel");
-
-    /* - - - -
-    ND range
-    - - - - */
-    printf(">>> NDrange configuration...\n");
-
-    #define WORK_DIM 1
-
-    // Describe the number of global work-items in work_dim dimensions that will execute the kernel function
-    size_t global0 = 100;
-    if (argc > 3)
-      global0 = (size_t) atoi(argv[3]);
-    const size_t global[WORK_DIM] = {  global0 };
-
-    // Describe the number of work-items that make up a work-group (also referred to as the size of the work-group).
-    // local_work_size can also be a NULL value in which case the OpenCL implementation will determine how to be break the global work-items into appropriate work-group instances.
-    size_t local0 = global0 / 10;
-    if (argc > 4)
-      local0 = (size_t) atoi(argv[4]);
-    const size_t local[WORK_DIM] = { local0 };
-
-    printf("Global work size: %zu \n", global[0]);
-    printf("Local work size: %zu \n", local[0]);
+    cl_mem d_a = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_HOST_READ_ONLY, a_bytes, NULL, &err);
+    check_error(err,"cclCreateBuffer");
 
     /* - - - -
     Execute
     - - - - */
-    for (int i = 0; i < N; i++) {
+    for (unsigned int i = 0; i < NLOOP; i++) {
       printf(">>> Kernel Execution [%d]...\n", i);
 
-      clSetKernelArg(kernel, 0, sizeof(int), &i);
-
       // trace only odd executions
       if (i % 2 == 1)
          thapi_start_tracing();
 
-      err  = clEnqueueNDRangeKernel(queue, kernel, WORK_DIM, NULL, global, local, 0, NULL, NULL);
-      check_error(err,"clEnqueueNDRangeKernel");
+    // (cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
+      err  = clEnqueueWriteBuffer(queue, d_a, CL_FALSE, 0, a_bytes, h_a, 0, NULL, NULL);
+      check_error(err,"clEnqueueWriteBuffer");
 
       /* - - -
       Sync & check
@@ -184,9 +128,8 @@ int main(int argc, char* argv[]) {
     free(devices);
     free(platforms);
     clReleaseCommandQueue(queue);
+    clReleaseMemObject(d_a);
     clReleaseContext(context);
-    clReleaseProgram(program);
-    clReleaseKernel(kernel);
 
     // Exit
     return 0;