OpenCL (JOCL) - 2D-исчисление над двумя массивами в ядре

Question

Я спрашиваю об этом здесь, потому что я думал, что понял, как работает OpenCL, но ... Я думаю, что есть несколько вещей, которые я не получаю.

То, что я хочу сделать, это получить разницу между всеми значениями двух массивов, а затем вычислить hypot и, наконец, получить максимальное значение hypot, поэтому если у меня есть:

double[] arrA = new double[]{1,2,3} 
double[] arrB = new double[]{6,7,8} 

Calculate 
dx1 = 1 - 1; dx2 = 2 - 1; dx3 = 3 - 1, dx4= 1 - 2;... dxLast = 3 - 3 
dy1 = 6 - 6; dy2 = 7 - 6; dy3 = 8 - 6, dy4= 6 - 7;... dyLast = 8 - 8 

(Extreme dx and dy will get 0, but i don't care about ignoring those cases by now) 

Затем рассчитать каждый hypot основанный на hypot (дх (I), ау (я)) И когда все эти значения, где получены, получить максимальное значение hypot

Итак, у меня есть следующий определенный Kernel:

String programSource = 
    "#ifdef cl_khr_fp64 \n" 
+ " #pragma OPENCL EXTENSION cl_khr_fp64 : enable \n" 
+ "#elif defined(cl_amd_fp64) \n" 
+ " #pragma OPENCL EXTENSION cl_amd_fp64 : enable \n" 
+ "#else " 
+ " #error Double precision floating point not supported by OpenCL implementation.\n" 
+ "#endif \n" 
+ "__kernel void " 
+ "sampleKernel(__global const double *bufferX," 
+ "    __global const double *bufferY," 
+ "    __local double* scratch," 
+ "    __global double* result," 
+ "    __const int lengthX," 
+ "    __const int lengthY){" 
+ " const int index_a = get_global_id(0);"//Get the global indexes for 2D reference 
+ " const int index_b = get_global_id(1);" 
+ " const int local_index = get_local_id(0);"//Current thread id -> Should be the same as index_a * index_b + index_b; 
+ " if (local_index < (lengthX * lengthY)) {"// Load data into local memory 
+ "  if(index_a < lengthX && index_b < lengthY)" 
+ "  {" 
+ "   double dx = (bufferX[index_b] - bufferX[index_a]);" 
+ "   double dy = (bufferY[index_b] - bufferY[index_a]);" 
+ "   scratch[local_index] = hypot(dx, dy);" 
+ "  }" 
+ " } " 
+ " else {" 
+ "  scratch[local_index] = 0;"// Infinity is the identity element for the min operation 
+ " }" 
//Make a Barrier to make sure all values were set into the local array 
+ " barrier(CLK_LOCAL_MEM_FENCE);" 
//If someone can explain to me the offset thing I'll really apreciate that... 
//I just know there is alway a division by 2 
+ " for(int offset = get_local_size(0)/2; offset > 0; offset >>= 1) {" 
+ "  if (local_index < offset) {" 
+ "   float other = scratch[local_index + offset];" 
+ "   float mine = scratch[local_index];" 
+ "   scratch[local_index] = (mine > other) ? mine : other;" 
+ "  }" 
+ "  barrier(CLK_LOCAL_MEM_FENCE);" 
//A barrier to make sure that all values where checked 
+ " }" 
+ " if (local_index == 0) {" 
+ "  result[get_group_id(0)] = scratch[0];" 
+ " }" 
+ "}"; 

Для этого случая определенный размер GWG равен (100, 100, 0) и размер LWI (10, 10, 0).

Таким образом, для данного примера, оба массива имеет размера 10 и GWG и L получаются следующим образом:

//clGetKernelWorkGroupInfo(kernel, device, CL.CL_KERNEL_WORK_GROUP_SIZE, Sizeof.size_t, Pointer.to(buffer), null); 
long kernel_work_group_size = OpenClUtil.getKernelWorkGroupSize(kernel, device.getCl_device_id(), 3); 
//clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, Sizeof.size_t * numValues, Pointer.to(buffer), null); 
long[] maxSize = device.getMaximumSizes(); 

maxSize[0] = (kernel_work_group_size > maxSize[0] ? maxSize[0] : kernel_work_group_size); 
maxSize[1] = (kernel_work_group_size > maxSize[1] ? maxSize[1] : kernel_work_group_size); 
maxSize[2] = (kernel_work_group_size > maxSize[2] ? maxSize[2] : kernel_work_group_size); 
// maxSize[2] = 

long xMaxSize = (x > maxSize[0] ? maxSize[0] : x); 
long yMaxSize = (y > maxSize[1] ? maxSize[1] : y); 
long zMaxSize = (z > maxSize[2] ? maxSize[2] : z); 

long local_work_size[] = new long[] { xMaxSize, yMaxSize, zMaxSize }; 

int numWorkGroupsX = 0; 
int numWorkGroupsY = 0; 
int numWorkGroupsZ = 0; 

if(local_work_size[0] != 0) 
    numWorkGroupsX = (int) ((total + local_work_size[0] - 1)/local_work_size[0]); 

if(local_work_size[1] != 0) 
    numWorkGroupsY = (int) ((total + local_work_size[1] - 1)/local_work_size[1]); 

if(local_work_size[2] != 0) 
    numWorkGroupsZ = (int) ((total + local_work_size[2] - 1)/local_work_size[2]); 

long global_work_size[] = new long[] { numWorkGroupsX * local_work_size[0], 
    numWorkGroupsY * local_work_size[1], numWorkGroupsZ * local_work_size[2]}; 

Дело в том, что я не получает espected значения, так что я решил сделать некоторые тесты основанный на меньшем ядре и изменения [VARIABLE ИСПЫТАНИЯ зНАЧЕНИЯ] объект, возвращаемый в результате массив:

/** 
* The source code of the OpenCL program to execute 
*/ 
private static String programSourceA = 
    "#ifdef cl_khr_fp64 \n" 
+ " #pragma OPENCL EXTENSION cl_khr_fp64 : enable \n" 
+ "#elif defined(cl_amd_fp64) \n" 
+ " #pragma OPENCL EXTENSION cl_amd_fp64 : enable \n" 
+ "#else " 
+ " #error Double precision floating point not supported by OpenCL implementation.\n" 
+ "#endif \n" 
+ "__kernel void " 
+ "sampleKernel(__global const double *bufferX," 
+ "    __global const double *bufferY," 
+ "    __local double* scratch," 
+ "    __global double* result," 
+ "    __const int lengthX," 
+ "    __const int lengthY){" 
//Get the global indexes for 2D reference 
+ " const int index_a = get_global_id(0);" 
+ " const int index_b = get_global_id(1);" 
//Current thread id -> Should be the same as index_a * index_b + index_b; 
+ " const int local_index = get_local_id(0);" 
// Load data into local memory 
//Only print values if index_a < ArrayA length 
//Only print values if index_b < ArrayB length 
//Only print values if local_index < (lengthX * lengthY) 
//Only print values if this is the first work group. 
+ " if (local_index < (lengthX * lengthY)) {" 
+ "  if(index_a < lengthX && index_b < lengthY)" 
+ "  {" 
+ "   double dx = (bufferX[index_b] - bufferX[index_a]);" 
+ "   double dy = (bufferY[index_b] - bufferY[index_a]);" 
+ "   result[local_index] = hypot(dx, dy);" 
+ "  }" 
+ " } " 
+ " else {" 
// Infinity is the identity element for the min operation 
+ "  result[local_index] = 0;" 
+ " }" 

возвращаемые значения далеки того, чтобы быть espected, но, если [ПЕРЕМЕННОГО ИСПЫТАНИЯ зНАЧЕНИЯ] является (index_a * index_b) + index_a, почти каждое значение возвращаемого массива имеет правильное значение (Index_a * index_b) + значение index_a, я имею в виду:

result[0] -> 0 
result[1] -> 1 
result[2] -> 2 
.... 
result[97] -> 97 
result[98] -> 98 
result[99] -> 99 

, но некоторые значения: -3.350700319577517E-308 ....

Что я не делаю правильно ???

Я надеюсь, что это хорошо объяснено и не большой, чтобы вы сердитесь на меня ....

Огромное спасибо !!!!!

TomRacer

Answer 1

У вас есть много проблем в вашем коде, и некоторые из них понятие связаны между собой. Я думаю, вы должны полностью прочитать стандарт или OpenCL guide, прежде чем начинать код. Потому что некоторые из системных вызовов, которые вы используете, имеют другое поведение, которое вы ожидаете.

1. Рабочие группы и рабочие элементы НЕ относятся к CUDA. Если вы хотите 100x100 рабочих элементов, разделенных на 10x10 рабочих групп, которые вы используете как глобальные (100x100) и локальные (10x10). В отличие от CUDA, где глобальный рабочий элемент умножается на локальный размер внутри.

   1.1. В тестовом коде, если вы используете 10x10 с 10x10. Тогда вы не заполняете все пространство, в незаполненной области по-прежнему будет мусор, такой как -X.xxxxxE-308.

2. Вы не должны использовать lengthX и lengthY и вставлять много кода if. OpenCL имеет метод вызова ядер с смещениями и с определенным количеством элементов, поэтому вы можете контролировать это со стороны хоста. BTW делает это потеря производительности и никогда не является хорошей практикой, поскольку код менее читабельен.

3. get_local_size(0) дает локальный размер оси 0 (10 в вашем случае). Что вы не понимаете в этом призыве? Почему вы разделите его на 2 всегда?

Надеюсь, это поможет вам в процессе отладки. Cheers

Answer 2

Благодарим вас за ответ, в первую очередь этот код ядра основан на коде коммутативного сокращения, описанном здесь: http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/. Итак, я использую этот код, но добавил некоторые вещи, такие как 2D-операции.

Что касается до точки вы упомянули раньше:

1.1- Собственно глобальный размер рабочей группы составляет (100, 100, 0) ... Это 100 является результатом умножения 10 х 10, где 10 является текущим размер массива, поэтому мой глобальный размер рабочей группы основан на этом правиле ... тогда размер местного рабочего элемента (10, 10, 0). Размер глобальной рабочей группы должен быть кратным размеру местного рабочего элемента, я прочитал это во многих примерах, и я думаю, что все в порядке.

1.2- В моем тестовом коде я использую те же массивы, фактически, если я изменил размер GWG размера массива, а размер LWI изменится динамически.

2.1- Там не так много «если», есть только 3 «если», первый проверяет, когда я должен вычислить гипотезу() на основе объектов массива или заполнить этот объект нулем. Второе и третье «если» - это лишь часть алгоритма сокращения, который кажется прекрасным.

2.2- Относительно длины X и длины. Да, вы правы, но у меня пока нет этого, как я должен использовать это?

3.1- Да, я знаю это, но я понял, что я не использую идентификатор оси Y, поэтому, возможно, здесь есть еще одна проблема.

3.2- Алгоритм сокращения выполняет итерацию по каждой паре элементов, хранящихся в переменной нуля, и проверяет максимальное значение между ними, поэтому для каждого «для», что он делает это, уменьшается количество элементов, которые будут вычисляться до половины предыдущее количество.

Также я опубликую некоторые изменения в главном коде ядра и в тестовом коде ядра, потому что там, где есть некоторые ошибки.

Приветствия ... !!!