OpenCL (JOCL) - 2D исчисление по двум массивам в ядре
Я спрашиваю об этом здесь, потому что я думал, что понял, как работает OpenCL, но... я думаю, что есть несколько вещей, которые я не понимаю.
То, что я хочу сделать, это получить разницу между всеми значениями двух массивов, затем вычислить гипотезу и, наконец, получить максимальное значение гипотеты, поэтому, если у меня есть:
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)
Затем рассчитайте каждую гипотезу на основе гипотезы (dx(i), dy(i)). И как только все эти значения будут получены, получите максимальное значение гипотеты.
Итак, у меня есть следующее определенное ядро:
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 и LWI получаются следующим образом:
//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]};
Дело в том, что я не получаю ожидаемые значения, поэтому я решил сделать несколько тестов на основе меньшего ядра и изменить объект [VARIABLE TO TEST VALUES], возвращаемый в массиве результатов:
/**
* 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;"
+ " }"
Возвращаемые значения далеки от ожидаемых, но, если [VARIABLE TO TEST VALUES] равно (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
2 ответа
В вашем коде много проблем, и некоторые из них связаны с концепцией. Я думаю, что вы должны полностью прочитать стандарт или руководство OpenCL, прежде чем приступить к написанию кода. Потому что некоторые системные вызовы, которые вы используете, ведут себя иначе, чем вы ожидаете.
Рабочие группы и рабочие элементы НЕ похожи на CUDA. Если вы хотите, чтобы рабочие элементы размером 100x100 были разделены на рабочие группы 10x10, вы используете глобальный размер (100x100) и локальный размер (10x10). В отличие от CUDA, где глобальный рабочий элемент внутренне умножается на локальный размер.
1.1. В вашем тестовом коде, если вы используете 10x10 с 10x10. Тогда вы не заполняете все пространство, в незаполненной области все равно будет мусор
-X.xxxxxE-308,Вы не должны использовать lengthX и lengthY и помещать много if в ваш код. В OpenCL есть метод для вызова ядер со смещением и с определенным количеством элементов, поэтому вы можете управлять этим со стороны хоста. Кстати, это приводит к потере производительности и никогда не является хорошей практикой, поскольку код менее читабелен.
get_local_size(0)дает вам локальный размер оси 0 (10 в вашем случае). Что вы не понимаете в этом звонке? Почему вы делите это на 2 всегда?
Я надеюсь, что это может помочь вам в процессе отладки. ура
Спасибо за ваш ответ, прежде всего этот код ядра основан на коде коммутативного сокращения, объясненном здесь: http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/. Поэтому я использую этот код, но я добавил некоторые вещи, такие как 2D-операции.
Что касается пункта, который вы упомянули ранее:
1.1- На самом деле глобальный размер рабочей группы равен (100, 100, 0)... Эти 100 являются результатом умножения 10 x 10, где 10 - текущий размер массива, поэтому мой глобальный размер рабочей группы основан на этом правиле.. тогда размер локального рабочего элемента равен (10, 10, 0). Глобальный размер рабочей группы должен быть кратным локальному размеру рабочего элемента, я читал это во многих примерах и думаю, что это нормально.
1.2- В моем тестовом коде я использую те же массивы, фактически, если я изменю размер массива, размер GWG и размер LWI изменятся динамически.
2.1- Там не так много "если", есть только 3 "если", первый проверяет, когда я должен вычислить hypot() на основе объектов массива или заполнить этот объект нулем. Второе и третье "если" - это только часть алгоритма редукции, что, кажется, хорошо.
2.2- Что касается длины X и длины Y, да, вы правы, но у меня этого еще нет, как мне это использовать?
3.1- Да, я знаю это, но я понял, что я не использую идентификатор оси Y, так что, возможно, здесь есть другая проблема.
3.2- Алгоритм редукции перебирает каждую пару элементов, хранящихся в переменной нуля, и проверяет максимальное значение между ними, поэтому для каждого "для", которое он делает, он сокращает вычисляемые элементы до половины предыдущей величины.
Также я собираюсь опубликовать некоторые изменения в основном коде ядра и в тестовом коде ядра, потому что там, где есть некоторые ошибки.
Привет...!!!