P4 to Git Change 1083437 by bsumner@bsumner-lin-opencl on 2014/10/01 20:06:47
ECR #304775 - move runtime kernels mostly into library Affected files ... ... //depot/stg/opencl/drivers/opencl/library/common.hsa/src/blitKernels.cl#1 add ... //depot/stg/opencl/drivers/opencl/library/common/src/blitKernels.cl#1 add ... //depot/stg/opencl/drivers/opencl/library/hsa/hsail/src/devenq/devenq.h#4 edit ... //depot/stg/opencl/drivers/opencl/library/hsa/hsail/src/devenq/schedule.cl#1 add ... //depot/stg/opencl/drivers/opencl/library/x86/gen/build/Makefile.gen#17 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/blitcl.cpp#5 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuschedcl.cpp#30 edit
This commit is contained in:
foreman
@@ -7,430 +7,160 @@ namespace device {
|
||||
#define BLIT_KERNELS(...) #__VA_ARGS__
|
||||
|
||||
const char* BlitSourceCode = BLIT_KERNELS(
|
||||
\n
|
||||
|
||||
extern void __amd_copyBufferToImage(
|
||||
__global uint*, __write_only image2d_array_t, ulong4,
|
||||
int4, int4, uint4, ulong4);
|
||||
|
||||
extern void __amd_copyImageToBuffer(
|
||||
__read_only image2d_array_t, __global uint*, __global ushort*,
|
||||
__global uchar*, int4, ulong4, int4, uint4, ulong4);
|
||||
|
||||
extern void __amd_copyImage(
|
||||
__read_only image2d_array_t, __write_only image2d_array_t,
|
||||
int4, int4, int4);
|
||||
|
||||
extern void __amd_copyImage1DA(
|
||||
__read_only image2d_array_t, __write_only image2d_array_t,
|
||||
int4, int4, int4);
|
||||
|
||||
extern void __amd_copyBufferRect(
|
||||
__global uchar*, __global uchar*,
|
||||
ulong4, ulong4, ulong4);
|
||||
|
||||
extern void __amd_copyBufferRectAligned(
|
||||
__global uint*, __global uint*,
|
||||
ulong4, ulong4, ulong4);
|
||||
|
||||
extern void __amd_copyBuffer(
|
||||
__global uchar*, __global uchar*,
|
||||
ulong, ulong, ulong);
|
||||
|
||||
extern void __amd_copyBufferAligned(
|
||||
__global uint*, __global uint*,
|
||||
ulong, ulong, ulong, uint);
|
||||
|
||||
extern void __amd_fillBuffer(
|
||||
__global uchar*, __global uint*, __constant uchar*,
|
||||
uint, ulong, ulong);
|
||||
|
||||
extern void __amd_fillImage(
|
||||
__write_only image2d_array_t,
|
||||
float4, int4, uint4, int4, int4, uint);
|
||||
|
||||
|
||||
__kernel void copyBufferToImage(
|
||||
__global uint* src,
|
||||
__write_only image2d_array_t dst,
|
||||
ulong4 srcOrigin,
|
||||
int4 dstOrigin,
|
||||
int4 size,
|
||||
uint4 format,
|
||||
ulong4 pitch)
|
||||
__global uint* src,
|
||||
__write_only image2d_array_t dst,
|
||||
ulong4 srcOrigin,
|
||||
int4 dstOrigin,
|
||||
int4 size,
|
||||
uint4 format,
|
||||
ulong4 pitch)
|
||||
{
|
||||
ulong idxSrc;
|
||||
int4 coordsDst;
|
||||
uint4 pixel;
|
||||
__global uint* srcUInt = src;
|
||||
__global ushort* srcUShort = (__global ushort*)src;
|
||||
__global uchar* srcUChar = (__global uchar*)src;
|
||||
ushort tmpUShort;
|
||||
uint tmpUInt;
|
||||
|
||||
coordsDst.x = get_global_id(0);
|
||||
coordsDst.y = get_global_id(1);
|
||||
coordsDst.z = get_global_id(2);
|
||||
coordsDst.w = 0;
|
||||
|
||||
if ((coordsDst.x >= size.x) ||
|
||||
(coordsDst.y >= size.y) ||
|
||||
(coordsDst.z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
idxSrc = (coordsDst.z * pitch.y +
|
||||
coordsDst.y * pitch.x + coordsDst.x) *
|
||||
format.z + srcOrigin.x;
|
||||
|
||||
coordsDst.x += dstOrigin.x;
|
||||
coordsDst.y += dstOrigin.y;
|
||||
coordsDst.z += dstOrigin.z;
|
||||
|
||||
// Check components
|
||||
switch (format.x) {
|
||||
case 1:
|
||||
// Check size
|
||||
if (format.y == 1) {
|
||||
pixel.x = (uint)srcUChar[idxSrc];
|
||||
}
|
||||
else if (format.y == 2) {
|
||||
pixel.x = (uint)srcUShort[idxSrc];
|
||||
}
|
||||
else {
|
||||
pixel.x = srcUInt[idxSrc];
|
||||
}
|
||||
break;
|
||||
case 2:
|
||||
// Check size
|
||||
if (format.y == 1) {
|
||||
tmpUShort = srcUShort[idxSrc];
|
||||
pixel.x = (uint)(tmpUShort & 0xff);
|
||||
pixel.y = (uint)(tmpUShort >> 8);
|
||||
}
|
||||
else if (format.y == 2) {
|
||||
tmpUInt = srcUInt[idxSrc];
|
||||
pixel.x = (tmpUInt & 0xffff);
|
||||
pixel.y = (tmpUInt >> 16);
|
||||
}
|
||||
else {
|
||||
pixel.x = srcUInt[idxSrc++];
|
||||
pixel.y = srcUInt[idxSrc];
|
||||
}
|
||||
break;
|
||||
case 4:
|
||||
// Check size
|
||||
if (format.y == 1) {
|
||||
tmpUInt = srcUInt[idxSrc];
|
||||
pixel.x = tmpUInt & 0xff;
|
||||
pixel.y = (tmpUInt >> 8) & 0xff;
|
||||
pixel.z = (tmpUInt >> 16) & 0xff;
|
||||
pixel.w = (tmpUInt >> 24) & 0xff;
|
||||
}
|
||||
else if (format.y == 2) {
|
||||
tmpUInt = srcUInt[idxSrc++];
|
||||
pixel.x = tmpUInt & 0xffff;
|
||||
pixel.y = (tmpUInt >> 16);
|
||||
tmpUInt = srcUInt[idxSrc];
|
||||
pixel.z = tmpUInt & 0xffff;
|
||||
pixel.w = (tmpUInt >> 16);
|
||||
}
|
||||
else {
|
||||
pixel.x = srcUInt[idxSrc++];
|
||||
pixel.y = srcUInt[idxSrc++];
|
||||
pixel.z = srcUInt[idxSrc++];
|
||||
pixel.w = srcUInt[idxSrc];
|
||||
}
|
||||
break;
|
||||
}
|
||||
// Write the final pixel
|
||||
write_imageui(dst, coordsDst, pixel);
|
||||
__amd_copyBufferToImage(src, dst, srcOrigin, dstOrigin, size, format, pitch);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyImageToBuffer(
|
||||
__read_only image2d_array_t src,
|
||||
__global uint* dstUInt,
|
||||
__global ushort* dstUShort,
|
||||
__global uchar* dstUChar,
|
||||
int4 srcOrigin,
|
||||
ulong4 dstOrigin,
|
||||
int4 size,
|
||||
uint4 format,
|
||||
ulong4 pitch)
|
||||
__read_only image2d_array_t src,
|
||||
__global uint* dstUInt,
|
||||
__global ushort* dstUShort,
|
||||
__global uchar* dstUChar,
|
||||
int4 srcOrigin,
|
||||
ulong4 dstOrigin,
|
||||
int4 size,
|
||||
uint4 format,
|
||||
ulong4 pitch)
|
||||
{
|
||||
ulong idxDst;
|
||||
int4 coordsSrc;
|
||||
uint4 texel;
|
||||
|
||||
coordsSrc.x = get_global_id(0);
|
||||
coordsSrc.y = get_global_id(1);
|
||||
coordsSrc.z = get_global_id(2);
|
||||
coordsSrc.w = 0;
|
||||
|
||||
if ((coordsSrc.x >= size.x) ||
|
||||
(coordsSrc.y >= size.y) ||
|
||||
(coordsSrc.z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
idxDst = (coordsSrc.z * pitch.y + coordsSrc.y * pitch.x +
|
||||
coordsSrc.x) * format.z + dstOrigin.x;
|
||||
|
||||
coordsSrc.x += srcOrigin.x;
|
||||
coordsSrc.y += srcOrigin.y;
|
||||
coordsSrc.z += srcOrigin.z;
|
||||
|
||||
texel = read_imageui(src, coordsSrc);
|
||||
|
||||
// Check components
|
||||
switch (format.x) {
|
||||
case 1:
|
||||
// Check size
|
||||
switch (format.y) {
|
||||
case 1:
|
||||
dstUChar[idxDst] = (uchar)texel.x;
|
||||
break;
|
||||
case 2:
|
||||
dstUShort[idxDst] = (ushort)texel.x;
|
||||
break;
|
||||
case 4:
|
||||
dstUInt[idxDst] = texel.x;
|
||||
break;
|
||||
}
|
||||
break;
|
||||
case 2:
|
||||
// Check size
|
||||
switch (format.y) {
|
||||
case 1:
|
||||
dstUShort[idxDst] = (ushort)texel.x |
|
||||
((ushort)texel.y << 8);
|
||||
break;
|
||||
case 2:
|
||||
dstUInt[idxDst] = texel.x | (texel.y << 16);
|
||||
break;
|
||||
case 4:
|
||||
dstUInt[idxDst++] = texel.x;
|
||||
dstUInt[idxDst] = texel.y;
|
||||
break;
|
||||
}
|
||||
break;
|
||||
case 4:
|
||||
// Check size
|
||||
switch (format.y) {
|
||||
case 1:
|
||||
dstUInt[idxDst] = (uint)texel.x |
|
||||
(texel.y << 8) |
|
||||
(texel.z << 16) |
|
||||
(texel.w << 24);
|
||||
break;
|
||||
case 2:
|
||||
dstUInt[idxDst++] = texel.x | (texel.y << 16);
|
||||
dstUInt[idxDst] = texel.z | (texel.w << 16);
|
||||
break;
|
||||
case 4:
|
||||
dstUInt[idxDst++] = texel.x;
|
||||
dstUInt[idxDst++] = texel.y;
|
||||
dstUInt[idxDst++] = texel.z;
|
||||
dstUInt[idxDst] = texel.w;
|
||||
break;
|
||||
}
|
||||
break;
|
||||
}
|
||||
__amd_copyImageToBuffer(src, dstUInt, dstUShort, dstUChar,
|
||||
srcOrigin, dstOrigin, size, format, pitch);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyImage(
|
||||
__read_only image2d_array_t src,
|
||||
__write_only image2d_array_t dst,
|
||||
int4 srcOrigin,
|
||||
int4 dstOrigin,
|
||||
int4 size)
|
||||
int4 srcOrigin,
|
||||
int4 dstOrigin,
|
||||
int4 size)
|
||||
{
|
||||
int4 coordsDst;
|
||||
int4 coordsSrc;
|
||||
|
||||
coordsDst.x = get_global_id(0);
|
||||
coordsDst.y = get_global_id(1);
|
||||
coordsDst.z = get_global_id(2);
|
||||
coordsDst.w = 0;
|
||||
|
||||
if ((coordsDst.x >= size.x) ||
|
||||
(coordsDst.y >= size.y) ||
|
||||
(coordsDst.z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
coordsSrc = srcOrigin + coordsDst;
|
||||
coordsDst += dstOrigin;
|
||||
|
||||
uint4 texel;
|
||||
texel = read_imageui(src, coordsSrc);
|
||||
write_imageui(dst, coordsDst, texel);
|
||||
__amd_copyImage(src, dst, srcOrigin, dstOrigin, size);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyImage1DA(
|
||||
__read_only image2d_array_t src,
|
||||
__write_only image2d_array_t dst,
|
||||
int4 srcOrigin,
|
||||
int4 dstOrigin,
|
||||
int4 size)
|
||||
__read_only image2d_array_t src,
|
||||
__write_only image2d_array_t dst,
|
||||
int4 srcOrigin,
|
||||
int4 dstOrigin,
|
||||
int4 size)
|
||||
{
|
||||
int4 coordsDst;
|
||||
int4 coordsSrc;
|
||||
|
||||
coordsDst.x = get_global_id(0);
|
||||
coordsDst.y = get_global_id(1);
|
||||
coordsDst.z = get_global_id(2);
|
||||
coordsDst.w = 0;
|
||||
|
||||
if ((coordsDst.x >= size.x) ||
|
||||
(coordsDst.y >= size.y) ||
|
||||
(coordsDst.z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
coordsSrc = srcOrigin + coordsDst;
|
||||
coordsDst += dstOrigin;
|
||||
if (srcOrigin.w != 0) {
|
||||
coordsSrc.z = coordsSrc.y;
|
||||
coordsSrc.y = 0;
|
||||
}
|
||||
if (dstOrigin.w != 0) {
|
||||
coordsDst.z = coordsDst.y;
|
||||
coordsDst.y = 0;
|
||||
}
|
||||
|
||||
uint4 texel;
|
||||
texel = read_imageui(src, coordsSrc);
|
||||
write_imageui(dst, coordsDst, texel);
|
||||
__amd_copyImage1DA(src, dst, srcOrigin, dstOrigin, size);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyBufferRect(
|
||||
__global uchar* src,
|
||||
__global uchar* dst,
|
||||
ulong4 srcRect,
|
||||
ulong4 dstRect,
|
||||
ulong4 size)
|
||||
__global uchar* src,
|
||||
__global uchar* dst,
|
||||
ulong4 srcRect,
|
||||
ulong4 dstRect,
|
||||
ulong4 size)
|
||||
{
|
||||
ulong x = get_global_id(0);
|
||||
ulong y = get_global_id(1);
|
||||
ulong z = get_global_id(2);
|
||||
|
||||
if ((x >= size.x) ||
|
||||
(y >= size.y) ||
|
||||
(z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
ulong offsSrc = srcRect.z + x + y * srcRect.x + z * srcRect.y;
|
||||
ulong offsDst = dstRect.z + x + y * dstRect.x + z * dstRect.y;
|
||||
|
||||
dst[offsDst] = src[offsSrc];
|
||||
__amd_copyBufferRect(src, dst, srcRect, dstRect, size);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyBufferRectAligned(
|
||||
__global uint* src,
|
||||
__global uint* dst,
|
||||
ulong4 srcRect,
|
||||
ulong4 dstRect,
|
||||
ulong4 size)
|
||||
__global uint* src,
|
||||
__global uint* dst,
|
||||
ulong4 srcRect,
|
||||
ulong4 dstRect,
|
||||
ulong4 size)
|
||||
{
|
||||
ulong x = get_global_id(0);
|
||||
ulong y = get_global_id(1);
|
||||
ulong z = get_global_id(2);
|
||||
|
||||
if ((x >= size.x) ||
|
||||
(y >= size.y) ||
|
||||
(z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
ulong offsSrc = srcRect.z + x + y * srcRect.x + z * srcRect.y;
|
||||
ulong offsDst = dstRect.z + x + y * dstRect.x + z * dstRect.y;
|
||||
|
||||
if (size.w == 16) {
|
||||
__global uint4* src4 = (__global uint4*)src;
|
||||
__global uint4* dst4 = (__global uint4*)dst;
|
||||
dst4[offsDst] = src4[offsSrc];
|
||||
}
|
||||
else {
|
||||
dst[offsDst] = src[offsSrc];
|
||||
}
|
||||
__amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyBuffer(
|
||||
__global uchar* src,
|
||||
__global uchar* dst,
|
||||
ulong srcOrigin,
|
||||
ulong dstOrigin,
|
||||
ulong size)
|
||||
__global uchar* src,
|
||||
__global uchar* dst,
|
||||
ulong srcOrigin,
|
||||
ulong dstOrigin,
|
||||
ulong size)
|
||||
{
|
||||
ulong id = get_global_id(0);
|
||||
|
||||
if (id >= size) {
|
||||
return;
|
||||
}
|
||||
|
||||
ulong offsSrc = id + srcOrigin;
|
||||
ulong offsDst = id + dstOrigin;
|
||||
|
||||
dst[offsDst] = src[offsSrc];
|
||||
__amd_copyBuffer(src, dst, srcOrigin, dstOrigin, size);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void copyBufferAligned(
|
||||
__global uint* src,
|
||||
__global uint* dst,
|
||||
ulong srcOrigin,
|
||||
ulong dstOrigin,
|
||||
ulong size,
|
||||
uint alignment)
|
||||
__global uint* src,
|
||||
__global uint* dst,
|
||||
ulong srcOrigin,
|
||||
ulong dstOrigin,
|
||||
ulong size,
|
||||
uint alignment)
|
||||
{
|
||||
ulong id = get_global_id(0);
|
||||
|
||||
if (id >= size) {
|
||||
return;
|
||||
}
|
||||
|
||||
ulong offsSrc = id + srcOrigin;
|
||||
ulong offsDst = id + dstOrigin;
|
||||
|
||||
if (alignment == 16) {
|
||||
__global uint4* src4 = (__global uint4*)src;
|
||||
__global uint4* dst4 = (__global uint4*)dst;
|
||||
dst4[offsDst] = src4[offsSrc];
|
||||
}
|
||||
else {
|
||||
dst[offsDst] = src[offsSrc];
|
||||
}
|
||||
__amd_copyBufferAligned(src, dst, srcOrigin, dstOrigin, size, alignment);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void fillBuffer(
|
||||
__global uchar* bufUChar,
|
||||
__global uint* bufUInt,
|
||||
__constant uchar* pattern,
|
||||
uint patternSize,
|
||||
ulong offset,
|
||||
ulong size)
|
||||
__global uchar* bufUChar,
|
||||
__global uint* bufUInt,
|
||||
__constant uchar* pattern,
|
||||
uint patternSize,
|
||||
ulong offset,
|
||||
ulong size)
|
||||
{
|
||||
ulong id = get_global_id(0);
|
||||
|
||||
if (id >= size) {
|
||||
return;
|
||||
}
|
||||
|
||||
if (bufUInt) {
|
||||
__global uint* element = &bufUInt[offset + id * patternSize];
|
||||
__constant uint* pt = (__constant uint*)pattern;
|
||||
|
||||
for (uint i = 0; i < patternSize; ++i) {
|
||||
element[i] = pt[i];
|
||||
}
|
||||
}
|
||||
else {
|
||||
__global uchar* element = &bufUChar[offset + id * patternSize];
|
||||
|
||||
for (uint i = 0; i < patternSize; ++i) {
|
||||
element[i] = pattern[i];
|
||||
}
|
||||
}
|
||||
__amd_fillBuffer(bufUChar, bufUInt, pattern, patternSize, offset, size);
|
||||
}
|
||||
\n
|
||||
|
||||
__kernel void fillImage(
|
||||
__write_only image2d_array_t image,
|
||||
float4 patternFLOAT4,
|
||||
int4 patternINT4,
|
||||
uint4 patternUINT4,
|
||||
int4 origin,
|
||||
int4 size,
|
||||
uint type)
|
||||
__write_only image2d_array_t image,
|
||||
float4 patternFLOAT4,
|
||||
int4 patternINT4,
|
||||
uint4 patternUINT4,
|
||||
int4 origin,
|
||||
int4 size,
|
||||
uint type)
|
||||
{
|
||||
int4 coords;
|
||||
|
||||
coords.x = get_global_id(0);
|
||||
coords.y = get_global_id(1);
|
||||
coords.z = get_global_id(2);
|
||||
coords.w = 0;
|
||||
|
||||
if ((coords.x >= size.x) ||
|
||||
(coords.y >= size.y) ||
|
||||
(coords.z >= size.z)) {
|
||||
return;
|
||||
}
|
||||
|
||||
coords += origin;
|
||||
|
||||
// Check components
|
||||
switch (type) {
|
||||
case 0:
|
||||
write_imagef(image, coords, patternFLOAT4);
|
||||
break;
|
||||
case 1:
|
||||
write_imagei(image, coords, patternINT4);
|
||||
break;
|
||||
case 2:
|
||||
write_imageui(image, coords, patternUINT4);
|
||||
break;
|
||||
}
|
||||
__amd_fillImage(image, patternFLOAT4, patternINT4, patternUINT4,
|
||||
origin, size, type);
|
||||
}
|
||||
\n
|
||||
\n
|
||||
|
||||
);
|
||||
|
||||
} // namespace device
|
||||
|
||||
Verwijs in nieuw issue
Block a user