My intention is to use two OpenCL devices to run the same kernel, same kernel in a sense of the same .cl file not the same cl_kernel.
I have two decomposed domains (2D data), p0 and p1, which would require for them to exchange boundary information for every time step.
My strategy is to do the following.
- Get the boundary info of b0 and b1 for p0 and p1.
- Exchange boundary information: p0 uses b1 and p1 uses b0.
- Do some numerical operation for p0 and p1.
- And repeat.
For the step 2, I used clEnqueueMigrateMemObjects() function to migrate the buffer ownership one to another, and ran a kernel that would perform the boundary exchange.
From what I understand, in OpenCL, when you write data on buffers (using clEnqueueWriteBuffer()) you have to specify the command queue (let's say for device 0) that would be associated with. So, a written buffer should ONLY work for that specific command queue, right?
But, look at this example...
#include <stdio.h>
#include <stdlib.h>
#include "mult_funcs.h"
#define CL_TARGET_OPENCL_VERSION 200
#include <CL/cl.h>
#include <cl_tools.h>
#define clerr(arg, e) {printf(" %-40s : %s\n",arg,geterrstr(e));}
int main()
{
int i,j,t,ctr=0;
int nt = 1,
ny = 12,
nx = 12,
sy = ny,
sx = nx/2,
py = sy+2,
px = sx+2,
oy = 2,
ox = 2;
float u = 3.2,
v = 3.2,
c = 0.1;
float an0 [ny*nx], an1 [ny*nx],
sn0p0[sy*sx], sn0p1[sy*sx],
sn1p0[sy*sx], sn1p1[sy*sx],
pn0p0[py*px], pn0p1[py*px],
pn1p0[py*px], pn1p1[py*px];
float e0[sy], w0[sy], e1[sy], w1[sy];
// netcdf variables
int ec, nid, y_did, x_did, t_did, ovid,
dids[ond];
size_t start[ond] = {0, 0, 0},
count[ond] = {1,ny,nx};
// opencl related variables
size_t psiz = sizeof(float) * py * px,
bsiz = sizeof(float) * sy,
glbwksiz[2] = {sx,sy},
lclwksiz[2] = {16,16};
cl_int cle;
cl_uint nplf, ndev;
cl_context ctx;
cl_device_id devs[4];
cl_platform_id plf[1];
cl_command_queue cq0, cq1;
cl_program getb_pg_d0 , getb_pg_d1 ,
bexchg_pg_d0, bexchg_pg_d1;
cl_kernel getb_kn_n0p0 , getb_kn_n0p1 ,
getb_kn_n1p0 , getb_kn_n1p1 ,
bexchg_kn_n0p0, bexchg_kn_n0p1,
bexchg_kn_n1p0, bexchg_kn_n1p1;
cl_event getb_ev_n0p0 , getb_ev_n0p1 ,
getb_ev_n1p0 , getb_ev_n1p1 ,
mim_ev_d00 , mim_ev_d01 ,
mim_ev_d10 , mim_ev_d11 ,
bexchg_ev_n0p0, bexchg_ev_n0p1,
bexchg_ev_n1p0, bexchg_ev_n1p1;
cl_mem bn0p0, bn0p1, bn1p0, bn1p1,
be0 , bw0 , be1 , bw1 ;
// initializing OpenCL things : platforms and devices
cle = clGetPlatformIDs(0, NULL, &nplf);
cle = clGetPlatformIDs(nplf, plf, NULL);
cle = clGetDeviceIDs(plf[0], CL_DEVICE_TYPE_GPU, 0, NULL, &ndev);
cle = clGetDeviceIDs(plf[0], CL_DEVICE_TYPE_GPU, ndev, devs, NULL);
// initializing OpenCL things : context and command queues
ctx = clCreateContext(NULL, ndev, devs, NULL, NULL, &cle);
cq0 = clCreateCommandQueueWithProperties(ctx, devs[0], 0, &cle);
cq1 = clCreateCommandQueueWithProperties(ctx, devs[1], 0, &cle);
// buffers
bn0p0 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, psiz , NULL, &cle);
bn0p1 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, psiz , NULL, &cle);
be0 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bsiz, NULL, &cle);
bw0 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bsiz, NULL, &cle);
be1 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bsiz, NULL, &cle);
bw1 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bsiz, NULL, &cle);
// initialize array
arr_init0(an0,ny,nx);
arr_init0(an1,ny,nx);
// decompose array into two pieces
dcmp_data_2pc(an0,sn0p0,sn0p1,ny,nx,sy,sx);
dcmp_data_2pc(an1,sn1p0,sn1p1,ny,nx,sy,sx);
// padding zeros around the arrays
pad0(sn0p0,pn0p0,sy,sx); // ap part0
pad0(sn0p1,pn0p1,sy,sx); // ap part1
pad0(sn1p0,pn1p0,sy,sx); // ac part0
pad0(sn1p1,pn1p1,sy,sx); // ac part1
// write main arrays into buffers
// device 0
cle = clEnqueueWriteBuffer(cq0, bn0p0, CL_TRUE, 0, psiz, pn0p0, 0, NULL, NULL);
// device 1
cle = clEnqueueWriteBuffer(cq1, bn0p1, CL_TRUE, 0, psiz, pn0p1, 0, NULL, NULL);
// write boundary arrays into buffers
// device 0
cle = clEnqueueWriteBuffer(cq0, be0, CL_TRUE, 0, bsiz, e0, 0, NULL, NULL);
cle = clEnqueueWriteBuffer(cq0, bw0, CL_TRUE, 0, bsiz, w0, 0, NULL, NULL);
// device 1
cle = clEnqueueWriteBuffer(cq1, be1, CL_TRUE, 0, bsiz, e1, 0, NULL, NULL);
cle = clEnqueueWriteBuffer(cq1, bw1, CL_TRUE, 0, bsiz, w1, 0, NULL, NULL);
// creating program
getb_pg_d0 = createprogram(ctx, ndev, devs, "get_we_binfo_2d.cl" , 0);
getb_pg_d1 = createprogram(ctx, ndev, devs, "get_we_binfo_2d.cl" , 1);
bexchg_pg_d0 = createprogram(ctx, ndev, devs, "bexchg_we_2d.cl" , 0);
bexchg_pg_d1 = createprogram(ctx, ndev, devs, "bexchg_we_2d.cl" , 1);
// device 0
getb_kn_n0p0 = clCreateKernel(getb_pg_d0, "get_we_binfo_2d", &cle);
bexchg_kn_n0p0 = clCreateKernel(bexchg_pg_d0,"bexchg_we_2d" , &cle);
// device 1
getb_kn_n0p1 = clCreateKernel(getb_pg_d1, "get_we_binfo_2d", &cle);
bexchg_kn_n0p1 = clCreateKernel(bexchg_pg_d0,"bexchg_we_2d" , &cle);
// setting kernel arguments
// get_we_binfo for device 0
cle = clSetKernelArg(getb_kn_n0p0, 0, sizeof(cl_mem), &bn0p0);
cle = clSetKernelArg(getb_kn_n0p0, 1, sizeof(cl_int), &sx);
cle = clSetKernelArg(getb_kn_n0p0, 2, sizeof(cl_mem), &be0);
cle = clSetKernelArg(getb_kn_n0p0, 3, sizeof(cl_mem), &bw0);
// get_we_binfo for device 1
cle = clSetKernelArg(getb_kn_n0p1, 0, sizeof(cl_mem), &bn0p1);
cle = clSetKernelArg(getb_kn_n0p1, 1, sizeof(cl_int), &sx);
cle = clSetKernelArg(getb_kn_n0p1, 2, sizeof(cl_mem), &be1);
cle = clSetKernelArg(getb_kn_n0p1, 3, sizeof(cl_mem), &bw1);
// bexchg for device 0
cle = clSetKernelArg(bexchg_kn_n0p0, 0, sizeof(cl_mem), &bn0p0);
cle = clSetKernelArg(bexchg_kn_n0p0, 1, sizeof(cl_int), &sy);
cle = clSetKernelArg(bexchg_kn_n0p0, 2, sizeof(cl_int), &sx);
cle = clSetKernelArg(bexchg_kn_n0p0, 3, sizeof(cl_mem), &be1);
cle = clSetKernelArg(bexchg_kn_n0p0, 4, sizeof(cl_mem), &bw1);
// bexchg for device 1
cle = clSetKernelArg(bexchg_kn_n0p1, 0, sizeof(cl_mem), &bn0p1);
cle = clSetKernelArg(bexchg_kn_n0p1, 1, sizeof(cl_int), &sy);
cle = clSetKernelArg(bexchg_kn_n0p1, 2, sizeof(cl_int), &sx);
cle = clSetKernelArg(bexchg_kn_n0p1, 3, sizeof(cl_mem), &be0);
cle = clSetKernelArg(bexchg_kn_n0p1, 4, sizeof(cl_mem), &bw0);
while(ctr < nt)
{
cle = clEnqueueNDRangeKernel(cq0, getb_kn_n0p0, 1, NULL, &glbwksiz[1], NULL, 0, NULL, &getb_ev_n0p0);
cle = clEnqueueNDRangeKernel(cq1, getb_kn_n0p1, 1, NULL, &glbwksiz[1], NULL, 0, NULL, &getb_ev_n0p1);
cle = clEnqueueReadBuffer(cq0, be0, CL_TRUE, 0, bsiz, e0, 1, &getb_ev_n0p0, NULL);
cle = clEnqueueReadBuffer(cq0, bw0, CL_TRUE, 0, bsiz, w0, 1, &getb_ev_n0p0, NULL);
cle = clEnqueueReadBuffer(cq1, be1, CL_TRUE, 0, bsiz, e1, 1, &getb_ev_n0p1, NULL);
cle = clEnqueueReadBuffer(cq1, bw1, CL_TRUE, 0, bsiz, w1, 1, &getb_ev_n0p1, NULL);
cle = clEnqueueNDRangeKernel(cq0, bexchg_kn_n0p0, 2, NULL, glbwksiz, NULL, 0, NULL, &bexchg_ev_n0p0);
cle = clEnqueueNDRangeKernel(cq1, bexchg_kn_n0p1, 2, NULL, glbwksiz, NULL, 0, NULL, &bexchg_ev_n0p1);
ctr = ctr + 1;
}
return 0;
}
In the code above, I write some data on the buffer be0 and bw0 for command queue cq0 which is for device 0, and declared the buffer be1 and bw1 for device 1. And I declared the kernel getb_kn_n0p0for device 0 (also for device 1, for later use). Lastly, I set the kernel arguments bn0p0,sx, ,sy BUT I used be1 and bw1 instead of be0 and bw0. From how I understand it, this should not work. But, not only the code compiles without any errors but also it gives me the correct results, arrays with exchanged boundary information that are fetched from different devices. And, just like in the code, I did not have to use the clEnqueueMigrateMemObjects() function to change the ownership of the buffers.
Why is this allowed? or is this something that is allowed?
Thanks.