/* /\* */
/*  * testconv_extras.cu */
/*  * */
/*  *  Created on: Nov 10, 2009 */
/*  *      Author: Alex Krizhevsky (akrizhevsky@gmail.com) */
/*  *\/ */

/* #include <assert.h> */
/* #include "testconv_extras.cuh" */
/* #include "convCPU.h" */

/* void test_conv_bw_fit_dyn_2per(int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/*     int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */
/*     images.addScalar(1); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, 0, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     filters.print(3,3); */
/*     dim3 threads(8,8,8); */
/*     dim3 blocks(numCases, numFilters / 16); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     int shmem = 4*((filterSize + threads.x - 1)*(filterSize + threads.y - 1) + 2*threads.z * filterSize * filterSize); */
/*     conv_bw_fit_dyn_2per<9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), imgSize); */
/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, 0, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */


/* void test_conv_bw_nofit_dyn_1per(int imgSize, int filterSize, int threadsY, int threadsX, int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/* //    int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     assert(numOutputsX % threadsX == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, 0, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     filters.print(3,3); */
/*     int threadsZ = int(512.0 / (threadsX * threadsY)); */
/*     int blocksY = int(ceil(float(numFilters) / (threadsZ))); */
/*     bool checkFilterBounds = numOutputsX % filterSize != 0; */
/*     assert(threadsZ > 0); */
/*     dim3 threads(threadsX,threadsY,threadsZ); */
/*     dim3 blocks(numCases, blocksY); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     int shmem = 4*((2*threads.x - 1)*(2*threads.y - 1) + threads.z * threads.x * threads.y); */
/*     printf("Running %dx%d grid with %dx%dx%d blocks and %d bytes of shmem.\n", blocks.x, blocks.y, threads.x, threads.y, threads.z, shmem); */
/*     if(checkFilterBounds) { */
/*         conv_bw_nofit_dyn_1per<true,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                     imgSize, filterSize, numFilters); */
/*     } else { */
/*         conv_bw_nofit_dyn_1per<false,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                      imgSize, filterSize, numFilters); */
/*     } */

/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, 0, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */


/* void test_conv_bw_nofit_dyn_2per(int imgSize, int filterSize, int threadsY, int threadsX, int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/* //    int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     assert(numOutputsX % threadsX == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, 0, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     filters.print(3,3); */
/*     int threadsZ = int(512.0 / (threadsX * threadsY)); */
/*     int blocksY = int(ceil(float(numFilters) / (2*threadsZ))); */
/*     assert((numFilters % (threadsZ*2)) % 2 == 0); */
/*     bool checkFilterBounds = numOutputsX % filterSize != 0; */
/*     assert(threadsZ > 0); */
/*     dim3 threads(threadsX,threadsY,threadsZ); */
/*     dim3 blocks(numCases, blocksY); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     int shmem = 4*((2*threads.x - 1)*(2*threads.y - 1) + 2*threads.z * threads.x * threads.y); */
/*     printf("Running %dx%d grid with %dx%dx%d blocks and %d bytes of shmem.\n", blocks.x, blocks.y, threads.x, threads.y, threads.z, shmem); */
/*     assert(threadsX == 9 && threadsY == 9); */
/*     if(checkFilterBounds) { */
/*         conv_bw_nofit_dyn_2per<true,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                     imgSize, filterSize, numFilters); */
/*     } else { */
/*         conv_bw_nofit_dyn_2per<false,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                      imgSize, filterSize, numFilters); */
/*     } */

/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, 0, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */


/* void test_conv_bw_nofit_4x16_dynfilter_2per(int imgSize, int filterSize, int filterCacheY, int filterCacheX, int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/* //    int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     assert(filterSize % filterCacheX == 0); */
/*     assert(filterSize % filterCacheY == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, numOutputsX-10, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */


/*     bool checkOutputBounds = numOutputsX % 16 != 0; */
/*     int blocksY = numFilters / 16; */
/*     dim3 threads(16,4,8); */
/*     dim3 blocks(numCases, blocksY); */
/*     assert(filterCacheY == 8); */
/*     assert(filterCacheX == 16); */
/*     int shmem = (16 + filterCacheX - 1)*(4 + filterCacheY - 1) + 16*filterCacheX*filterCacheY; */
/*     assert(shmem < 4096); */
/*     printf("Filter cache size: %dx%d\n", filterCacheX, filterCacheY); */
/*     printf("Using %d bytes of shared memory\n", shmem*4); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */

/*     if(checkOutputBounds) { */
/*         conv_bw_nofit_4x16_dynfilter_2per<true,8,16, 1><<<blocks, threads>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                     imgSize, filterSize); */
/*     } else { */
/*         conv_bw_nofit_4x16_dynfilter_2per<false,8,16, 1><<<blocks, threads>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                      imgSize, filterSize); */
/*     } */

/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, numOutputsX-10, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */

	.file	"tmpxft_003777d8_00000000-6_testconv_extras.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
/* /\* */
/*  * testconv_extras.cu */
/*  * */
/*  *  Created on: Nov 10, 2009 */
/*  *      Author: Alex Krizhevsky (akrizhevsky@gmail.com) */
/*  *\/ */

/* #include <assert.h> */
/* #include "testconv_extras.cuh" */
/* #include "convCPU.h" */

/* void test_conv_bw_fit_dyn_2per(int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/*     int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */
/*     images.addScalar(1); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, 0, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     filters.print(3,3); */
/*     dim3 threads(8,8,8); */
/*     dim3 blocks(numCases, numFilters / 16); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     int shmem = 4*((filterSize + threads.x - 1)*(filterSize + threads.y - 1) + 2*threads.z * filterSize * filterSize); */
/*     conv_bw_fit_dyn_2per<9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), imgSize); */
/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, 0, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */


/* void test_conv_bw_nofit_dyn_1per(int imgSize, int filterSize, int threadsY, int threadsX, int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/* //    int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     assert(numOutputsX % threadsX == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, 0, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     filters.print(3,3); */
/*     int threadsZ = int(512.0 / (threadsX * threadsY)); */
/*     int blocksY = int(ceil(float(numFilters) / (threadsZ))); */
/*     bool checkFilterBounds = numOutputsX % filterSize != 0; */
/*     assert(threadsZ > 0); */
/*     dim3 threads(threadsX,threadsY,threadsZ); */
/*     dim3 blocks(numCases, blocksY); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     int shmem = 4*((2*threads.x - 1)*(2*threads.y - 1) + threads.z * threads.x * threads.y); */
/*     printf("Running %dx%d grid with %dx%dx%d blocks and %d bytes of shmem.\n", blocks.x, blocks.y, threads.x, threads.y, threads.z, shmem); */
/*     if(checkFilterBounds) { */
/*         conv_bw_nofit_dyn_1per<true,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                     imgSize, filterSize, numFilters); */
/*     } else { */
/*         conv_bw_nofit_dyn_1per<false,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                      imgSize, filterSize, numFilters); */
/*     } */

/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, 0, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */


/* void test_conv_bw_nofit_dyn_2per(int imgSize, int filterSize, int threadsY, int threadsX, int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/* //    int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     assert(numOutputsX % threadsX == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, 0, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     filters.print(3,3); */
/*     int threadsZ = int(512.0 / (threadsX * threadsY)); */
/*     int blocksY = int(ceil(float(numFilters) / (2*threadsZ))); */
/*     assert((numFilters % (threadsZ*2)) % 2 == 0); */
/*     bool checkFilterBounds = numOutputsX % filterSize != 0; */
/*     assert(threadsZ > 0); */
/*     dim3 threads(threadsX,threadsY,threadsZ); */
/*     dim3 blocks(numCases, blocksY); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     int shmem = 4*((2*threads.x - 1)*(2*threads.y - 1) + 2*threads.z * threads.x * threads.y); */
/*     printf("Running %dx%d grid with %dx%dx%d blocks and %d bytes of shmem.\n", blocks.x, blocks.y, threads.x, threads.y, threads.z, shmem); */
/*     assert(threadsX == 9 && threadsY == 9); */
/*     if(checkFilterBounds) { */
/*         conv_bw_nofit_dyn_2per<true,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                     imgSize, filterSize, numFilters); */
/*     } else { */
/*         conv_bw_nofit_dyn_2per<false,9,9, 1><<<blocks, threads, shmem>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                      imgSize, filterSize, numFilters); */
/*     } */

/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, 0, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */


/* void test_conv_bw_nofit_4x16_dynfilter_2per(int imgSize, int filterSize, int filterCacheY, int filterCacheX, int boardNum) { */
/*     cudaSetDevice(boardNum > -1 ? boardNum : cutGetMaxGflopsDeviceId()); */
/*     cublasInit(); */
/*     NVMatrix::initDeviceProps(); */
/*     NVMatrix::initRandom(7); */
/*     uint timer = 0; */
/*     cutilCheckError( cutCreateTimer( &timer)); */

/* //    int imgSize = 32, filterSize = 9; */
/*     int numFilters = 64, numCases = 128; */
/*     int filterPixels = filterSize * filterSize; */
/*     int imgPixels = imgSize * imgSize; */
/*     int numOutputsX = imgSize - filterSize + 1; */
/*     int numOutputs = numOutputsX * numOutputsX; */
/*     assert(numFilters % 16 == 0); */
/*     assert(filterSize % filterCacheX == 0); */
/*     assert(filterSize % filterCacheY == 0); */
/*     printf("Images: %d, filters: %d\n", numCases, numFilters); */
/*     printf("Image size: %dx%d, filter size: %dx%d\n", imgSize, imgSize, filterSize, filterSize); */
/*     printf("Color: no\n"); */

/*     Matrix filters(numFilters, filterPixels); */
/*     Matrix images(numCases, imgPixels); */
/*     Matrix targets(numCases, numFilters * numOutputs); */
/*     filters.randomizeUniform(); */
/*     images.randomizeUniform(); */
/*     targets.apply(Matrix::ZERO); */

/*     NVMatrix nvFilters(filters, true); */
/*     NVMatrix nvImages(images, true); */
/*     NVMatrix nvTargets(targets, true); // eh why not */

/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */
/*     //:IS:convCPU(images.getData(), filters.getData(), targets.getData(), imgSize, filterSize, numCases, numFilters); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("CPU (partial) result:\n"); */
/*     targets.print(0, 3, numOutputsX-10, 6); */
/*     printf("CPU time: %.6f msec\n", cutGetTimerValue(timer)); */


/*     bool checkOutputBounds = numOutputsX % 16 != 0; */
/*     int blocksY = numFilters / 16; */
/*     dim3 threads(16,4,8); */
/*     dim3 blocks(numCases, blocksY); */
/*     assert(filterCacheY == 8); */
/*     assert(filterCacheX == 16); */
/*     int shmem = (16 + filterCacheX - 1)*(4 + filterCacheY - 1) + 16*filterCacheX*filterCacheY; */
/*     assert(shmem < 4096); */
/*     printf("Filter cache size: %dx%d\n", filterCacheX, filterCacheY); */
/*     printf("Using %d bytes of shared memory\n", shmem*4); */
/*     cutilCheckError( cutResetTimer( timer)); */
/*     cutilCheckError( cutStartTimer( timer)); */

/*     if(checkOutputBounds) { */
/*         conv_bw_nofit_4x16_dynfilter_2per<true,8,16, 1><<<blocks, threads>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                     imgSize, filterSize); */
/*     } else { */
/*         conv_bw_nofit_4x16_dynfilter_2per<false,8,16, 1><<<blocks, threads>>>(nvImages.getDevData(), nvFilters.getDevData(), nvTargets.getDevData(), */
/*                                                                      imgSize, filterSize); */
/*     } */

/*     cudaThreadSynchronize(); */
/*     cutilCheckError( cutStopTimer( timer)); */
/*     printf("GPU (partial) result:\n"); */
/*     nvTargets.print(0, 3, numOutputsX-10, 6); */
/*     printf("GPU time: %.6f msec\n", cutGetTimerValue(timer)); */

/*     // Compare results */
/*     Matrix cpuNVTargets(targets); */
/*     nvTargets.copyToHost(cpuNVTargets); */
/*     cpuNVTargets.subtract(targets); */
/*     cpuNVTargets.apply(Matrix::ABS); */
/*     printf("Max diff between CPU/GPU: %.6f\n", cpuNVTargets.max()); */
/* } */

	.text
	.file	"testconv_extras.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
