echo"__device__ void load_global_register_size${amountToLoad}x${amountToLoad}(const float* in, float registers[$amountToLoad][$amountToLoad], int startX, int startY, int width, int height)
echo"__device__ void load_texCache_register_size${amountToLoad}x${amountToLoad}(const float* in, float registers[$amountToLoad][$amountToLoad], int startX, int startY, int width, int height)
__device__ void store_register_global_size${amountToLoad}x${amountToLoad}(float* out, float registers[$amountToLoad][$amountToLoad], int startX, int startY, int width, int height)
for((kernelSize=2; kernelSize<8; kernelSize++))do#convolution filter dim
if [ $inDim-ge$kernelSize ]
then
echo"
//@param height,width = dims of 'in' array
__device__ void convolutionDevice_size${inDim}x${inDim}_kernel${kernelSize}x${kernelSize}(float in[$inDim][$inDim], float *out, int startX, int startY)
sqrtConvsPerThread=$(($amountToLoad-$kernelSize+1))#square of this is total number of convs per thread [TODO: check correctness]
if [ $sqrtConvsPerThread-ge 1 ]
then
echo"__global__ void convolutionKernel_${memoryScheme}_size${amountToLoad}x${amountToLoad}_kernel${kernelSize}x${kernelSize}(const float* in, float* out, const int width, const int height)
{
int globalX = $sqrtConvsPerThread*(blockIdx.x*blockDim.x + threadIdx.x);
int globalY = $sqrtConvsPerThread*(blockIdx.y*blockDim.y + threadIdx.y);
load_${memoryScheme}_size${amountToLoad}x${amountToLoad}(in, registers, globalX-$((($kernelSize+1)/2-1)), globalY-$((($kernelSize+1)/2-1)), width, height); //includes offset so ghost zone is loaded
for(int y=0; y<$sqrtConvsPerThread; y++) //(int, y) = top left of region to convolve
echo"__global__ void convolutionKernel_${memoryScheme}_BenchmarkComputeTime_size${amountToLoad}x${amountToLoad}_kernel${kernelSize}x${kernelSize}(const float* in, float* out, const int width, const int height, float* fpuTime)
{
int globalX = $sqrtConvsPerThread*(blockIdx.x*blockDim.x + threadIdx.x);
int globalY = $sqrtConvsPerThread*(blockIdx.y*blockDim.y + threadIdx.y);
load_${memoryScheme}_size${amountToLoad}x${amountToLoad}(in, registers, globalX-$((($kernelSize+1)/2-1)), globalY-$((($kernelSize+1)/2-1)), width, height); //includes offset so ghost zone is loaded
__syncthreads();
double start = clock64();
//for(int y=0; y<$sqrtConvsPerThread; y++) //(int, y) = top left of region to convolve
__global__ void convolutionKernel_${memoryScheme}_size${amountToLoad}x${amountToLoad}_kernel${kernelSize}x${kernelSize}(const float* in, float* out, const int width, const int height)
{
int globalX = $sqrtConvsPerThread*(blockIdx.x*blockDim.x + threadIdx.x);
int globalY = $sqrtConvsPerThread*(blockIdx.y*blockDim.y + threadIdx.y);
#for(int outIdxX=globalX; outIdxX < globalX+$sqrtConvsPerThread; outIdxX++) //TODO: unroll this in script.
for((xx=0; xx<$sqrtConvsPerThread; xx++))do
echo"
outIdxX = globalX+$xx;
outIdxY = globalY+yy;
out[outIdxY*width + outIdxX]="
for((y=-$startOffset; y<$(($kernelSize-$startOffset)); y++)) do
for((x=-$startOffset; x<$(($kernelSize-$startOffset)); x++)) do
echo" in[clamp_addr(outIdxX+$x, outIdxY+$y, width, height)] * filter +"#TODO: use currX and currY
done
echo""
done
echo" 0;"
done
echo"
}
}
}"
done
done
#texCache only kernel, hard-coded bounds
memoryScheme="texCache_only"
for((kernelSize=2; kernelSize<8; kernelSize++))do
amountToLoad=$kernelSize
echo"
//@param height,width = dims of 'in' array
__global__ void convolutionKernel_${memoryScheme}_size${amountToLoad}x${amountToLoad}_kernel${kernelSize}x${kernelSize}(const float* in, float* out, const int width, const int height)
{
int globalX = (blockIdx.x*blockDim.x + threadIdx.x);
int globalY = (blockIdx.y*blockDim.y + threadIdx.y);
echo"//configuration directly pulled from simpleTexture in nvidia sdk
void setTexCacheParams()
{
tex.addressMode[0] = cudaAddressModeClamp;
tex.addressMode[1] = cudaAddressModeClamp;
tex.filterMode = cudaFilterModeLinear;
tex.normalized = false;
}"
echo"float convolutionWrapper(float* hImg, const int width, const int height, int amountToLoad, int kernelSize, string memoryScheme, bool outputImgFlag, string outFilename)
{
dim3 grid;
dim3 block;
block.x = 16;
block.y = 16;
int sqrtConvsPerThread = amountToLoad-kernelSize+1;
int nx = width / (block.x*sqrtConvsPerThread); //magic number is for persistant kernels (e.g. if persistant6x6, divide by 4)
int ny = height / (block.y*sqrtConvsPerThread);
grid.x = (width % block.x == 0) ? nx : nx+1;
grid.y = (height % block.y == 0) ? ny : ny+1;
float *dImg;
CHECK_CUDART(cudaMalloc((void**)&dImg, sizeof(float)*width*height)); //for input data (possibly use texture cache?)
printf("COMPUTE time on GPU, %dx%d filter: \n", kernelSize, kernelSize);
for(int i=0; i<(width*height); i++)
{
printf("%f\n", hFpuTime[i]);
}
}
CHECK_CUDART(cudaFree(dResult));
CHECK_CUDART(cudaFree(dFpuTime));
return responseTime;
}'
echo"float convolutionWrapper_texCache(float* hImg, const int width, const int height, int amountToLoad, int kernelSize, string memoryScheme, bool outputImgFlag, string outFilename)
{
dim3 grid;
dim3 block;
block.x = 16;
block.y = 16;
int sqrtConvsPerThread = amountToLoad-kernelSize+1;
int nx = width / (block.x*sqrtConvsPerThread); //magic number is for persistant kernels (e.g. if persistant6x6, divide by 4)