SPIR-V -> NV50 IR: Texture R/W

ArchivedPublic

Actions

__kernel void rw_image(__read_only image2d_t in, __write_only image2d_t out)
{
  int2 coords = (int2) (get_global_id(0), get_global_id(1));
  if (coords.x > 4 || coords.y > 4)
    return;

  sampler_t const smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
  float4 val = read_imagef(in, smp, coords);
  write_imagef(out, coords, val * 2.0f);
}

#include <stdio.h>
#include <stdlib.h>

#include <CL/cl.h>

#include "cl_util.h"


int main(int argc, char ** argv)
{
   cl_int error;

   cl_uint total_platforms;
   cl_platform_id platform_id;

   cl_uint total_gpu_devices;
   cl_device_id device_id;

   cl_context context;

   cl_command_queue command_queue;

   cl_program program;

   cl_kernel kernel;

   FILE *input;
   void *il;
   size_t length;

   cl_image_format format;

   float in_data[64];
   float out_data[64];
   cl_mem in_image;
   cl_mem out_image;
   size_t global_work_size[2] = { 4, 4 };
   size_t local_work_size[2] = { 1, 1 };

   size_t origin[3] = { 0, 0, 0 };
   size_t region[3] = { 4, 4, 1 };

   cl_int status;

   error = clGetPlatformIDs(
                           1, /* Max number of platform IDs to return */
                           &platform_id, /* Pointer to platform_id */
                           &total_platforms); /* Total number of platforms
                                               * found on the system */

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clGetPlatformIDs() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "There are %u platforms.\n", total_platforms);



   error = clGetDeviceIDs(platform_id,
                          CL_DEVICE_TYPE_GPU,
                          1,
                          &device_id,
                          &total_gpu_devices);

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clGetDeviceIDs() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "There are %u GPU devices.\n", total_gpu_devices);

   context = clCreateContext(NULL, /* Properties */
                           1, /* Number of devices */
                           &device_id, /* Device pointer */
                           NULL, /* Callback for reporting errors */
                           NULL, /* User data to pass to error callback */
                           &error); /* Error code */

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clCreateContext() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clCreateContext() succeeded.\n");

   command_queue = clCreateCommandQueue(context,
                                        device_id,
                                        0, /* Command queue properties */
                                        &error); /* Error code */

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clCreateCommandQueue() failed: %s\n",
                      clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clCreateCommandQueue() succeeded.\n");

   input = fopen("image2d_rw.spv", "rb");
   if (!input) {
     fprintf(stderr, "failed to open image2d_rw.spv\n");
     return EXIT_FAILURE;
   }

   fseek(input, 0l, SEEK_END);
   length = ftell(input);
   rewind(input);

   il = calloc(length, sizeof(char));
   fread(il, length, 1, input);

//   program = clCreateProgramWithBinary(context, 1, &device_id, &length, (const unsigned char**)&il, &status, &error);
   program = clCreateProgramWithIL(context, il, length, &error);

   if (error != CL_SUCCESS) {
//     fprintf(stderr, "clCreateProgramWithBinary() failed: %s\n",
      fprintf(stderr, "clCreateProgramWithIL() failed: %s\n",
                      clUtilErrorString(error));
      return EXIT_FAILURE;
   }

//   fprintf(stderr, "clCreateProgramWithBinary() succeeded\n");
   fprintf(stderr, "clCreateProgramWithIL() suceeded.\n");

   error = clBuildProgram(program,
                          1, /* Number of devices */
                          &device_id,
                          NULL, /* options */
                          NULL, /* callback function when compile is complete */
                          NULL); /* user data for callback */


   if (error != CL_SUCCESS) {
      char build_str[1000000];
      error = clGetProgramBuildInfo(program,
                                    device_id,
                                    CL_PROGRAM_BUILD_LOG,
                                    1000000, /* Size of output string */
                                    build_str, /* pointer to write the log to */
                                    NULL); /* Number of bytes written to the log */
      if (error != CL_SUCCESS) {
         fprintf(stderr, "clGetProgramBuildInfo() failed: %s\n",
                          clUtilErrorString(error));
      } else {
         fprintf(stderr, "Build Log: \n%s\n\n", build_str);
      }
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clBuildProgram() suceeded.\n");

   kernel = clCreateKernel(program, "rw_image", &error);

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clCreateKernel() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clCreateKernel() suceeded.\n");

   format.image_channel_order = CL_RGBA;
   format.image_channel_data_type = CL_FLOAT;

   for (int i = 0; i < 64; ++i)
     out_data[i] = 0.0f;

   out_image = clCreateImage2D(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, &format, 4, 4, 0, out_data, &error);

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clCreateImage2D() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clCreateImage2D() succeeded.\n");

   error = clSetKernelArg(kernel,
                          0, /* Arg index */
                          sizeof(cl_mem),
                          &out_image);

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clSetKernelArg failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clSetKernelArg() succeeded.\n");

   for (int i = 0; i < 64; ++i)
     in_data[i] = (float) i;

   in_image = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &format, 4, 4, 0, in_data, &error);

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clCreateImage2D() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clCreateImage2D() succeeded.\n");

   error = clSetKernelArg(kernel,
                          1, /* Arg index */
                          sizeof(cl_mem),
                          &in_image);

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clSetKernelArg failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clSetKernelArg() succeeded.\n");


   error = clEnqueueNDRangeKernel(command_queue,
                                  kernel,
                                  2, /* Number of dimensions */
                                  NULL, /* Global work offset */
                                  global_work_size,
                                  local_work_size, /* local work size */
                                  0, /* Events in wait list */
                                  NULL, /* Wait list */
                                  NULL); /* Event object for this event */

   if (error != CL_SUCCESS) {
      fprintf(stderr, "clEnqueueNDRangeKernel() failed: %s\n",
                      clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clEnqueueNDRangeKernel() suceeded.\n");

   error = clFinish(command_queue);
   if (error != CL_SUCCESS) {
      fprintf(stderr, "clFinish() failed: %s\n", clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clFinish() succeeded.\n");

   error = clEnqueueReadImage(command_queue,
                                out_image,
                                CL_TRUE, /* TRUE means it is a blocking read. */
                                origin,
                                region,
                                0,
                                0,
                                out_data, /* Pointer to store the data */
                                0, /* Events in wait list */
                                NULL, /* Wait list */
                                NULL); /* Event object */


   if (error != CL_SUCCESS) {
      fprintf(stderr, "clEnqueueReadImage() failed: %s\n",
                      clUtilErrorString(error));
      return EXIT_FAILURE;
   }

   fprintf(stderr, "clEnqueueReadImage() suceeded.\n");

   for (int i = 0; i < 64; ++i)
     if (out_data[i] != 2.0f * in_data[i])
       fprintf(stderr, "out_data[%d] = %f, expected %f\n", i, out_data[i], (2.0f * in_data[i]));
   return EXIT_SUCCESS;
}