summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rwxr-xr-xcuda-kernels/Makefile1
-rw-r--r--cuda-kernels/genericMatrixMultiply.cu6
-rwxr-xr-xcuda-kernels/gpgpusim.config2
-rwxr-xr-xcuda-kernels/scripts118
-rw-r--r--cuda-kernels/v16p_kernel.cu4
-rw-r--r--src/Makefile2
-rw-r--r--src/cuda-sim/Makefile2
-rw-r--r--src/cuda-sim/cuda-sim.cc8
-rw-r--r--src/gpgpu-sim/Makefile2
-rw-r--r--src/gpgpu-sim/gpu-sim.cc4
-rw-r--r--src/gpgpu-sim/shader.cc16
-rw-r--r--src/gpgpu-sim/shader.h2
12 files changed, 131 insertions, 36 deletions
diff --git a/cuda-kernels/Makefile b/cuda-kernels/Makefile
index 8effd11..b33519f 100755
--- a/cuda-kernels/Makefile
+++ b/cuda-kernels/Makefile
@@ -7,4 +7,5 @@ clean:
rm _cuob*
rm gpgpusim_power*
rm gpgpu_inst_stats.txt
+ rm gpgpusim_visualizer*
# nvcc -arch=sm_70 --gpu-architecture=compute_50 --gpu-code=compute_50 -lcudart -g -o tensor_core tensor_core.cu
diff --git a/cuda-kernels/genericMatrixMultiply.cu b/cuda-kernels/genericMatrixMultiply.cu
index 6a5d33f..8b96483 100644
--- a/cuda-kernels/genericMatrixMultiply.cu
+++ b/cuda-kernels/genericMatrixMultiply.cu
@@ -40,9 +40,9 @@ void cudaErrCheck_(cudaError_t stat, const char *file, int line) {
using namespace nvcuda;
// Must be multiples of 16 for wmma code to work
-#define MATRIX_M (32)
-#define MATRIX_N (32)
-#define MATRIX_K (32)
+#define MATRIX_M (64)
+#define MATRIX_N (64)
+#define MATRIX_K (64)
diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config
index 3daa539..5f1be25 100755
--- a/cuda-kernels/gpgpusim.config
+++ b/cuda-kernels/gpgpusim.config
@@ -138,7 +138,7 @@
-gpgpu_memlatency_stat 14
-gpgpu_runtime_stat 500
-enable_ptx_file_line_stats 1
--visualizer_enabled 0
+-visualizer_enabled 1
# power model configs
-power_simulation_enabled 1
diff --git a/cuda-kernels/scripts b/cuda-kernels/scripts
index b9a3b50..a24fb72 100755
--- a/cuda-kernels/scripts
+++ b/cuda-kernels/scripts
@@ -10,11 +10,16 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp4_256
./vp4_128>vp4_log128
./vp4_256>vp4_log256
-grep "kernel_name.*vp_example" -A 589 vp4_log16 >vp4_16_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log32 >vp4_32_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log64 >vp4_64_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log128 >vp4_128_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log256 >vp4_256_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log16 >vp4_16_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log32 >vp4_32_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log64 >vp4_64_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log128 >vp4_128_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log256 >vp4_256_summary
+###grep "kernel_name.*vp_example" -A 589 vp4_log16 >>vp4_16_summary
+###grep "kernel_name.*vp_example" -A 589 vp4_log32 >>vp4_32_summary
+###grep "kernel_name.*vp_example" -A 589 vp4_log64 >>vp4_64_summary
+#grep "kernel_name.*vp_example" -A 589 vp4_log128 >>vp4_128_summary
+#grep "kernel_name.*vp_example" -A 589 vp4_log256 >>vp4_256_summary
nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp8_16 v8p_genericMatrixMultiply.cu -DSIZE=16
@@ -28,11 +33,16 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp8_256
./vp8_128>vp8_log128
./vp8_256>vp8_log256
-grep "kernel_name.*vp_example" -A 589 vp8_log16 >vp8_16_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log32 >vp8_32_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log64 >vp8_64_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log128 >vp8_128_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log256 >vp8_256_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log16 >vp8_16_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log32 >vp8_32_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log64 >vp8_64_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log128 >vp8_128_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log256 >vp8_256_summary
+##grep "kernel_name.*vp_example" -A 589 vp8_log16 >>vp8_16_summary
+##grep "kernel_name.*vp_example" -A 589 vp8_log32 >>vp8_32_summary
+##grep "kernel_name.*vp_example" -A 589 vp8_log64 >>vp8_64_summary
+#grep "kernel_name.*vp_example" -A 589 vp8_log128 >>vp8_128_summary
+#grep "kernel_name.*vp_example" -A 589 vp8_log256 >>vp8_256_summary
nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_16 v16p_genericMatrixMultiply.cu -DSIZE=16
nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_32 v16p_genericMatrixMultiply.cu -DSIZE=32
@@ -45,11 +55,85 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_256
./vp16_128>vp16_log128
./vp16_256>vp16_log256
-grep "kernel_name.*vp_example" -A 589 vp16_log16 >vp16_16_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log32 >vp16_32_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log64 >vp16_64_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log128 >vp16_128_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log256 >vp16_256_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log16 >vp16_16_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log32 >vp16_32_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log64 >vp16_64_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log128 >vp16_128_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log256 >vp16_256_summary
+##grep "kernel_name.*vp_example" -A 589 vp16_log16 >>vp16_16_summary
+##grep "kernel_name.*vp_example" -A 589 vp16_log32 >>vp16_32_summary
+##grep "kernel_name.*vp_example" -A 589 vp16_log64 >>vp16_64_summary
+#grep "kernel_name.*vp_example" -A 589 vp16_log128 >>vp16_128_summary
+#grep "kernel_name.*vp_example" -A 589 vp16_log256 >>vp16_256_summary
-mv vp* VPlog/
-make clean
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log16 > kernel_log_vp4_log16
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log32 > kernel_log_vp4_log32
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log64 > kernel_log_vp4_log64
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log128 > kernel_log_vp4_log128
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log256 > kernel_log_vp4_log256
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log16 > kernel_log_vp8_log16
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log32 > kernel_log_vp8_log32
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log64 > kernel_log_vp8_log64
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log128 > kernel_log_vp8_log128
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log256 > kernel_log_vp8_log256
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log16 > kernel_log_vp16_log16
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log32 > kernel_log_vp16_log32
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log64 > kernel_log_vp16_log64
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log128 > kernel_log_vp16_log128
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log256 > kernel_log_vp16_log256
+###
+###grep memtocore:Icnt kernel_log_vp4_log16>memtocore_vp4_log16
+###grep coretomem:Icnt kernel_log_vp4_log16>coretomem_vp4_log16
+###grep MEM_TXN kernel_log_vp4_log16> gen_mem_vp4_log16
+###grep memtocore:Icnt kernel_log_vp4_log32>memtocore_vp4_log32
+###grep coretomem:Icnt kernel_log_vp4_log32>coretomem_vp4_log32
+###grep MEM_TXN kernel_log_vp4_log32> gen_mem_vp4_log32
+###grep memtocore:Icnt kernel_log_vp4_log64>memtocore_vp4_log64
+###grep coretomem:Icnt kernel_log_vp4_log64>coretomem_vp4_log64
+###grep MEM_TXN kernel_log_vp4_log64> gen_mem_vp4_log64
+####grep memtocore:Icnt kernel_log_vp4_log128>memtocore_vp4_log128
+####grep coretomem:Icnt kernel_log_vp4_log128>coretomem_vp4_log128
+####grep MEM_TXN kernel_log_vp4_log128> gen_mem_vp4_log128
+####grep memtocore:Icnt kernel_log_vp4_log256>memtocore_vp4_log256
+####grep coretomem:Icnt kernel_log_vp4_log256>coretomem_vp4_log256
+####grep MEM_TXN kernel_log_vp4_log256> gen_mem_vp4_log256
+###
+###grep memtocore:Icnt kernel_log_vp8_log16>memtocore_vp8_log16
+###grep coretomem:Icnt kernel_log_vp8_log16>coretomem_vp8_log16
+###grep MEM_TXN kernel_log_vp8_log16> gen_mem_vp8_log16
+###grep memtocore:Icnt kernel_log_vp8_log32>memtocore_vp8_log32
+###grep coretomem:Icnt kernel_log_vp8_log32>coretomem_vp8_log32
+###grep MEM_TXN kernel_log_vp8_log32> gen_mem_vp8_log32
+###grep memtocore:Icnt kernel_log_vp8_log64>memtocore_vp8_log64
+###grep coretomem:Icnt kernel_log_vp8_log64>coretomem_vp8_log64
+###grep MEM_TXN kernel_log_vp8_log64> gen_mem_vp8_log64
+####grep memtocore:Icnt kernel_log_vp8_log128>memtocore_vp8_log128
+####grep coretomem:Icnt kernel_log_vp8_log128>coretomem_vp8_log128
+####grep MEM_TXN kernel_log_vp8_log128> gen_mem_vp8_log128
+####grep memtocore:Icnt kernel_log_vp8_log256>memtocore_vp8_log256
+####grep coretomem:Icnt kernel_log_vp8_log256>coretomem_vp8_log256
+####grep MEM_TXN kernel_log_vp8_log256> gen_mem_vp8_log256
+###
+###grep memtocore:Icnt kernel_log_vp16_log16>memtocore_vp16_log16
+###grep coretomem:Icnt kernel_log_vp16_log16>coretomem_vp16_log16
+###grep MEM_TXN kernel_log_vp16_log16> gen_mem_vp16_log16
+###grep memtocore:Icnt kernel_log_vp16_log32>memtocore_vp16_log32
+###grep coretomem:Icnt kernel_log_vp16_log32>coretomem_vp16_log32
+###grep MEM_TXN kernel_log_vp16_log32> gen_mem_vp16_log32
+###grep memtocore:Icnt kernel_log_vp16_log64>memtocore_vp16_log64
+###grep coretomem:Icnt kernel_log_vp16_log64>coretomem_vp16_log64
+###grep MEM_TXN kernel_log_vp16_log64> gen_mem_vp16_log64
+####grep memtocore:Icnt kernel_log_vp16_log128>memtocore_vp16_log128
+####grep coretomem:Icnt kernel_log_vp16_log128>coretomem_vp16_log128
+####grep MEM_TXN kernel_log_vp16_log128> gen_mem_vp16_log128
+####grep memtocore:Icnt kernel_log_vp16_log256>memtocore_vp16_log256
+####grep coretomem:Icnt kernel_log_vp16_log256>coretomem_vp16_log256
+####grep MEM_TXN kernel_log_vp16_log256> gen_mem_vp16_log256
+###
+###mv kernel_* MemTraffic/
+###mv memtocore* MemTraffic/
+###mv coretomem* MemTraffic/
+###mv gen_mem* MemTraffic/
+###
+###mv vp* VPlog/
+###make clean
diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu
index 864597d..d84328f 100644
--- a/cuda-kernels/v16p_kernel.cu
+++ b/cuda-kernels/v16p_kernel.cu
@@ -31,7 +31,7 @@ const int WMMA_M = 16;
const int WMMA_N = 16;
const int WMMA_K = 16;
-__global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int M, int N, int K) {
+__global__ void v16p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int M, int N, int K) {
int registers_a[8];
int registers_b[8];
@@ -294,7 +294,7 @@ int main(int argc, char* argv[]) {
//AAMIR
//AAMIR printf("Running with wmma...\n");
cudaErrCheck(cudaEventRecord(startWMMA));
- v4p_example <<< 1, 32>>> (a_int32, b_int16, c_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K);
+ v16p_example <<< 1, 32>>> (a_int32, b_int16, c_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K);
cudaErrCheck(cudaEventRecord(stopWMMA));
cudaErrCheck(cudaEventSynchronize(stopWMMA));
diff --git a/src/Makefile b/src/Makefile
index 6001669..09194f3 100644
--- a/src/Makefile
+++ b/src/Makefile
@@ -46,7 +46,7 @@ ifeq ($(TRACE),1)
endif
ifneq ($(DEBUG),1)
- OPTFLAGS += -O3
+ OPTFLAGS += -O0
else
CXXFLAGS +=
endif
diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
index e8977a6..d08a696 100644
--- a/src/cuda-sim/Makefile
+++ b/src/cuda-sim/Makefile
@@ -42,7 +42,7 @@ include ../../version_detection.mk
OUTPUT_DIR=$(SIM_OBJ_FILES_DIR)/cuda-sim
-OPT := -O3 -g3 -Wall -Wno-unused-function -Wno-sign-compare
+OPT := -O0 -g3 -Wall -Wno-unused-function -Wno-sign-compare
ifeq ($(DEBUG),1)
OPT := -g3 -Wall -Wno-unused-function -Wno-sign-compare
endif
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 9735a0e..dd7edd9 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -809,7 +809,7 @@ void ptx_instruction::set_opcode_and_latency()
}
if(get_wmma_type()==VP_MMA8)
{
- latency = 128;
+ latency =128;
initiation_interval = 128;
}
if(get_wmma_type()==VP_MMA16)
@@ -1295,12 +1295,14 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
addr_t pc = next_instr();
assert( pc == inst.pc ); // make sure timing model and functional model are in sync
const ptx_instruction *pI = m_func_info->get_instruction(pc);
- #if 0
+ //#if 0
if(lane_id==0){
+ //printf("EXECUTION_FLOW:LINE_NUM:%d\n",pI->source_line());
printf("EXECUTION_FLOW:LINE_NUM:%d\n",pI->source_line());
printf("EXECUTION_FLOW:INST:%s\n",pI->get_source());
+ printf("EXECUTION_FLOW:PC%d\n",pc);
}
- #endif
+ //#endif
set_npc( pc + pI->inst_size() );
diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile
index f10a8a4..4f77699 100644
--- a/src/gpgpu-sim/Makefile
+++ b/src/gpgpu-sim/Makefile
@@ -48,7 +48,7 @@ ifeq ($(GNUC_CPP0X), 1)
endif
ifneq ($(DEBUG),1)
- OPTFLAGS += -O3
+ OPTFLAGS += -O0
else
CXXFLAGS +=
endif
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 7a797b5..eba6f54 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -1469,7 +1469,9 @@ void gpgpu_sim::cycle()
if( g_single_step && ((gpu_sim_cycle+gpu_tot_sim_cycle) >= g_single_step) ) {
asm("int $03");
}
- gpu_sim_cycle++;
+ printf("gpu_sim_cycle=%d\n",gpu_sim_cycle);
+ gpu_sim_cycle++;
+
if( g_interactive_debugger_enabled )
gpgpu_debug();
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index e745f03..eeb59bb 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -1321,10 +1321,16 @@ void ldst_unit::get_L1T_sub_stats(struct cache_sub_stats &css) const{
void shader_core_ctx::warp_inst_complete(const warp_inst_t &inst)
{
- #if 0
- printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%#x @ time=%llu issued@%llu\n",
- inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_tot_sim_cycle + gpu_sim_cycle, inst.get_issue_cycle());
- #endif
+ //#if 0
+ //printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%#x @ time=%llu issued@%llu\n",
+ //const ptx_instruction *pI = m_func_info->get_instruction(inst.pc);
+ printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%d @ time=%llu \n",
+ inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_sim_cycle);
+ //printf("^instruction:%s",(pI->m_source).c_str());
+ //inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_tot_sim_cycle + gpu_sim_cycle, inst.get_issue_cycle());
+ //#endif
+
+
if(inst.op_pipe==SP__OP)
m_stats->m_num_sp_committed[m_sid]++;
else if(inst.op_pipe==SFU__OP)
@@ -1510,7 +1516,6 @@ bool ldst_unit::memory_cycle( warp_inst_t &inst, mem_stage_stall_type &stall_rea
if (m_core->get_config()->gmem_skip_L1D)
bypassL1D = true;
}
-
if( bypassL1D ) {
// bypass L1 cache
unsigned control_size = inst.is_store() ? WRITE_PACKET_SIZE : READ_PACKET_SIZE;
@@ -3472,6 +3477,7 @@ void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf)
case CONST_ACC_R: m_stats->gpgpu_n_mem_const++; break;
case TEXTURE_ACC_R: m_stats->gpgpu_n_mem_texture++; break;
case GLOBAL_ACC_R: m_stats->gpgpu_n_mem_read_global++; break;
+ //case GLOBAL_ACC_R: m_stats->gpgpu_n_mem_read_global++; printf("read_global%d\n",m_stats->gpgpu_n_mem_read_global); break;
case GLOBAL_ACC_W: m_stats->gpgpu_n_mem_write_global++; break;
case LOCAL_ACC_R: m_stats->gpgpu_n_mem_read_local++; break;
case LOCAL_ACC_W: m_stats->gpgpu_n_mem_write_local++; break;
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index d9558b0..99abb63 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1328,7 +1328,7 @@ struct shader_core_config : public core_config
gpgpu_operand_collector_num_in_ports_tensor_core=8;
gpgpu_operand_collector_num_out_ports_tensor_core=8;
- max_vp_core_latency = 64;
+ max_vp_core_latency = 512;
gpgpu_num_vp_core_units=8;
gpgpu_operand_collector_num_units_vp_core=24;
gpgpu_operand_collector_num_in_ports_vp_core=8;
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 01:28:33 +0000