
// /*
//  * Matrix multiplication experiments.
//  */
// #include "matrix_io.cuh"
// #include "mm.cuh"
// #include <stdlib.h>
// #include <stdio.h>
// #include <cublas.h>
// #include <cusparse.h>
// #include <unistd.h>
// #include <ctype.h>

// // #define DEBUG

// int main(int argc, char *argv[])
// {
//     struct Matrix matrix1, matrix2, matrixResCPU, matrixResGPU;
//     int n_elements;
//     const char *filename1;
//     const char *filename2;
//     double time_taken;
//     clock_t start, end;
//     float diff;

//     int time_flag = 0;
//     int repetation_n = 0;
//     int correctness_check_flag = 0;
//     char *alg_type_flag = NULL;

//     filename1 = argv[optind];
//     filename2 = argv[optind + 1];

//     cudaFree(0);
//     read_matrix_dims(filename1, &matrix1, &n_elements);
//     matrix1.vals = (float *)calloc(n_elements, sizeof(float));
//     read_matrix_vals(filename1, &matrix1, 1);

//     read_matrix_dims(filename2, &matrix2, &n_elements);
//     matrix2.vals = (float *)calloc(n_elements, sizeof(float));
//     read_matrix_vals(filename2, &matrix2, 1);

//     int m = matrix1.dims[2];
//     int k = matrix1.dims[3];
//     int n = matrix2.dims[3];

//     if (correctness_check_flag)
//     {
//         //cpu_mm
//         initiliaze2dMatrix(&matrixResCPU, m, n);
//         start = clock();
//         cpu_mm(&matrix1, &matrix2, &matrixResCPU, m, n, k);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for the cpu_mm is %lf\n", time_taken);
//     }

//     if (strcmp(alg_type_flag, "denseblas") == 0)
//     {
//         ///gpu_mm_dense
//         // Create a handle for CUBLAS
//         initiliaze2dMatrix(&matrixResGPU, m, n);
//         cublasHandle_t handleBLAS;
//         cublasCreate(&handleBLAS);

//         start = clock();
//         gpu_mm_dense(&matrix1, &matrix2, &matrixResGPU, m, n, k, handleBLAS, time_flag, repetation_n);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for the gpu_mm_dense is %lf\n", time_taken);
// #ifdef DEBUG
//         printf("CuBLAS result:\n");
//         print_matrix(&matrixResGPU);
// #endif
//         if (correctness_check_flag)
//         {
//             diff = calculateDistanceMatrix(&matrixResCPU, &matrixResGPU);
//             if (diff > 1e-7)
//             {
//                 printf("Diff is %.8f\n", diff);
//                 printf("There might be a problem\n");
//             }
//             else
//             {
//                 printf("Diff is less then 1e-7\n", diff);
//             }
//         }
//         destroyMatrix(&matrixResGPU);
//         cublasDestroy(handleBLAS);
//     }
//     else if (strcmp(alg_type_flag, "cusparse") == 0)
//     {
//         ///gpu_mm_sparse
//         // Initialize cusparse library
//         initiliaze2dMatrix(&matrixResGPU, m, n);
//         cusparseHandle_t handleSparse;
//         cusparseCreate(&handleSparse);

//         start = clock();
//         gpu_mm_sparse(&matrix1, &matrix2, &matrixResGPU, m, n, k, handleSparse, time_flag, repetation_n);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for gpu_mm_sparse is %lf\n", time_taken);
// #ifdef DEBUG
//         printf("cuSparse result:\n");
//         print_matrix(&matrixResGPU);
// #endif
//         if (correctness_check_flag)
//         {
//             diff = calculateDistanceMatrix(&matrixResCPU, &matrixResGPU);
//             if (diff > 1e-7)
//             {
//                 printf("Diff is %.8f\n", diff);
//                 printf("There might be a problem\n");
//             }
//             else
//             {
//                 printf("Diff is less then 1e-7\n", diff);
//             }
//         }
//         cusparseDestroy(handleSparse);
//         destroyMatrix(&matrixResGPU);
//     }
//     else if (strcmp(alg_type_flag, "sparseimp") == 0)
//     {
//         //gpu_mm_sparse_ProjectImp
//         initiliaze2dMatrix(&matrixResGPU, m, n);
//         cusparseHandle_t handleSparse;
//         cusparseCreate(&handleSparse);
//         start = clock();
//         gpu_mm_sparse_ProjectImp(&matrix1, &matrix2, &matrixResGPU, m, n, k, handleSparse, time_flag, repetation_n);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for gpu_mm_sparse_ProjectImp is %lf\n", time_taken);
// #ifdef DEBUG
//         printf("cuSparse result:\n");
//         print_matrix(&matrixResGPU);
// #endif
//         if (correctness_check_flag)
//         {
//             diff = calculateDistanceMatrix(&matrixResCPU, &matrixResGPU);
//             if (diff > 1e-7)
//             {
//                 printf("Diff is %.8f\n", diff);
//                 printf("There might be a problem\n");
//             }
//             else
//             {
//                 printf("Diff is less then 1e-7\n", diff);
//             }
//         }
//         cusparseDestroy(handleSparse);
//         destroyMatrix(&matrixResGPU);
//     }
//     else
//     {
//         printf("Use denseblas/cusparse/sparseimp with -a flag.\n");
//     }

//     destroyMatrix(&matrix1);
//     destroyMatrix(&matrix2);
//     if (correctness_check_flag)
//     {
//         destroyMatrix(&matrixResCPU);
//     }
// }
	.file	"tmpxft_003463de_00000000-6_sdnn_blas.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>

// /*
//  * Matrix multiplication experiments.
//  */
// #include "matrix_io.cuh"
// #include "mm.cuh"
// #include <stdlib.h>
// #include <stdio.h>
// #include <cublas.h>
// #include <cusparse.h>
// #include <unistd.h>
// #include <ctype.h>

// // #define DEBUG

// int main(int argc, char *argv[])
// {
//     struct Matrix matrix1, matrix2, matrixResCPU, matrixResGPU;
//     int n_elements;
//     const char *filename1;
//     const char *filename2;
//     double time_taken;
//     clock_t start, end;
//     float diff;

//     int time_flag = 0;
//     int repetation_n = 0;
//     int correctness_check_flag = 0;
//     char *alg_type_flag = NULL;

//     filename1 = argv[optind];
//     filename2 = argv[optind + 1];

//     cudaFree(0);
//     read_matrix_dims(filename1, &matrix1, &n_elements);
//     matrix1.vals = (float *)calloc(n_elements, sizeof(float));
//     read_matrix_vals(filename1, &matrix1, 1);

//     read_matrix_dims(filename2, &matrix2, &n_elements);
//     matrix2.vals = (float *)calloc(n_elements, sizeof(float));
//     read_matrix_vals(filename2, &matrix2, 1);

//     int m = matrix1.dims[2];
//     int k = matrix1.dims[3];
//     int n = matrix2.dims[3];

//     if (correctness_check_flag)
//     {
//         //cpu_mm
//         initiliaze2dMatrix(&matrixResCPU, m, n);
//         start = clock();
//         cpu_mm(&matrix1, &matrix2, &matrixResCPU, m, n, k);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for the cpu_mm is %lf\n", time_taken);
//     }

//     if (strcmp(alg_type_flag, "denseblas") == 0)
//     {
//         ///gpu_mm_dense
//         // Create a handle for CUBLAS
//         initiliaze2dMatrix(&matrixResGPU, m, n);
//         cublasHandle_t handleBLAS;
//         cublasCreate(&handleBLAS);

//         start = clock();
//         gpu_mm_dense(&matrix1, &matrix2, &matrixResGPU, m, n, k, handleBLAS, time_flag, repetation_n);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for the gpu_mm_dense is %lf\n", time_taken);
// #ifdef DEBUG
//         printf("CuBLAS result:\n");
//         print_matrix(&matrixResGPU);
// #endif
//         if (correctness_check_flag)
//         {
//             diff = calculateDistanceMatrix(&matrixResCPU, &matrixResGPU);
//             if (diff > 1e-7)
//             {
//                 printf("Diff is %.8f\n", diff);
//                 printf("There might be a problem\n");
//             }
//             else
//             {
//                 printf("Diff is less then 1e-7\n", diff);
//             }
//         }
//         destroyMatrix(&matrixResGPU);
//         cublasDestroy(handleBLAS);
//     }
//     else if (strcmp(alg_type_flag, "cusparse") == 0)
//     {
//         ///gpu_mm_sparse
//         // Initialize cusparse library
//         initiliaze2dMatrix(&matrixResGPU, m, n);
//         cusparseHandle_t handleSparse;
//         cusparseCreate(&handleSparse);

//         start = clock();
//         gpu_mm_sparse(&matrix1, &matrix2, &matrixResGPU, m, n, k, handleSparse, time_flag, repetation_n);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for gpu_mm_sparse is %lf\n", time_taken);
// #ifdef DEBUG
//         printf("cuSparse result:\n");
//         print_matrix(&matrixResGPU);
// #endif
//         if (correctness_check_flag)
//         {
//             diff = calculateDistanceMatrix(&matrixResCPU, &matrixResGPU);
//             if (diff > 1e-7)
//             {
//                 printf("Diff is %.8f\n", diff);
//                 printf("There might be a problem\n");
//             }
//             else
//             {
//                 printf("Diff is less then 1e-7\n", diff);
//             }
//         }
//         cusparseDestroy(handleSparse);
//         destroyMatrix(&matrixResGPU);
//     }
//     else if (strcmp(alg_type_flag, "sparseimp") == 0)
//     {
//         //gpu_mm_sparse_ProjectImp
//         initiliaze2dMatrix(&matrixResGPU, m, n);
//         cusparseHandle_t handleSparse;
//         cusparseCreate(&handleSparse);
//         start = clock();
//         gpu_mm_sparse_ProjectImp(&matrix1, &matrix2, &matrixResGPU, m, n, k, handleSparse, time_flag, repetation_n);
//         end = clock();
//         time_taken = ((double)(end - start)) / CLOCKS_PER_SEC;
//         printf("Time taken for gpu_mm_sparse_ProjectImp is %lf\n", time_taken);
// #ifdef DEBUG
//         printf("cuSparse result:\n");
//         print_matrix(&matrixResGPU);
// #endif
//         if (correctness_check_flag)
//         {
//             diff = calculateDistanceMatrix(&matrixResCPU, &matrixResGPU);
//             if (diff > 1e-7)
//             {
//                 printf("Diff is %.8f\n", diff);
//                 printf("There might be a problem\n");
//             }
//             else
//             {
//                 printf("Diff is less then 1e-7\n", diff);
//             }
//         }
//         cusparseDestroy(handleSparse);
//         destroyMatrix(&matrixResGPU);
//     }
//     else
//     {
//         printf("Use denseblas/cusparse/sparseimp with -a flag.\n");
//     }

//     destroyMatrix(&matrix1);
//     destroyMatrix(&matrix2);
//     if (correctness_check_flag)
//     {
//         destroyMatrix(&matrixResCPU);
//     }
// }
	.text
	.file	"sdnn_blas.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
