Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add missing const qualifiers for inputs #220

Open
wants to merge 1 commit into
base: develop
Choose a base branch
from
Open

Add missing const qualifiers for inputs #220

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions src/library/generator.copy.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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";
}
}

Expand Down
54 changes: 27 additions & 27 deletions src/library/generator.stockham.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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 += ", ";
}
}

Expand Down Expand Up @@ -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
{
Expand Down Expand Up @@ -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";

}
}
Expand Down Expand Up @@ -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";
}
}

Expand Down
16 changes: 8 additions & 8 deletions src/tests/test_constants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include <stdexcept>

//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 \
Expand All @@ -36,22 +36,22 @@
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 \
float2 ret = *((__global float2*)in + offset) * scalar; \n \
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 \
Expand All @@ -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 \
Expand All @@ -69,22 +69,22 @@
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 \
return ret; \n \
}

//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 \
Expand Down