diff --git a/src/library/generator.copy.cpp b/src/library/generator.copy.cpp index b55f5992..a80bc775 100644 --- a/src/library/generator.copy.cpp +++ b/src/library/generator.copy.cpp @@ -258,12 +258,12 @@ namespace CopyGenerator // input if(inIlvd) { - str += "__global "; str += r2Type; str += " *lwbIn;\n\t"; + str += "__global const "; str += r2Type; str += " *lwbIn;\n\t"; } else { - str += "__global "; str += rType; str += " *lwbInRe;\n\t"; - str += "__global "; str += rType; str += " *lwbInIm;\n\t"; + str += "__global const "; str += rType; str += " *lwbInRe;\n\t"; + str += "__global const "; str += rType; str += " *lwbInIm;\n\t"; } } diff --git a/src/library/generator.stockham.cpp b/src/library/generator.stockham.cpp index 012b59cd..5b1e5f67 100644 --- a/src/library/generator.stockham.cpp +++ b/src/library/generator.stockham.cpp @@ -2123,20 +2123,20 @@ namespace StockhamGenerator { if(inInterleaved) { - passStr += "__global "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; - if(!rcSimple) { passStr += "__global "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe2; passStr += ", "; } + passStr += "__global const "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + if(!rcSimple) { passStr += "__global const "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe2; passStr += ", "; } } else if(inReal) { - passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; - if(!rcSimple) { passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe2; passStr += ", "; } + passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + if(!rcSimple) { passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe2; passStr += ", "; } } else { - passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; - if(!rcSimple) { passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe2; passStr += ", "; } - passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm; passStr += ", "; - if(!rcSimple) { passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm2; passStr += ", "; } + passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + if(!rcSimple) { passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe2; passStr += ", "; } + passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm; passStr += ", "; + if(!rcSimple) { passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm2; passStr += ", "; } } } else @@ -2177,24 +2177,24 @@ namespace StockhamGenerator { if(inInterleaved) { - passStr += "__global "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + passStr += "__global const "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; } else { - passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; - passStr += "__global "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm; passStr += ", "; + passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + passStr += "__global const "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm; passStr += ", "; } } else { if(inInterleaved) { - passStr += "__local "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + passStr += "__local const "; passStr += regB2Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; } else { - passStr += "__local "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; - passStr += "__local "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm; passStr += ", "; + passStr += "__local const "; passStr += regB1Type; passStr += " *"; passStr += bufferInRe; passStr += ", "; + passStr += "__local const "; passStr += regB1Type; passStr += " *"; passStr += bufferInIm; passStr += ", "; } } @@ -3625,11 +3625,11 @@ namespace StockhamGenerator { if(inInterleaved) { - str += "__global "; str += r2Type; str += " * restrict gbIn, "; + str += "__global const "; str += r2Type; str += " * restrict gbIn, "; } else if(inReal) { - str += "__global "; str += rType; str += " * restrict gbIn, "; + str += "__global const "; str += rType; str += " * restrict gbIn, "; } else { @@ -3736,22 +3736,22 @@ namespace StockhamGenerator { if(inInterleaved) { - if(!rcSimple) { str += "__global "; str += r2Type; str += " *lwbIn2;\n\t"; } - str += "__global "; str += r2Type; str += " *lwbIn;\n\t"; + if(!rcSimple) { str += "__global const "; str += r2Type; str += " *lwbIn2;\n\t"; } + str += "__global const "; str += r2Type; str += " *lwbIn;\n\t"; } else if(inReal) { - if(!rcSimple) { str += "__global "; str += rType; str += " *lwbIn2;\n\t"; } - str += "__global "; str += rType; str += " *lwbIn;\n\t"; + if(!rcSimple) { str += "__global const "; str += rType; str += " *lwbIn2;\n\t"; } + str += "__global const "; str += rType; str += " *lwbIn;\n\t"; } else { - if(!rcSimple) { str += "__global "; str += rType; str += " *lwbInRe2;\n\t"; } - if(!rcSimple) { str += "__global "; str += rType; str += " *lwbInIm2;\n\t"; } + if(!rcSimple) { str += "__global const "; str += rType; str += " *lwbInRe2;\n\t"; } + if(!rcSimple) { str += "__global const "; str += rType; str += " *lwbInIm2;\n\t"; } - str += "__global "; str += rType; str += " *lwbInRe;\n\t"; - str += "__global "; str += rType; str += " *lwbInIm;\n\t"; + str += "__global const "; str += rType; str += " *lwbInRe;\n\t"; + str += "__global const "; str += rType; str += " *lwbInIm;\n\t"; } } @@ -3815,12 +3815,12 @@ namespace StockhamGenerator { if(inInterleaved) { - str += "__global "; str += r2Type; str += " *lwbIn;\n\t"; + str += "__global const "; str += r2Type; str += " *lwbIn;\n\t"; } else { - str += "__global "; str += rType; str += " *lwbInRe;\n\t"; - str += "__global "; str += rType; str += " *lwbInIm;\n\t"; + str += "__global const "; str += rType; str += " *lwbInRe;\n\t"; + str += "__global const "; str += rType; str += " *lwbInIm;\n\t"; } } diff --git a/src/tests/test_constants.h b/src/tests/test_constants.h index b882ead4..3d689162 100644 --- a/src/tests/test_constants.h +++ b/src/tests/test_constants.h @@ -24,7 +24,7 @@ #include //Pre-callback function strings -#define PRE_MULVAL float2 mulval_pre(__global void* in, uint offset, __global void* userdata)\n \ +#define PRE_MULVAL float2 mulval_pre(__global const void* in, uint offset, __global void* userdata)\n \ { \n \ float scalar = *((__global float*)userdata + offset); \n \ float2 ret = *((__global float2*)in + offset) * scalar; \n \ @@ -36,7 +36,7 @@ float scalar1; \ float scalar2; \ } USER_DATA; \n \ - float2 mulval_pre(__global void* in, uint offset, __global void* userdata)\n \ + float2 mulval_pre(__global const void* in, uint offset, __global void* userdata)\n \ { \n \ __global USER_DATA *data = ((__global USER_DATA *)userdata + offset); \n \ float scalar = data->scalar1 * data->scalar2; \n \ @@ -44,14 +44,14 @@ return ret; \n \ } -#define PRE_MULVAL_DP double2 mulval_pre(__global void* in, uint offset, __global void* userdata)\n \ +#define PRE_MULVAL_DP double2 mulval_pre(__global const void* in, uint offset, __global void* userdata)\n \ { \n \ double scalar = *((__global double*)userdata + offset); \n \ double2 ret = *((__global double2*)in + offset) * scalar; \n \ return ret; \n \ } -#define PRE_MULVAL_PLANAR float2 mulval_pre(__global void* inRe, __global void* inIm, uint offset, __global void* userdata)\n \ +#define PRE_MULVAL_PLANAR float2 mulval_pre(__global const void* inRe, __global const void* inIm, uint offset, __global void* userdata)\n \ { \n \ float scalar = *((__global float*)userdata + offset); \n \ float2 ret; \n \ @@ -60,7 +60,7 @@ return ret; \n \ } -#define PRE_MULVAL_PLANAR_DP double2 mulval_pre(__global void* inRe, __global void* inIm, uint offset, __global void* userdata)\n \ +#define PRE_MULVAL_PLANAR_DP double2 mulval_pre(__global const void* inRe, __global const void* inIm, uint offset, __global void* userdata)\n \ { \n \ double scalar = *((__global double*)userdata + offset); \n \ double2 ret; \n \ @@ -69,14 +69,14 @@ return ret; \n \ } -#define PRE_MULVAL_REAL float mulval_pre(__global void* in, uint offset, __global void* userdata)\n \ +#define PRE_MULVAL_REAL float mulval_pre(__global const void* in, uint offset, __global void* userdata)\n \ { \n \ float scalar = *((__global float*)userdata + offset); \n \ float ret = *((__global float*)in + offset) * scalar; \n \ return ret; \n \ } -#define PRE_MULVAL_REAL_DP double mulval_pre(__global void* in, uint offset, __global void* userdata)\n \ +#define PRE_MULVAL_REAL_DP double mulval_pre(__global const void* in, uint offset, __global void* userdata)\n \ { \n \ double scalar = *((__global double*)userdata + offset); \n \ double ret = *((__global double*)in + offset) * scalar; \n \ @@ -84,7 +84,7 @@ } //Precallback test for LDS - works when 1 WI works on one input element -#define PRE_MULVAL_LDS float2 mulval_pre(__global void* in, uint offset, __global void* userdata, __local void* localmem)\n \ +#define PRE_MULVAL_LDS float2 mulval_pre(__global const void* in, uint offset, __global void* userdata, __local void* localmem)\n \ { \n \ uint lid = get_local_id(0); \n \ __local float* lds = (__local float*)localmem + lid; \n \