Javier Luis Lopez Ответов: 1

Async_work_group_copy float4 to float


Можно ли скопировать float-массивы в float4?
Я не знаю, можно ли выровнять элементы массива float4 с массивом float для их копирования.

Я попробовал это сделать, но не смог скомпилировать:

Что я уже пробовал:

#define WD2 WIDTH/4
__global float A[WIDTH*HEIGHT];
...
__local float4 B[WD2];
barrier(CLK_LOCAL_MEM_FENCE);
async_work_group_copy(B,A+j*WIDTH,WD2,0);

Jochen Arndt

Если вы посмотрите на объявление async_work_group_copy (), то увидите, что аргументы src и dst должны иметь один и тот же тип, но вы передаете float4* и float*.

Вы можете попробовать передать & (B[0].x), что должно сделать компилятор счастливым. Но не спрашивайте меня, разрешено ли это или рекомендуется.

Javier Luis Lopez

Соответственно Хронос float4 должен быть выровнен с float2 и float1 (float3 с float4), как можно видеть здесь: https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/dataTypes.html

К сожалению, что-то не так, потому что это не сработало, я попробовал это с учарами:
__global uchar* image0;
.....
если (pix= = 0)
{
__local uchar16 В. В.;
((__местного Учар *) &ампер async_work_group_copy;ВВ&амп;imagen0[0],16,0);
printf ("===GPU vv: %4v16i \n", vv);
}
если (pix< 16)
printf ("imagen0[%2i]=%3i", pix, imagen0[pix]);


Но результат был неправильным, только первое значение было правильным. Результат можно увидеть здесь:
https://photos.app.goo.gl/QrgjyaNRvxnnKN672

Jochen Arndt

Просто убедитесь, что выравнивание правильное. С вашим исходным кодом это должно быть так, потому что float4 всегда правильно выровнен, а глобальный float должен быть по крайней мере 32-битным выровнен компилятором (64-битный с 64-битными приложениями).

Однако я мало что знал об OpenCL и поэтому не знал, разрешен ли он, рекомендован или работает.

Javier Luis Lopez

Спасибо, Яхен, но, как вы говорите, это может быть неправильно в зависимости от компилятора (и драйвера gpu).

Можно ли использовать vload или vload4 для загрузки массивов данных?

Karthik_Mahalingam

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

1 Ответов

Рейтинг:
0

Javier Luis Lopez

Наконец-то vloadn работает, но, к сожалению, мне приходится копировать только один вектор, а не массив из них:

__global float*  imagen0
...
long pix = get_global_id(0);
if (pix==0)
{
	float16 vv=vload16(0,imagen0);
	printf("===GPU vv: %6v16f \n",vv);
}
if (pix<16)
	printf("imagen0[%2i]=%6f",pix,imagen0[pix]);