diff --git a/.DS_Store b/.DS_Store deleted file mode 100644 index 62bf31b..0000000 Binary files a/.DS_Store and /dev/null differ diff --git a/HW4_102062111.zip b/HW4_102062111.zip deleted file mode 100644 index 6dbf965..0000000 Binary files a/HW4_102062111.zip and /dev/null differ diff --git a/HW4_102062111/.DS_Store b/HW4_102062111/.DS_Store deleted file mode 100644 index a0b61bd..0000000 Binary files a/HW4_102062111/.DS_Store and /dev/null differ diff --git a/HW4_102062111/HW4_102062111_cuda.cu b/HW4_102062111/HW4_102062111_cuda.cu deleted file mode 100644 index a697578..0000000 --- a/HW4_102062111/HW4_102062111_cuda.cu +++ /dev/null @@ -1,296 +0,0 @@ -#include -#include -#include -#include -#include -#include - -#define INF 1e9 -#define H2D cudaMemcpyHostToDevice -#define D2H cudaMemcpyDeviceToHost -#define timer(type) \ - cudaEventCreate(&type##_start); \ - cudaEventCreate(&type##_stop); - -#define record(type) \ - cudaEventRecord(type) - -#define elapsed(type, start, stop) \ - cudaEventElapsedTime(&type, start, stop) - -#define sync(type) \ - cudaEventSynchronize(type) - -#define MASK_N 2 -#define MASK_X 5 -#define MASK_Y 5 -#define SCALE 8 - -#define cudaCheckErrors(msg) \ - do { \ - cudaError_t __err = cudaGetLastError(); \ - if (__err != cudaSuccess) { \ - fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ - msg, cudaGetErrorString(__err), \ - __FILE__, __LINE__); \ - fprintf(stderr, "*** FAILED - ABORTING\n"); \ - exit(1); \ - } \ - } while (0) - -using namespace std; - -void DisplayHeader() -{ - const int kb = 1024; - const int mb = kb * kb; - wcout << "NBody.GPU" << endl << "=========" << endl << endl; - - //wcout << "CUDA version: v" << CUDART_VERSION << endl; - //wcout << "Thrust version: v" << THRUST_MAJOR_VERSION << "." << THRUST_MINOR_VERSION << endl << endl; - - int devCount; - cudaGetDeviceCount(&devCount); - wcout << "CUDA Devices: " << endl << endl; - - for(int i = 0; i < devCount; ++i) - { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; - - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; - } -} - - -__global__ void floyd_warshall(int *d, int blockSize, int length, - int XIndex, int YIndex, int kk) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -} - - -int main(int argc, char **argv) { - cudaEvent_t total_start, total_stop; - cudaEvent_t com_start, com_stop; - cudaEvent_t mem_start, mem_stop; - cudaEvent_t io_start, io_stop; - - timer(total); - timer(com); - timer(mem); - timer(io); - - cudaEventRecord(total_start); - - float total, compute, memory, IO; - float mem_part, IO_part; - total = compute = memory = IO = 0; - - FILE *fin = fopen(argv[1], "r"); - FILE *fout = fopen(argv[2], "w"); - int blockSize = atoi(argv[3]); - - int N, M; - int *edge; - int *cuda_edge; - - fscanf(fin, "%d %d", &N, &M); - int gridSize = N % blockSize ? N / blockSize + 1 : N / blockSize; - int length = blockSize * gridSize + 1; - edge = new int[length * length]; - - //fprintf(stderr, "grid = %d, block = %d\n", gridSize, blockSize); - - for (int i = 0; i < length; ++i) { - for (int j = 0; j < length; ++j) - edge[i * length + j] = INF; - edge[i * length + i] = 0; - } - - - record(io_start); - while (M--) { - int a, b, w; - fscanf(fin, "%d %d %d", &a, &b, &w); - a = a - 1; - b = b - 1; - edge[a * length + b] = w; - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - cudaSetDevice(0); - - cudaMalloc((void**)&cuda_edge, sizeof(int) * length * length); - cudaCheckErrors("malloc cuda_edge"); - - record(mem_start); - cudaMemcpy(cuda_edge, edge, sizeof(int) * length * length, H2D); - cudaCheckErrors("copy cuda_edge"); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - // Now only hangle N = 3200 testcase - // size_t sharedSize = 8 * 8; - int blockNum = (N + blockSize - 1) / blockSize; - int blockDimension = 8; - gridSize = blockSize / blockDimension; - int gridFactor = 1024 / blockSize ; - gridFactor *= gridFactor; - - - dim3 blocks(gridSize, gridSize); - dim3 threads(blockDimension, blockDimension); - - dim3 blockCol(gridSize * gridFactor, gridSize); - dim3 blockRow(gridSize, gridSize * gridFactor); - int remainBegin = (blockNum / gridFactor) * gridFactor ; - int remain = blockNum - remainBegin; - dim3 blockColRemain(gridSize * remain, gridSize); - dim3 blockRowRemain(gridSize, gridSize * remain); - //wcout << "blocknum = " << blockNum << endl; - - cudaEventRecord(com_start); - for (int k = 0; k < blockNum; ++k) { - //wcout << "k = " << k << endl; - // phase one - { - for (int cur = 0; cur < blockSize; ++cur) { - //wcout << "(" << cur << "/" << blockSize << endl; - floyd_warshall<<>> - (cuda_edge, blockSize, length, k, k, k * blockSize + cur); - - cudaCheckErrors("phase one"); - } - } - // phase two - { - // Column - for (int i = 0; i < blockNum - remain; i = i + gridFactor) { - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, i, k, k * blockSize + cur); - - cudaCheckErrors("phase two column main"); - } - } - - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, remainBegin, k, k * blockSize + cur); - cudaCheckErrors("phase two column remain"); - } - // Row - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, k, j, k * blockSize + cur); - cudaCheckErrors("phase two row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, k, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase two row remain"); - } - } - - //phase three - { - - for (int i = 0; i < blockNum; i++) { - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, i, j, k * blockSize + cur); - cudaCheckErrors("phase three row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, i, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase three row remain"); - } - } - - /* - for (int i = 0; i < blockNum; ++i) { - for (int j = 0; j < blockNum; ++j) - if (i != k && j != k) - for (int cur = 0; cur < blockSize; ++cur) - floyd_warshall<<>> - (cuda_edge, blockSize, length, i, j, k * blockSize + cur); - } - */ - - } - - } - cudaDeviceSynchronize(); - cudaEventRecord(com_stop); - cudaEventSynchronize(com_stop); - cudaEventElapsedTime(&compute, com_start, com_stop); - - - record(mem_start); - cudaMemcpy(edge, cuda_edge, sizeof(int) * length * length, D2H); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - - record(io_start); - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N - 1; ++j) { - if (edge[i * length + j] == INF) - fprintf(fout, "INF "); - else - fprintf(fout, "%d ", edge[i * length + j]); - } - if (edge[i * length + N - 1] == INF) - fprintf(fout, "INF\n"); - else - fprintf(fout, "%d\n", edge[i * length + N - 1]); - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - cudaEventRecord(total_stop); - cudaEventSynchronize(total_stop); - cudaEventElapsedTime(&total, total_start, total_stop); - - fprintf(stderr, "\n\n"); - fprintf(stderr, "TOTAL = %f\n", total); - fprintf(stderr, "COMPUTE = %f\n", compute); - fprintf(stderr, "MEMORY = %f\n", memory); - fprintf(stderr, "IO = %f\n", IO); - - return 0; -} diff --git a/HW4_102062111/HW4_102062111_fast.cu b/HW4_102062111/HW4_102062111_fast.cu deleted file mode 100644 index a65f84d..0000000 --- a/HW4_102062111/HW4_102062111_fast.cu +++ /dev/null @@ -1,311 +0,0 @@ -#include -#include -#include -#include -#include -#include - -#define INF 1e9 -#define H2D cudaMemcpyHostToDevice -#define D2H cudaMemcpyDeviceToHost -#define timer(type) \ - cudaEventCreate(&type##_start); \ - cudaEventCreate(&type##_stop); - -#define record(type) \ - cudaEventRecord(type) - -#define elapsed(type, start, stop) \ - cudaEventElapsedTime(&type, start, stop) - -#define sync(type) \ - cudaEventSynchronize(type) - -#define MASK_N 2 -#define MASK_X 5 -#define MASK_Y 5 -#define SCALE 8 - -#define cudaCheckErrors(msg) \ - do { \ - cudaError_t __err = cudaGetLastError(); \ - if (__err != cudaSuccess) { \ - fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ - msg, cudaGetErrorString(__err), \ - __FILE__, __LINE__); \ - fprintf(stderr, "*** FAILED - ABORTING\n"); \ - exit(1); \ - } \ - } while (0) - -using namespace std; - -void DisplayHeader() -{ - const int kb = 1024; - const int mb = kb * kb; - wcout << "NBody.GPU" << endl << "=========" << endl << endl; - - //wcout << "CUDA version: v" << CUDART_VERSION << endl; - //wcout << "Thrust version: v" << THRUST_MAJOR_VERSION << "." << THRUST_MINOR_VERSION << endl << endl; - - int devCount; - cudaGetDeviceCount(&devCount); - wcout << "CUDA Devices: " << endl << endl; - - for(int i = 0; i < devCount; ++i) - { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; - - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; - } -} - - -__global__ void floyd_warshall(int *d, int blockSize, int length, - int XIndex, int YIndex, int kk) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -} - -__global__ void floyd_warshall_whole(int *d, int blockSize, int length, - int XIndex, int YIndex, int ZIndex) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - //__shared__ int dij[8][8]; - - // dij[threadIdx.x][threadIdx.y] = d[ii * length + jj]; - for (int cur = 0; cur < blockSize; ++cur) { - int kk = ZIndex + cur; - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -/* - if (dij[threadIdx.x][threadIdx.y] > dik + dkj) - dij[threadIdx.x][threadIdx.y] = dik + dkj; -*/ - } - // d[ii * length + jj] = dij[threadIdx.x][threadIdx.y]; -} - - -int main(int argc, char **argv) { - cudaEvent_t total_start, total_stop; - cudaEvent_t com_start, com_stop; - cudaEvent_t mem_start, mem_stop; - cudaEvent_t io_start, io_stop; - - timer(total); - timer(com); - timer(mem); - timer(io); - - cudaEventRecord(total_start); - - float total, compute, memory, IO; - float mem_part, IO_part; - total = compute = memory = IO = 0; - - FILE *fin = fopen(argv[1], "r"); - FILE *fout = fopen(argv[2], "w"); - int blockSize = atoi(argv[3]); - - int N, M; - int *edge; - int *cuda_edge; - - fscanf(fin, "%d %d", &N, &M); - int gridSize = N % blockSize ? N / blockSize + 1 : N / blockSize; - int length = blockSize * gridSize + 1; - edge = new int[length * length]; - - //fprintf(stderr, "grid = %d, block = %d\n", gridSize, blockSize); - - for (int i = 0; i < length; ++i) { - for (int j = 0; j < length; ++j) - edge[i * length + j] = INF; - edge[i * length + i] = 0; - } - - - record(io_start); - while (M--) { - int a, b, w; - fscanf(fin, "%d %d %d", &a, &b, &w); - a = a - 1; - b = b - 1; - edge[a * length + b] = w; - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - cudaSetDevice(0); - - cudaMalloc((void**)&cuda_edge, sizeof(int) * length * length); - cudaCheckErrors("malloc cuda_edge"); - - record(mem_start); - cudaMemcpy(cuda_edge, edge, sizeof(int) * length * length, H2D); - cudaCheckErrors("copy cuda_edge"); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - // Now only hangle N = 3200 testcase - // size_t sharedSize = 8 * 8; - int blockNum = (N + blockSize - 1) / blockSize; - int blockDimension = 8; - gridSize = blockSize / blockDimension; - int gridFactor = 1024 / blockSize ; - gridFactor *= gridFactor; - - - dim3 blocks(gridSize, gridSize); - dim3 threads(blockDimension, blockDimension); - - dim3 blockCol(gridSize * gridFactor, gridSize); - dim3 blockRow(gridSize, gridSize * gridFactor); - int remainBegin = (blockNum / gridFactor) * gridFactor ; - int remain = blockNum - remainBegin; - dim3 blockColRemain(gridSize * remain, gridSize); - dim3 blockRowRemain(gridSize, gridSize * remain); - //wcout << "blocknum = " << blockNum << endl; - - cudaEventRecord(com_start); - for (int k = 0; k < blockNum; ++k) { - //wcout << "k = " << k << endl; - // phase one - { - for (int cur = 0; cur < blockSize; ++cur) { - //wcout << "(" << cur << "/" << blockSize << endl; - floyd_warshall<<>> - (cuda_edge, blockSize, length, k, k, k * blockSize + cur); - - cudaCheckErrors("phase one"); - } - } - // phase two - { - // Column - for (int i = 0; i < blockNum - remain; i = i + gridFactor) { - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, i, k, k * blockSize + cur); - - cudaCheckErrors("phase two column main"); - } - } - - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, remainBegin, k, k * blockSize + cur); - cudaCheckErrors("phase two column remain"); - } - // Row - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, k, j, k * blockSize + cur); - cudaCheckErrors("phase two row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (cuda_edge, blockSize, length, k, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase two row remain"); - } - } - - //phase three - { - - for (int i = 0; i < blockNum; i++) { - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - int cur = 0; - //for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall_whole<<>> - (cuda_edge, blockSize, length, i, j, k * blockSize + cur); - cudaCheckErrors("phase three row main"); - //} - } - if (remainBegin < blockNum) { - int cur = 0; - //for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall_whole<<>> - (cuda_edge, blockSize, length, i, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase three row remain"); - //} - } - } - } - - } - cudaDeviceSynchronize(); - cudaEventRecord(com_stop); - cudaEventSynchronize(com_stop); - cudaEventElapsedTime(&compute, com_start, com_stop); - - - record(mem_start); - cudaMemcpy(edge, cuda_edge, sizeof(int) * length * length, D2H); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - - record(io_start); - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N - 1; ++j) { - if (edge[i * length + j] == INF) - fprintf(fout, "INF "); - else - fprintf(fout, "%d ", edge[i * length + j]); - } - if (edge[i * length + N - 1] == INF) - fprintf(fout, "INF\n"); - else - fprintf(fout, "%d\n", edge[i * length + N - 1]); - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - cudaEventRecord(total_stop); - cudaEventSynchronize(total_stop); - cudaEventElapsedTime(&total, total_start, total_stop); - - fprintf(stderr, "\n\n"); - fprintf(stderr, "TOTAL = %f\n", total); - fprintf(stderr, "COMPUTE = %f\n", compute); - fprintf(stderr, "MEMORY = %f\n", memory); - fprintf(stderr, "IO = %f\n", IO); - - return 0; -} diff --git a/HW4_102062111/HW4_102062111_fast_mpi.cu b/HW4_102062111/HW4_102062111_fast_mpi.cu deleted file mode 100644 index 64b95ce..0000000 --- a/HW4_102062111/HW4_102062111_fast_mpi.cu +++ /dev/null @@ -1,389 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include -#include - -#define INF 1e9 -#define H2D cudaMemcpyHostToDevice -#define D2H cudaMemcpyDeviceToHost -#define D2D cudaMemcpyDeviceToDevice -#define COMM MPI_COMM_WORLD - -#define send(buffer, count, dest) \ - MPI_Send(buffer, count, MPI_CHAR, dest, 0, COMM) - -#define recv(buffer, count, src, status)\ - MPI_Recv(buffer, count, MPI_CHAR, src, MPI_ANY_TAG, COMM, &status) - -#define MASK_N 2 -#define MASK_X 5 -#define MASK_Y 5 -#define SCALE 8 -#define ROOT 0 - -#define timer(type) \ - cudaEventCreate(&type##_start); \ - cudaEventCreate(&type##_stop); - -#define record(type) \ - cudaEventRecord(type) - -#define elapsed(type, start, stop) \ - cudaEventElapsedTime(&type, start, stop) - -#define sync(type) \ - cudaEventSynchronize(type) - -#define cudaCheckErrors(msg) \ - do { \ - cudaError_t __err = cudaGetLastError(); \ - if (__err != cudaSuccess) { \ - fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ - msg, cudaGetErrorString(__err), \ - __FILE__, __LINE__); \ - fprintf(stderr, "*** FAILED - ABORTING\n"); \ - exit(1); \ - } \ - } while (0) - -using namespace std; - -void DisplayHeader() -{ - const int kb = 1024; - const int mb = kb * kb; - wcout << "NBody.GPU" << endl << "=========" << endl << endl; - - //wcout << "CUDA version: v" << CUDART_VERSION << endl; - //wcout << "Thrust version: v" << THRUST_MAJOR_VERSION << "." << THRUST_MINOR_VERSION << endl << endl; - - int devCount; - cudaGetDeviceCount(&devCount); - wcout << "CUDA Devices: " << endl << endl; - - for(int i = 0; i < devCount; ++i) - { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; - - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; - } -} - -__global__ void floyd_warshall(int *d, int blockSize, int length, - int XIndex, int YIndex, int kk) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -} - -__global__ void floyd_warshall_whole(int *d, int blockSize, int length, - int XIndex, int YIndex, int ZIndex) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - //__shared__ int dij[8][8]; - - // dij[threadIdx.x][threadIdx.y] = d[ii * length + jj]; - for (int cur = 0; cur < blockSize; ++cur) { - int kk = ZIndex + cur; - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -/* - if (dij[threadIdx.x][threadIdx.y] > dik + dkj) - dij[threadIdx.x][threadIdx.y] = dik + dkj; -*/ - } - // d[ii * length + jj] = dij[threadIdx.x][threadIdx.y]; -} - -int main(int argc, char **argv) { - - int rank, size; - MPI_Init(&argc, &argv); - MPI_Comm_rank(MPI_COMM_WORLD, &rank); - MPI_Comm_size(MPI_COMM_WORLD, &size); - - cudaSetDevice(rank); - - cudaEvent_t total_start, total_stop; - cudaEvent_t com_start, com_stop; - cudaEvent_t mem_start, mem_stop; - cudaEvent_t io_start, io_stop; - cudaEvent_t trans_start, trans_stop; - - timer(total); - timer(com); - timer(mem); - timer(io); - timer(trans); - - cudaEventRecord(total_start); - - float total, compute, memory, IO, trans; - float mem_part, IO_part, com_part, trans_part; - total = compute = memory = IO = trans = 0; - - // fprintf(stderr, "rank = %d, size = %d\n", rank, size); - - MPI_Status status; - - int blockSize = atoi(argv[3]); - - int N, M; - int *edge; - int *gpu[2]; - int gridSize; - int length; - const int deviceNum = 2; - - if (rank == ROOT) { - FILE *fin = fopen(argv[1], "r"); - - fscanf(fin, "%d %d", &N, &M); - gridSize = N % blockSize ? N / blockSize + 1 : N / blockSize; - length = blockSize * gridSize + 1; - cudaMallocHost((void**)&edge, sizeof(int) * length * length); - for (int i = 0; i < length; ++i) { - for (int j = 0; j < length; ++j) - edge[i * length + j] = INF; - edge[i * length + i] = 0; - } - - record(io_start); - while (M--) { - int a, b, w; - fscanf(fin, "%d %d %d", &a, &b, &w); - a = a - 1; - b = b - 1; - edge[a * length + b] = w; - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - send(&N, sizeof(int), 1); - send(&gridSize, sizeof(int), 1); - send(&length, sizeof(int), 1); - send(edge, sizeof(int) * length * length, 1); - - fclose(fin); - } else { - recv(&N, sizeof(int), ROOT, status); - recv(&gridSize, sizeof(int), ROOT, status); - recv(&length, sizeof(int), ROOT, status); - cudaMallocHost((void**)&edge, sizeof(int) * length * length); - recv(edge, sizeof(int) * length * length, ROOT, status); - //fprintf(stderr, "RANK %d, length = %d, gridSize = %d\n", rank, length, gridSize); - } - - cudaSetDevice(rank); - cudaMalloc((void**)&gpu[rank], sizeof(int) * length * length); - - if(rank == 0) record(mem_start); - cudaMemcpy(gpu[rank], edge, sizeof(int) * length * length, H2D); - cudaCheckErrors("memcpy error"); - if (rank == 0) { - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - } - size_t sharedSize = 8 * 8; - int blockNum = (N + blockSize - 1) / blockSize; - int blockDimension = 8; - gridSize = blockSize / blockDimension; - int gridFactor = 1024 / blockSize ; - gridFactor *= gridFactor; - - dim3 blocks(gridSize, gridSize); - dim3 threads(blockDimension, blockDimension); - - dim3 blockCol(gridSize * gridFactor, gridSize); - dim3 blockRow(gridSize, gridSize * gridFactor); - int remainBegin = (blockNum / gridFactor) * gridFactor ; - int remain = blockNum - remainBegin; - dim3 blockColRemain(gridSize * remain, gridSize); - dim3 blockRowRemain(gridSize, gridSize * remain); - - // fprintf(stderr, "rank %d, gpu = %p \n", rank, gpu[rank]); - - - for (int k = 0; k < blockNum; ++k) { - record(com_start); - //wcout << "rank = " << rank << "k = " << k << ", blocknum = " << blockNum << endl; - // phase one - { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, k, k, k * blockSize + cur); - cudaCheckErrors("phase one"); - } - } - // phase two - { - // Column - for (int i = 0; i < blockNum - remain; i = i + gridFactor) { - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, i, k, k * blockSize + cur); - cudaCheckErrors("phase two column main"); - } - } - - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, remainBegin, k, k * blockSize + cur); - cudaCheckErrors("phase two column remain"); - } - // Row - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, k, j, k * blockSize + cur); - cudaCheckErrors("phase two row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, k, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase two row remain"); - } - } - - //phase three - { - int thread = rank; - int begin, end; - cudaSetDevice(thread); - - if (thread == 0) { - begin = 0; - end = blockNum / 2; - } else { - begin = blockNum / 2; - end = blockNum; - } - for (int i = begin; i < end; i++) { - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < 1; ++cur) { - floyd_warshall_whole<<>> - (gpu[thread], blockSize, length, i, j, k * blockSize + cur); - cudaCheckErrors("phase three row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < 1; cur++) { - floyd_warshall_whole<<>> - (gpu[thread], blockSize, length, i, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase three row remain"); - } - } - } - - cudaDeviceSynchronize(); - - record(com_stop); - sync(com_stop); - elapsed(com_part, com_start, com_stop); - compute += com_part; - - int offset = (blockNum / 2) * blockSize * length ; - int copySize = length * length - offset; - if (rank == 0) { - record(mem_start); - cudaMemcpy(edge, gpu[0], sizeof(int) * offset, D2H); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - - record(trans_start); - send(edge, sizeof(int) * offset, 1); - recv(edge + offset, sizeof(int) * copySize, 1, status); - record(trans_stop); - sync(trans_stop); - elapsed(trans_part, trans_start, trans_stop); - trans += trans_part; - - record(mem_start); - cudaMemcpy(gpu[0] + offset, edge + offset, sizeof(int) * copySize, H2D); - cudaCheckErrors("rank ROOT memcpy D2D"); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - } else { - recv(edge, sizeof(int) * offset, ROOT, status); - cudaMemcpy(gpu[1], edge, sizeof(int) * offset, H2D); - cudaMemcpy(edge + offset, gpu[1] + offset, sizeof(int) * copySize, D2H); - cudaCheckErrors("rank 1 memcpy D2D"); - send(edge + offset, sizeof(int) * copySize, ROOT); - } - } - if (rank == 0) { - record(io_start); - FILE *fout = fopen(argv[2], "w"); - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N - 1; ++j) { - if (edge[i * length + j] == INF) - fprintf(fout, "INF "); - else - fprintf(fout, "%d ", edge[i * length + j]); - } - if (edge[i * length + N - 1] == INF) - fprintf(fout, "INF\n"); - else - fprintf(fout, "%d\n", edge[i * length + N - 1]); - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - fclose(fout); - } - - MPI_Barrier(MPI_COMM_WORLD); - - cudaEventRecord(total_stop); - cudaEventSynchronize(total_stop); - cudaEventElapsedTime(&total, total_start, total_stop); - if (rank == 0) { - fprintf(stderr, "\n\n"); - fprintf(stderr, "TOTAL = %f\n", total); - fprintf(stderr, "COMPUTE = %f\n", compute); - fprintf(stderr, "COMMUNICATE = %f\n", trans); - fprintf(stderr, "MEMORY = %f\n", memory); - fprintf(stderr, "IO = %f\n", IO); - } - MPI_Barrier(MPI_COMM_WORLD); - MPI_Finalize(); - return 0; -} diff --git a/HW4_102062111/HW4_102062111_fast_openmp.cu b/HW4_102062111/HW4_102062111_fast_openmp.cu deleted file mode 100644 index 8bbfad0..0000000 --- a/HW4_102062111/HW4_102062111_fast_openmp.cu +++ /dev/null @@ -1,350 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include - -#define INF 1e9 -#define H2D cudaMemcpyHostToDevice -#define D2H cudaMemcpyDeviceToHost -#define D2D cudaMemcpyDeviceToDevice -#define MASK_N 2 -#define MASK_X 5 -#define MASK_Y 5 -#define SCALE 8 - -#define timer(type) \ - cudaEventCreate(&type##_start); \ - cudaEventCreate(&type##_stop); - -#define record(type) \ - cudaEventRecord(type) - -#define elapsed(type, start, stop) \ - cudaEventElapsedTime(&type, start, stop) - -#define sync(type) \ - cudaEventSynchronize(type) - -#define cudaCheckErrors(msg) \ - do { \ - cudaError_t __err = cudaGetLastError(); \ - if (__err != cudaSuccess) { \ - fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ - msg, cudaGetErrorString(__err), \ - __FILE__, __LINE__); \ - fprintf(stderr, "*** FAILED - ABORTING\n"); \ - exit(1); \ - } \ - } while (0) - -using namespace std; - -void DisplayHeader() -{ - const int kb = 1024; - const int mb = kb * kb; - wcout << "NBody.GPU" << endl << "=========" << endl << endl; - - //wcout << "CUDA version: v" << CUDART_VERSION << endl; - //wcout << "Thrust version: v" << THRUST_MAJOR_VERSION << "." << THRUST_MINOR_VERSION << endl << endl; - - int devCount; - cudaGetDeviceCount(&devCount); - wcout << "CUDA Devices: " << endl << endl; - - for(int i = 0; i < devCount; ++i) - { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; - - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; - } -} - - -__global__ void floyd_warshall(int *d, int blockSize, int length, - int XIndex, int YIndex, int kk) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -} - -__global__ void floyd_warshall_whole(int *d, int blockSize, int length, - int XIndex, int YIndex, int ZIndex) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - //__shared__ int dij[8][8]; - - // dij[threadIdx.x][threadIdx.y] = d[ii * length + jj]; - for (int cur = 0; cur < blockSize; ++cur) { - int kk = ZIndex + cur; - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -/* - if (dij[threadIdx.x][threadIdx.y] > dik + dkj) - dij[threadIdx.x][threadIdx.y] = dik + dkj; -*/ - } - // d[ii * length + jj] = dij[threadIdx.x][threadIdx.y]; -} - -int main(int argc, char **argv) { - - cudaEvent_t total_start, total_stop; - cudaEvent_t com_start, com_stop; - cudaEvent_t mem_start, mem_stop; - cudaEvent_t io_start, io_stop; - - timer(total); - timer(com); - timer(mem); - timer(io); - - cudaEventRecord(total_start); - - float total, compute, memory, IO; - float mem_part, IO_part, com_part; - total = compute = memory = IO = 0; - - FILE *fin = fopen(argv[1], "r"); - FILE *fout = fopen(argv[2], "w"); - int blockSize = atoi(argv[3]); - - int N, M; - int *edge; - int *gpu[2]; - const int deviceNum = 2; - - fscanf(fin, "%d %d", &N, &M); - int gridSize = N % blockSize ? N / blockSize + 1 : N / blockSize; - int length = blockSize * gridSize + 1; - edge = new int[length * length]; - - //fprintf(stderr, "grid = %d, block = %d\n", gridSize, blockSize); - - for (int i = 0; i < length; ++i) { - for (int j = 0; j < length; ++j) - edge[i * length + j] = INF; - edge[i * length + i] = 0; - } - record(io_start); - while (M--) { - int a, b, w; - fscanf(fin, "%d %d %d", &a, &b, &w); - a = a - 1; - b = b - 1; - edge[a * length + b] = w; - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - - for (int i = 0; i < deviceNum; ++i) { - cudaSetDevice(i); - cudaMalloc((void**)&gpu[i], sizeof(int) * length * length); - - if (i == 0) record(mem_start); - cudaMemcpy(gpu[i], edge, sizeof(int) * length * length, H2D); - cudaCheckErrors("melloc & copy gpu"); - if (i == 0) { - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - } - } - // Now only hangle N = 3200 testcase - size_t sharedSize = 8 * 8; - int blockNum = (N + blockSize - 1) / blockSize; - int blockDimension = 8; - gridSize = blockSize / blockDimension; - int gridFactor = 1024 / blockSize ; - gridFactor *= gridFactor; - - - dim3 blocks(gridSize, gridSize); - dim3 threads(blockDimension, blockDimension); - - dim3 blockCol(gridSize * gridFactor, gridSize); - dim3 blockRow(gridSize, gridSize * gridFactor); - int remainBegin = (blockNum / gridFactor) * gridFactor ; - int remain = blockNum - remainBegin; - dim3 blockColRemain(gridSize * remain, gridSize); - dim3 blockRowRemain(gridSize, gridSize * remain); - //wcout << "blocknum = " << blockNum << endl; - - cudaSetDevice(0); - for (int k = 0; k < blockNum; ++k) { - record(com_start); - //wcout << "k = " << k << endl; - // phase one - cudaSetDevice(0); - { - for (int cur = 0; cur < blockSize; ++cur) { - //wcout << "(" << cur << "/" << blockSize << endl; - for (int id = 0; id < 2; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, k, k, k * blockSize + cur); - } - cudaCheckErrors("phase one"); - } - } - // phase two - { - // Column - for (int i = 0; i < blockNum - remain; i = i + gridFactor) { - for (int cur = 0; cur < blockSize; cur++) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, i, k, k * blockSize + cur); - - cudaCheckErrors("phase two column main"); - } - } - } - - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, remainBegin, k, k * blockSize + cur); - cudaCheckErrors("phase two column remain"); - } - } - // Row - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, k, j, k * blockSize + cur); - cudaCheckErrors("phase two row main"); - } - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, k, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase two row remain"); - } - } - } - - //phase three - { - #pragma omp parallel num_threads(2) - { - int thread = omp_get_thread_num(); - int begin, end; - cudaSetDevice(thread); - - if (thread == 0) { - begin = 0; - end = blockNum / 2; - } else { - begin = blockNum / 2; - end = blockNum; - } - for (int i = begin; i < end; i++) { - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < 1; ++cur) { - floyd_warshall_whole<<>> - (gpu[thread], blockSize, length, i, j, k * blockSize + cur); - cudaCheckErrors("phase three row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < 1; cur++) { - floyd_warshall_whole<<>> - (gpu[thread], blockSize, length, i, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase three row remain"); - } - } - } - } - - cudaDeviceSynchronize(); - - record(com_stop); - sync(com_stop); - elapsed(com_part, com_start, com_stop); - compute += com_part; - - int offset = (blockNum / 2) * blockSize * length ; - int copySize = length * length - offset; - record(mem_start); - cudaMemcpy(gpu[1], gpu[0], sizeof(int) * offset, D2D); - cudaMemcpy(gpu[0] + offset, gpu[1] + offset, sizeof(int) * copySize, D2D); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - } - cudaSetDevice(0); - - record(mem_start); - cudaMemcpy(edge, gpu[1], sizeof(int) * length * length, D2H); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - - record(io_start); - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N - 1; ++j) { - if (edge[i * length + j] == INF) - fprintf(fout, "INF "); - else - fprintf(fout, "%d ", edge[i * length + j]); - } - if (edge[i * length + N - 1] == INF) - fprintf(fout, "INF\n"); - else - fprintf(fout, "%d\n", edge[i * length + N - 1]); - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - cudaEventRecord(total_stop); - cudaEventSynchronize(total_stop); - cudaEventElapsedTime(&total, total_start, total_stop); - fprintf(stderr, "\n\n"); - fprintf(stderr, "TOTAL = %f\n", total); - fprintf(stderr, "COMPUTE = %f\n", compute); - fprintf(stderr, "MEMORY = %f\n", memory); - fprintf(stderr, "IO = %f\n", IO); - return 0; -} diff --git a/HW4_102062111/HW4_102062111_mpi.cu b/HW4_102062111/HW4_102062111_mpi.cu deleted file mode 100644 index 232e818..0000000 --- a/HW4_102062111/HW4_102062111_mpi.cu +++ /dev/null @@ -1,312 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include -#include - -#define INF 1e9 -#define H2D cudaMemcpyHostToDevice -#define D2H cudaMemcpyDeviceToHost -#define D2D cudaMemcpyDeviceToDevice -#define COMM MPI_COMM_WORLD - -#define send(buffer, count, dest) \ - MPI_Send(buffer, count, MPI_CHAR, dest, 0, COMM) - -#define recv(buffer, count, src, status)\ - MPI_Recv(buffer, count, MPI_CHAR, src, MPI_ANY_TAG, COMM, &status) - -#define MASK_N 2 -#define MASK_X 5 -#define MASK_Y 5 -#define SCALE 8 -#define ROOT 0 - -#define cudaCheckErrors(msg) \ - do { \ - cudaError_t __err = cudaGetLastError(); \ - if (__err != cudaSuccess) { \ - fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ - msg, cudaGetErrorString(__err), \ - __FILE__, __LINE__); \ - fprintf(stderr, "*** FAILED - ABORTING\n"); \ - exit(1); \ - } \ - } while (0) - -using namespace std; - -void DisplayHeader() -{ - const int kb = 1024; - const int mb = kb * kb; - wcout << "NBody.GPU" << endl << "=========" << endl << endl; - - //wcout << "CUDA version: v" << CUDART_VERSION << endl; - //wcout << "Thrust version: v" << THRUST_MAJOR_VERSION << "." << THRUST_MINOR_VERSION << endl << endl; - - int devCount; - cudaGetDeviceCount(&devCount); - wcout << "CUDA Devices: " << endl << endl; - - for(int i = 0; i < devCount; ++i) - { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; - - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; - } -} - -__global__ void floyd_warshall(int *d, int blockSize, int length, - int XIndex, int YIndex, int kk) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -} - - -int main(int argc, char **argv) { - - int rank, size; - MPI_Init(&argc, &argv); - MPI_Comm_rank(MPI_COMM_WORLD, &rank); - MPI_Comm_size(MPI_COMM_WORLD, &size); - fprintf(stderr, "rank = %d, size = %d\n", rank, size); - - MPI_Status status; - - int blockSize = atoi(argv[3]); - - int N, M; - int *edge; - int *gpu[2]; - int gridSize; - int length; - const int deviceNum = 2; - - if (rank == ROOT) { - FILE *fin = fopen(argv[1], "r"); - - fscanf(fin, "%d %d", &N, &M); - gridSize = N % blockSize ? N / blockSize + 1 : N / blockSize; - length = blockSize * gridSize + 1; - cudaMallocHost((void**)&edge, sizeof(int) * length * length); - for (int i = 0; i < length; ++i) { - for (int j = 0; j < length; ++j) - edge[i * length + j] = INF; - edge[i * length + i] = 0; - } - while (M--) { - int a, b, w; - fscanf(fin, "%d %d %d", &a, &b, &w); - a = a - 1; - b = b - 1; - edge[a * length + b] = w; - } - send(&N, sizeof(int), 1); - send(&gridSize, sizeof(int), 1); - send(&length, sizeof(int), 1); - send(edge, sizeof(int) * length * length, 1); - - fclose(fin); - } else { - recv(&N, sizeof(int), ROOT, status); - recv(&gridSize, sizeof(int), ROOT, status); - recv(&length, sizeof(int), ROOT, status); - cudaMallocHost((void**)&edge, sizeof(int) * length * length); - recv(edge, sizeof(int) * length * length, ROOT, status); - fprintf(stderr, "RANK %d, length = %d, gridSize = %d\n", rank, length, gridSize); - } - - cudaSetDevice(rank); - cudaMalloc((void**)&gpu[rank], sizeof(int) * length * length); - cudaMemcpy(gpu[rank], edge, sizeof(int) * length * length, H2D); - cudaCheckErrors("memcpy error"); - - size_t sharedSize = 8 * 8; - int blockNum = (N + blockSize - 1) / blockSize; - int blockDimension = 8; - gridSize = blockSize / blockDimension; - int gridFactor = 1024 / blockSize ; - gridFactor *= gridFactor; - - dim3 blocks(gridSize, gridSize); - dim3 threads(blockDimension, blockDimension); - - dim3 blockCol(gridSize * gridFactor, gridSize); - dim3 blockRow(gridSize, gridSize * gridFactor); - int remainBegin = (blockNum / gridFactor) * gridFactor ; - int remain = blockNum - remainBegin; - dim3 blockColRemain(gridSize * remain, gridSize); - dim3 blockRowRemain(gridSize, gridSize * remain); - - fprintf(stderr, "rank %d, gpu = %p \n", rank, gpu[rank]); -/* - if (rank == 0) { - send(&gpu[0], sizeof(int*), 1); - } - else { - recv(&gpu[0], sizeof(int*), 0, status); - } - if (rank != 0) { - fprintf(stderr, "gpu[0] = %p, gpu[1] = %p", gpu[0], gpu[1]); - floyd_warshall<<>>(gpu[0], blockSize, length, 0, 0, 0); - cudaCheckErrors("launch kernel"); - } -*/ -/* - cudaEvent_t start, stop; - cudaEventCreate(&start); - cudaEventCreate(&stop); - - cudaEventRecord(start); - int n = 5120 / 256; - for (int i = 0; i < n; ++i) { - if (rank == 0) { - cudaMemcpy(edge, gpu[0], sizeof(int) * length * length, D2H); - send(edge, sizeof(int) * length * length, 1); - recv(edge, sizeof(int) * length * length, 1, status); - cudaMemcpy(gpu[0], edge, sizeof(int) * length * length, H2D); - } else { - recv(edge, sizeof(int) * length * length, 0, status); - cudaMemcpy(gpu[1], edge, sizeof(int) * length * length, H2D); - cudaMemcpy(edge, gpu[1], sizeof(int) * length * length, D2H); - send(edge, sizeof(int) * length * length, 0); - } - cudaCheckErrors("memcpy D2D"); - } - cudaEventRecord(stop); - cudaEventSynchronize(stop); - - float time; - cudaEventElapsedTime(&time, start, stop); - fprintf(stderr, "time = %f\n", time); -*/ - for (int k = 0; k < blockNum; ++k) { - //wcout << "rank = " << rank << "k = " << k << ", blocknum = " << blockNum << endl; - // phase one - { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, k, k, k * blockSize + cur); - cudaCheckErrors("phase one"); - } - } - // phase two - { - // Column - for (int i = 0; i < blockNum - remain; i = i + gridFactor) { - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, i, k, k * blockSize + cur); - cudaCheckErrors("phase two column main"); - } - } - - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, remainBegin, k, k * blockSize + cur); - cudaCheckErrors("phase two column remain"); - } - // Row - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, k, j, k * blockSize + cur); - cudaCheckErrors("phase two row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[rank], blockSize, length, k, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase two row remain"); - } - } - - //phase three - { - int thread = rank; - int begin, end; - cudaSetDevice(thread); - - if (thread == 0) { - begin = 0; - end = blockNum / 2; - } else { - begin = blockNum / 2; - end = blockNum; - } - for (int i = begin; i < end; i++) { - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (gpu[thread], blockSize, length, i, j, k * blockSize + cur); - cudaCheckErrors("phase three row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[thread], blockSize, length, i, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase three row remain"); - } - } - } - int offset = (blockNum / 2) * blockSize * length ; - int copySize = length * length - offset; - if (rank == 0) { - cudaMemcpy(edge, gpu[0], sizeof(int) * offset, D2H); - send(edge, sizeof(int) * offset, 1); - recv(edge + offset, sizeof(int) * copySize, 1, status); - cudaMemcpy(gpu[0] + offset, edge + offset, sizeof(int) * copySize, H2D); - cudaCheckErrors("rank ROOT memcpy D2D"); - } else { - recv(edge, sizeof(int) * offset, ROOT, status); - cudaMemcpy(gpu[1], edge, sizeof(int) * offset, H2D); - cudaMemcpy(edge + offset, gpu[1] + offset, sizeof(int) * copySize, D2H); - cudaCheckErrors("rank 1 memcpy D2D"); - send(edge + offset, sizeof(int) * copySize, ROOT); - } - } - if (rank == 0) { - FILE *fout = fopen(argv[2], "w"); - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N - 1; ++j) { - if (edge[i * length + j] == INF) - fprintf(fout, "INF "); - else - fprintf(fout, "%d ", edge[i * length + j]); - } - if (edge[i * length + N - 1] == INF) - fprintf(fout, "INF\n"); - else - fprintf(fout, "%d\n", edge[i * length + N - 1]); - } - fclose(fout); - } - - MPI_Barrier(MPI_COMM_WORLD); - MPI_Finalize(); - return 0; -} diff --git a/HW4_102062111/HW4_102062111_openmp.cu b/HW4_102062111/HW4_102062111_openmp.cu deleted file mode 100644 index 08a0e68..0000000 --- a/HW4_102062111/HW4_102062111_openmp.cu +++ /dev/null @@ -1,323 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include - -#define INF 1e9 -#define H2D cudaMemcpyHostToDevice -#define D2H cudaMemcpyDeviceToHost -#define D2D cudaMemcpyDeviceToDevice -#define MASK_N 2 -#define MASK_X 5 -#define MASK_Y 5 -#define SCALE 8 - -#define timer(type) \ - cudaEventCreate(&type##_start); \ - cudaEventCreate(&type##_stop); - -#define record(type) \ - cudaEventRecord(type) - -#define elapsed(type, start, stop) \ - cudaEventElapsedTime(&type, start, stop) - -#define sync(type) \ - cudaEventSynchronize(type) - -#define cudaCheckErrors(msg) \ - do { \ - cudaError_t __err = cudaGetLastError(); \ - if (__err != cudaSuccess) { \ - fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ - msg, cudaGetErrorString(__err), \ - __FILE__, __LINE__); \ - fprintf(stderr, "*** FAILED - ABORTING\n"); \ - exit(1); \ - } \ - } while (0) - -using namespace std; - -void DisplayHeader() -{ - const int kb = 1024; - const int mb = kb * kb; - wcout << "NBody.GPU" << endl << "=========" << endl << endl; - - //wcout << "CUDA version: v" << CUDART_VERSION << endl; - //wcout << "Thrust version: v" << THRUST_MAJOR_VERSION << "." << THRUST_MINOR_VERSION << endl << endl; - - int devCount; - cudaGetDeviceCount(&devCount); - wcout << "CUDA Devices: " << endl << endl; - - for(int i = 0; i < devCount; ++i) - { - cudaDeviceProp props; - cudaGetDeviceProperties(&props, i); - wcout << i << ": " << props.name << ": " << props.major << "." << props.minor << endl; - wcout << " Global memory: " << props.totalGlobalMem / mb << "mb" << endl; - wcout << " Shared memory: " << props.sharedMemPerBlock / kb << "kb" << endl; - wcout << " Constant memory: " << props.totalConstMem / kb << "kb" << endl; - wcout << " Block registers: " << props.regsPerBlock << endl << endl; - - wcout << " Warp size: " << props.warpSize << endl; - wcout << " Threads per block: " << props.maxThreadsPerBlock << endl; - wcout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << endl; - wcout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << endl; - wcout << endl; - } -} - - -__global__ void floyd_warshall(int *d, int blockSize, int length, - int XIndex, int YIndex, int kk) { - int ii = blockSize * XIndex + blockIdx.x * blockDim.x + threadIdx.x; - int jj = blockSize * YIndex + blockIdx.y * blockDim.y + threadIdx.y; - - int dij = d[ii * length + jj]; - int dik = d[ii * length + kk]; - int dkj = d[kk * length + jj]; - - if (dij > dik + dkj) - d[ii * length + jj] = dik + dkj; -} - - -int main(int argc, char **argv) { - - cudaEvent_t total_start, total_stop; - cudaEvent_t com_start, com_stop; - cudaEvent_t mem_start, mem_stop; - cudaEvent_t io_start, io_stop; - - timer(total); - timer(com); - timer(mem); - timer(io); - - cudaEventRecord(total_start); - - float total, compute, memory, IO; - float mem_part, IO_part, com_part; - total = compute = memory = IO = 0; - - FILE *fin = fopen(argv[1], "r"); - FILE *fout = fopen(argv[2], "w"); - int blockSize = atoi(argv[3]); - - int N, M; - int *edge; - int *gpu[2]; - const int deviceNum = 2; - - fscanf(fin, "%d %d", &N, &M); - int gridSize = N % blockSize ? N / blockSize + 1 : N / blockSize; - int length = blockSize * gridSize + 1; - edge = new int[length * length]; - - //fprintf(stderr, "grid = %d, block = %d\n", gridSize, blockSize); - - for (int i = 0; i < length; ++i) { - for (int j = 0; j < length; ++j) - edge[i * length + j] = INF; - edge[i * length + i] = 0; - } - record(io_start); - while (M--) { - int a, b, w; - fscanf(fin, "%d %d %d", &a, &b, &w); - a = a - 1; - b = b - 1; - edge[a * length + b] = w; - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - - for (int i = 0; i < deviceNum; ++i) { - cudaSetDevice(i); - cudaMalloc((void**)&gpu[i], sizeof(int) * length * length); - - if (i == 0) record(mem_start); - cudaMemcpy(gpu[i], edge, sizeof(int) * length * length, H2D); - cudaCheckErrors("melloc & copy gpu"); - if (i == 0) { - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - } - } - // Now only hangle N = 3200 testcase - size_t sharedSize = 8 * 8; - int blockNum = (N + blockSize - 1) / blockSize; - int blockDimension = 8; - gridSize = blockSize / blockDimension; - int gridFactor = 1024 / blockSize ; - gridFactor *= gridFactor; - - - dim3 blocks(gridSize, gridSize); - dim3 threads(blockDimension, blockDimension); - - dim3 blockCol(gridSize * gridFactor, gridSize); - dim3 blockRow(gridSize, gridSize * gridFactor); - int remainBegin = (blockNum / gridFactor) * gridFactor ; - int remain = blockNum - remainBegin; - dim3 blockColRemain(gridSize * remain, gridSize); - dim3 blockRowRemain(gridSize, gridSize * remain); - //wcout << "blocknum = " << blockNum << endl; - - cudaSetDevice(0); - for (int k = 0; k < blockNum; ++k) { - record(com_start); - //wcout << "k = " << k << endl; - // phase one - cudaSetDevice(0); - { - for (int cur = 0; cur < blockSize; ++cur) { - //wcout << "(" << cur << "/" << blockSize << endl; - for (int id = 0; id < 2; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, k, k, k * blockSize + cur); - } - cudaCheckErrors("phase one"); - } - } - // phase two - { - // Column - for (int i = 0; i < blockNum - remain; i = i + gridFactor) { - for (int cur = 0; cur < blockSize; cur++) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, i, k, k * blockSize + cur); - - cudaCheckErrors("phase two column main"); - } - } - } - - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, remainBegin, k, k * blockSize + cur); - cudaCheckErrors("phase two column remain"); - } - } - // Row - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, k, j, k * blockSize + cur); - cudaCheckErrors("phase two row main"); - } - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - for (int id = 0; id < deviceNum; ++id) { - cudaSetDevice(id); - floyd_warshall<<>> - (gpu[id], blockSize, length, k, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase two row remain"); - } - } - } - - //phase three - { - #pragma omp parallel num_threads(2) - { - int thread = omp_get_thread_num(); - int begin, end; - cudaSetDevice(thread); - - if (thread == 0) { - begin = 0; - end = blockNum / 2; - } else { - begin = blockNum / 2; - end = blockNum; - } - for (int i = begin; i < end; i++) { - for (int j = 0; j < blockNum - remain; j = j + gridFactor) { - for (int cur = 0; cur < blockSize; ++cur) { - floyd_warshall<<>> - (gpu[thread], blockSize, length, i, j, k * blockSize + cur); - cudaCheckErrors("phase three row main"); - } - } - if (remainBegin < blockNum) - for (int cur = 0; cur < blockSize; cur++) { - floyd_warshall<<>> - (gpu[thread], blockSize, length, i, remainBegin, k * blockSize + cur); - cudaCheckErrors("phase three row remain"); - } - } - } - } - - cudaDeviceSynchronize(); - - record(com_stop); - sync(com_stop); - elapsed(com_part, com_start, com_stop); - compute += com_part; - - int offset = (blockNum / 2) * blockSize * length ; - int copySize = length * length - offset; - cudaMemcpy(gpu[1], gpu[0], sizeof(int) * offset, D2D); - cudaMemcpy(gpu[0] + offset, gpu[1] + offset, sizeof(int) * copySize, D2D); - - } - cudaSetDevice(0); - - record(mem_start); - record(mem_stop); - sync(mem_stop); - elapsed(mem_part, mem_start, mem_stop); - memory += mem_part; - cudaMemcpy(edge, gpu[1], sizeof(int) * length * length, D2H); - - record(io_start); - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N - 1; ++j) { - if (edge[i * length + j] == INF) - fprintf(fout, "INF "); - else - fprintf(fout, "%d ", edge[i * length + j]); - } - if (edge[i * length + N - 1] == INF) - fprintf(fout, "INF\n"); - else - fprintf(fout, "%d\n", edge[i * length + N - 1]); - } - record(io_stop); - sync(io_stop); - elapsed(IO_part, io_start, io_stop); - IO += IO_part; - - cudaEventRecord(total_stop); - cudaEventSynchronize(total_stop); - cudaEventElapsedTime(&total, total_start, total_stop); - fprintf(stderr, "\n\n"); - fprintf(stderr, "TOTAL = %f\n", total); - fprintf(stderr, "COMPUTE = %f\n", compute); - fprintf(stderr, "MEMORY = %f\n", memory); - fprintf(stderr, "IO = %f\n", IO); - return 0; -} diff --git a/HW4_102062111/Makefile b/HW4_102062111/Makefile deleted file mode 100644 index 4f982af..0000000 --- a/HW4_102062111/Makefile +++ /dev/null @@ -1,52 +0,0 @@ -NVFLAGS := -arch=sm_20 -CXXFLAGS := -fopenmp -API = cuda -NAME = HW4_102062111_$(API) -SRC = $(NAME).cu - -MPI_COMPILE_FLAGS = -I/usr/lib/openmpi/include -I/usr/lib/openmpi/include/openmpi -MPI_LINK_FLAGS = -L/usr/lib/openmpi/lib -lmpi_cxx -lmpi -ldl -# MPI_LIBS = $(MPI_COMPILE_FLAGS) $(MPI_LINK_FLAGS) -MPI_LIBS = -I/usr/include/mpich-x86_64 -L/usr/lib64/mpich/lib -lmpich - -all: cuda - -mpi: API = mpi - -fast: API = fast -fast: cuda - -debug: NVFLAGS += -G -g -debug: openmp - -fast_openmp: API = fast_openmp -fast_openmp: cuda - -fast_mpi: API = fast_mpi - -openmp: API = openmp -openmp: cuda - -gpu: API = gpu -gpu: mpi - -sobel: NVFLAGS += -G -g - -test: NAME = test -test: NVFLAGS += -G -g -test: cuda - -cuda: - nvcc -o $(NAME) $(NVFLAGS) -Xcompiler="$(CXXFLAGS)" $(LIBS) $(SRC) - -sobel: sobel.cu - nvcc -o sobel $(NVFLAGS) -Xcompiler="$(CXXFLAGS)" $(LIBS) sobel.cu - -mpi: - nvcc -o $(NAME) $(NVFLAGS) $(MPI_LIBS) $(SRC) - -fast_mpi: - nvcc -o $(NAME) $(NVFLAGS) $(MPI_LIBS) $(SRC) - -clean: - rm -f *cuda *openmp *mpi diff --git a/HW4_102062111/report/.DS_Store b/HW4_102062111/report/.DS_Store deleted file mode 100644 index 6fdf138..0000000 Binary files a/HW4_102062111/report/.DS_Store and /dev/null differ diff --git a/HW4_102062111/report/HW4_102062111_report.md b/HW4_102062111/report/HW4_102062111_report.md deleted file mode 100644 index 468b85d..0000000 --- a/HW4_102062111/report/HW4_102062111_report.md +++ /dev/null @@ -1,227 +0,0 @@ -#
Parallel Programming
- -###
[HW4] 102062111 林致民
- -## Implementation - -1. 首先把資料以 $B \times B$ 作為一個block切割，然後以固定threads個數去分配這些block，如下圖： - - ![img](./data.png) - - 假設黃色的部分之間沒有dependence，這個部分是一次CUDA kernel launch 固定做的範圍，單位時間最多可以同時運算的數量。 - - 再來是Multi-GPU資料切割的方式，我只有在phase 3的時候把資料切成兩半（因為我們只有兩張卡），分別丟到不同的GPU上，做完之後再互相交換彼此的運算結果，下圖是示意圖： ![img](./gpu.png) - - - -2. Multi-GPU 的implement方式則跟single GPU不太一樣，由於在做的時候發現phase one & two 其實影響整個跑分不大，大部分匯集中在phase 3，所以phase one & two 我讓各自的GPU自己算自己的，phase 3才把資料切給不同的GPU做。至於兩個GPU互相溝通的方式，由於使用的API不同，我分成openmp 以及 MPI 不同的方法來說明 - - * OpenMP : 使用`cudaMemcpy(gpu[0], gpu[1], count, cudaDeviceToDevice)`，直接讓device與device之間傳輸資料就好了 - * MPI : 先把`GPU[0]` 上的資料copy到main memory，再透過MPI_Send/Recv 把資料送到另外一個rank，在另外一個rank把資料copy到`GPU[1]`上，`GPU[1]`到`GPU[0]`也是同樣的方法。 - -3. Configuration : - * B = `8 16 32 64 128 256 512 1024` - * block = $1024 \times 1024$ - * threads = $8 \times 8$ - -## Experiment environment - -![img](./datasheet.png) - -```bash -$ nvidia-smi -Sun Jan 10 20:12:22 2016 -+------------------------------------------------------+ -| NVIDIA-SMI 352.63 Driver Version: 352.63 | -|-------------------------------+----------------------+----------------------+ -| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC | -| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. | -|===============================+======================+======================| -| 0 GRID K2 Off | 0000:05:00.0 Off | Off | -| N/A 41C P0 46W / 117W | 11MiB / 4095MiB | 0% Default | -+-------------------------------+----------------------+----------------------+ -| 1 GRID K2 Off | 0000:06:00.0 Off | Off | -| N/A 38C P0 38W / 117W | 11MiB / 4095MiB | 0% Default | -+-------------------------------+----------------------+----------------------+ -``` - -```bash -CPU : Intel(R) Xeon(R) CPU E5-2648L v2 @ 1.90GHz x 2 - 10 cores 20 threads) x 2 = (20 cores 40 threads) -Memory : 128 GB -Storage : 500 GB -Network : Ethernet -Operating System : Ubuntu 12.04 LTS, Linux 3.11.0-26-generic -Compiler : gcc-4.8, CUDA 6.5 -MPI : openmpi-1.5 -``` - -## Profiling - -### Single GPU - -```bash -$ nvprof ./HW4_102062111_cuda testcase/5120.txt cuda.txt 256 -==7356== Profiling result: -Time(%) Time Calls Avg Min Max Name - 99.76% 35.7972s 230400 155.37us 16.161us 446.94us floyd_warshall(int*, int, int, int, int, int) - 0.12% 43.219ms 3 14.406ms 1.0560us 43.216ms [CUDA memcpy HtoD] - 0.12% 41.666ms 1 41.666ms 41.666ms 41.666ms [CUDA memcpy DtoH] - -==7356== API calls: -Time(%) Time Calls Avg Min Max Name - 97.61% 35.1284s 230400 152.47us 8.7540us 5.0724ms cudaLaunch - 0.97% 347.98ms 1382400 251ns 190ns 622.07us cudaSetupArgument - 0.69% 248.64ms 4 62.159ms 9.5980us 204.90ms cudaMemcpy - 0.33% 117.17ms 4 29.292ms 8.9300us 116.97ms cudaMalloc - 0.22% 77.503ms 230406 336ns 277ns 615.26us cudaGetLastError - 0.18% 66.163ms 230400 287ns 253ns 65.080us cudaConfigureCall - 0.00% 830.96us 166 5.0050us 272ns 177.87us cuDeviceGetAttribute - 0.00% 104.15us 2 52.077us 50.666us 53.488us cuDeviceTotalMem - 0.00% 82.443us 2 41.221us 38.413us 44.030us cuDeviceGetName - 0.00% 12.078us 1 12.078us 12.078us 12.078us cudaSetDevice - 0.00% 3.7600us 2 1.8800us 402ns 3.3580us cuDeviceGetCount - 0.00% 1.9750us 4 493ns 346ns 802ns cuDeviceGet -``` - -### Multi-GPU-OpenMP - -```bash -$ nvprof ./HW4_102062111_openmp testcase/5120.txt cuda.txt 256 -==7631== Profiling result: -Time(%) Time Calls Avg Min Max Name - 98.55% 39.4087s 256000 153.94us 16.161us 446.65us floyd_warshall(int*, int, int, int, int, int) - 0.75% 298.38ms 42 7.1042ms 5.1370ms 43.405ms [CUDA memcpy HtoD] - 0.71% 282.57ms 41 6.8920ms 5.1500ms 70.235ms [CUDA memcpy DtoH] - -==7631== API calls: -Time(%) Time Calls Avg Min Max Name - 93.97% 34.2017s 256000 133.60us 8.6040us 4.8841ms cudaLaunch - 3.56% 1.29524s 43 30.122ms 23.544ms 193.92ms cudaMemcpy - 1.17% 426.72ms 1536000 277ns 188ns 611.98us cudaSetupArgument - 0.66% 241.91ms 2 120.96ms 117.15ms 124.76ms cudaMalloc - 0.25% 92.470ms 250882 368ns 264ns 537.29us cudaGetLastError - 0.22% 81.777ms 256000 319ns 240ns 490.32us cudaConfigureCall - 0.15% 54.061ms 51263 1.0540us 425ns 534.07us cudaSetDevice - 0.00% 828.30us 166 4.9890us 273ns 178.27us cuDeviceGetAttribute - 0.00% 102.26us 2 51.129us 50.642us 51.616us cuDeviceTotalMem - 0.00% 82.492us 2 41.246us 38.200us 44.292us cuDeviceGetName - 0.00% 3.8740us 2 1.9370us 411ns 3.4630us cuDeviceGetCount - 0.00% 1.9750us 4 493ns 310ns 804ns cuDeviceGet -``` - -### Multi-GPU-MPI - -```bash -==8525== Profiling result: -==8523== Profiling result: -Time(%) Time Calls Avg Min Max Name - 99.03% 19.7089s 128000 153.98us 16.161us 441.85us floyd_warshall(int*, int, int, int, int, int) - 0.51% 102.03ms 21 4.8585ms 4.5920ms 9.0901ms [CUDA memcpy HtoD] - 0.45% 90.489ms 20 4.5245ms 4.4990ms 4.8963ms [CUDA memcpy DtoH] -Time(%) Time Calls Avg Min Max Name - 99.03% 19.6893s 128000 153.82us 16.161us 444.12us floyd_warshall(int*, int, int, int, int, int) - 0.50% 99.940ms 21 4.7590ms 4.5099ms 9.1663ms [CUDA memcpy HtoD] - 0.46% 92.004ms 20 4.6002ms 4.5923ms 4.6455ms [CUDA memcpy DtoH] - -==8525== API calls: - -==8523== API calls: -Time(%) Time Calls Avg Min Max Name - 94.89% 18.8747s 128000 147.46us 8.4360us 9.2642ms cudaLaunch - 2.28% 453.71ms 41 11.066ms 4.5174ms 161.92ms cudaMemcpy - 1.00% 199.68ms 768000 259ns 192ns 622.54us cudaSetupArgument - 0.82% 162.43ms 1 162.43ms 162.43ms 162.43ms cudaMallocHost - 0.59% 117.09ms 1 117.09ms 117.09ms 117.09ms cudaMalloc - 0.21% 41.552ms 128021 324ns 269ns 27.909us cudaGetLastError - 0.20% 39.931ms 128000 311ns 253ns 521.01us cudaConfigureCall - 0.00% 839.29us 166 5.0550us 278ns 187.02us cuDeviceGetAttribute - 0.00% 587.89us 21 27.994us 1.5030us 550.34us cudaSetDevice - 0.00% 102.02us 2 51.009us 50.628us 51.390us cuDeviceTotalMem - 0.00% 82.815us 2 41.407us 38.461us 44.354us cuDeviceGetName - 0.00% 3.8980us 2 1.9490us 507ns 3.3910us cuDeviceGetCount - 0.00% 1.8960us 4 474ns 278ns 613ns cuDeviceGet -Time(%) Time Calls Avg Min Max Name - 80.41% 19.8362s 128000 154.97us 8.3230us 9.3118ms cudaLaunch - 17.69% 4.36418s 41 106.44ms 4.5969ms 482.63ms cudaMemcpy - 0.88% 216.77ms 768000 282ns 190ns 2.9469ms cudaSetupArgument - 0.66% 163.33ms 1 163.33ms 163.33ms 163.33ms cudaMallocHost - 0.18% 43.634ms 128021 340ns 258ns 18.379us cudaGetLastError - 0.16% 40.688ms 128000 317ns 251ns 17.879us cudaConfigureCall - 0.00% 1.2231ms 1 1.2231ms 1.2231ms 1.2231ms cudaMalloc - 0.00% 842.13us 166 5.0730us 278ns 189.03us cuDeviceGetAttribute - 0.00% 614.16us 21 29.245us 1.5330us 577.63us cudaSetDevice - 0.00% 102.34us 2 51.167us 51.017us 51.318us cuDeviceTotalMem - 0.00% 83.514us 2 41.757us 39.204us 44.310us cuDeviceGetName - 0.00% 4.0110us 2 2.0050us 386ns 3.6250us cuDeviceGetCount - 0.00% 1.8020us 4 450ns 327ns 644ns cuDeviceGet -``` - -上面三組profiling result 有個共同特徵是，cudaLaunch耗用的時間特別多，而其次是memcpy。不過cudaLaunch佔用了接近95%，之後優化的方向可以朝這裡努力。 - -## Experiment & Analysis - -以下的測量方式都是使用cudaEvent，針對特定的block去測量，並且累計總時間。CUDA blocks(128, 128), thread(8, 8) - -### Total runtime -![img](./total.png) - -這個實驗是測量執行程式的總時間，其實可以明顯看到MultiGPU的執行時間遠大於Single GPU，而openmp 的執行時間又小於MPI，由於MPI還需要communication，把算好的資料搬到另外一個process上，所以比openmp版本的時間多上一些。由於$Floyd\ Warshall$的時間複雜度是$O(N^3)$，所以可以看到時間呈現凹向上的狀態。 - -### Computating time (kernel runtime) -![img](./compute.png) - -這個實驗最主要是去除其他影響計算的因素，包括malloc, memcpy, IO... 等，單純只計算CUDA kernel執行的時間。以這張圖為例，single GPU的版本基本上比multi-GPU的版本的計算時間還要來的長。然後openmp和MPI的計算時間是差不多的，但是總時間openmp和MPI有些為的落差，下面其他的實驗會探討其原因。 - - -### Memory copy time -![img](./memory.png) - -針對memcpy，***OpenMP***的copy是`DeviceToDevice`， ***MPI*** 的copy方式是先copy到host memory，然後傳到其他process後，在copy進另外一個GPU的memory。在這個實驗看到有趣的地方是，使用`DeviceToDevice` copy 的時間比MPI copy的時間還來的長，Device之間傳輸的速度似乎沒有很快。當然Single GPU只有read & write有用到memory copy，所以當然比兩個GPU cpoy的時間還來的短。 - -### IO time -![img](./IO.png) - -IO的話三者是差不多的，畢竟只有一開始讀檔，還有把結果寫入檔案，三者的實作方式是一樣的。 - -### MPI Communication time -![img](./communicate.png) - -此實驗針對不同Input Size(N)，測量兩個process之間的communication，雖然理論上趨勢應該是會越來越多，但是在 N = 7168 執行時間就掉下去，在優化的版本也有類似的狀況發生，所以我猜測是因為每次傳輸資料的分配不平均，但是在某些狀況下，資料分配就很平均，所以傳輸的時間會快一點。 - -### Blocking Factor - -#### GFLOPS -Floyd Warshall 的時間複雜度是$O(N^3)$，因此假設我們最少需要計算$N^3$次，然後我們假設下面這段code 需要$3\ flops$，而且$N=1024$： - -```cpp -if (d[i][j] > d[i][k] + d[k][j]) - d[i][j] = d[i][k] + d[k][j]; -``` -所我以們需要的計算次數是：$GFLOPS=3 \times N^3 \div computing\ time \div 10^9$ - -以下是實驗結果： -![img](./gflops.png) - -在 blocksize = 32 到 64，GFLOPS會上生，$blockSize = 64$的時候會發現到，他的GFLOPS是最大的，在之後才會繼續往下掉。 - -#### BandWidth - -對於MPI以及OpenMP來說，Memcopy在每一個回合都需要和另外一個device溝通，每一次copy的資料量都是 N * N，然後需要做$N \div blockfactor$次，因此記方法如下： -$Bandwidth(GB/s) = N \div blockfactor * N * N \div memory\ copy\ time$ - -以下是實驗結果： - -![img](./bandwidth.png) - -由於我memory copy 的方式是每一回合直接對整份資料切成兩半，phase 3 做完在教換彼此沒有的資料。因此當blockFactor越小的時候，呼叫$cudaMemcpy$的次數就會變的比較多，相反的blockSize較大的時候，呼叫的次數就會比較少。OpenMP採用的是DeviceTODevice的memory Copy，不過跑出來的結果卻比MPI的DeviceToHost->HostToDevice還小。當然最高的memory bandwidth 可以接近2.5GB/s。不過我認為可能是因為資料量小，才會有較大的誤差。 - -## Optimization - -做完上述的版本之後，其實對於跑出來的效能不是很滿意，因此嘗試去做了優化，不過大多數都是在嘗試cuda block/thread之間參數的問題，這部分就不多加描述了。 - -前面我的profiling結果顯示，我的cudaLaunch佔了一大部分的時間，所以我嘗試優化這部分。我做完所有版本之後，回頭觀察一下演算法，發現phase 3 基本上是利用phase 2 以及phase 1的最佳解來迭代自己的答案，所以基本上不用去handle cuda block之間不能synchronize的問題，因此不用特別為了k值去synchronize所有的threads(包括不同的block)，也可以減少launch kernel的次數。優化後的版本如下圖： - -![img](./optimize.png) - -減少kernel launch 的時間就讓整個時間降了下來，甚至還比沒有優化過的openmp還快。其實也間接證明我的program的bottle neck是在kernel launch。OpenMP在極大的測資下，也可以在50秒左右跑完。 \ No newline at end of file diff --git a/HW4_102062111/report/HW4_102062111_report.pdf b/HW4_102062111/report/HW4_102062111_report.pdf deleted file mode 100644 index fd0bf21..0000000 Binary files a/HW4_102062111/report/HW4_102062111_report.pdf and /dev/null differ diff --git a/HW4_102062111/report/IO.png b/HW4_102062111/report/IO.png deleted file mode 100644 index cc78da0..0000000 Binary files a/HW4_102062111/report/IO.png and /dev/null differ diff --git a/HW4_102062111/report/bandwidth.png b/HW4_102062111/report/bandwidth.png deleted file mode 100644 index 063f289..0000000 Binary files a/HW4_102062111/report/bandwidth.png and /dev/null differ diff --git a/HW4_102062111/report/communicate.png b/HW4_102062111/report/communicate.png deleted file mode 100644 index 983dcd7..0000000 Binary files a/HW4_102062111/report/communicate.png and /dev/null differ diff --git a/HW4_102062111/report/compute.png b/HW4_102062111/report/compute.png deleted file mode 100644 index e245cf5..0000000 Binary files a/HW4_102062111/report/compute.png and /dev/null differ diff --git a/HW4_102062111/report/data.png b/HW4_102062111/report/data.png deleted file mode 100644 index e1fc857..0000000 Binary files a/HW4_102062111/report/data.png and /dev/null differ diff --git a/HW4_102062111/report/datasheet.png b/HW4_102062111/report/datasheet.png deleted file mode 100644 index 62146bc..0000000 Binary files a/HW4_102062111/report/datasheet.png and /dev/null differ diff --git a/HW4_102062111/report/gflops.png b/HW4_102062111/report/gflops.png deleted file mode 100644 index c022574..0000000 Binary files a/HW4_102062111/report/gflops.png and /dev/null differ diff --git a/HW4_102062111/report/gpu.png b/HW4_102062111/report/gpu.png deleted file mode 100644 index a2762df..0000000 Binary files a/HW4_102062111/report/gpu.png and /dev/null differ diff --git a/HW4_102062111/report/memory.png b/HW4_102062111/report/memory.png deleted file mode 100644 index 1cc47de..0000000 Binary files a/HW4_102062111/report/memory.png and /dev/null differ diff --git a/HW4_102062111/report/optimize.png b/HW4_102062111/report/optimize.png deleted file mode 100644 index 8f03c0a..0000000 Binary files a/HW4_102062111/report/optimize.png and /dev/null differ diff --git a/HW4_102062111/report/ploter.m b/HW4_102062111/report/ploter.m deleted file mode 100644 index a021d5e..0000000 --- a/HW4_102062111/report/ploter.m +++ /dev/null @@ -1,165 +0,0 @@ -cuda = load('report/HW4_102062111_cuda.log'); -openmp = load('report/HW4_102062111_openmp.log'); -mpi = load('report/HW4_102062111_mpi.log'); - -total(1,:) = cuda(:, 1) / 1000; -total(2,:) = openmp(:, 1) / 1000; -total(3,:) = mpi(:, 1) / 1000; - -plot(total'); -title('total time'); -legend('cuda', 'openmp', 'mpi'); -xlabel('testcase Size(N)', 'FontSize', 16); -ylabel('run time(seconds)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',[1024:1024:9216],... - 'XTick', [1:9]); - -%% -figure; -compute(1,:) = cuda(:, 2) / 1000; -compute(2,:) = openmp(:, 2) / 1000; -compute(3,:) = mpi(:, 2) / 1000; - -plot(compute'); -title('computing time'); -legend('cuda', 'openmp', 'mpi'); -xlabel('testcase Size(N)', 'FontSize', 16); -ylabel('run time(seconds)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',[1024:1024:9216],... - 'XTick', [1:9]); - - %% - - figure; -io(1,:) = cuda(:, 3) / 1000; -io(2,:) = openmp(:, 3) / 1000; -io(3,:) = mpi(:, 4) / 1000; - -plot(io'); -title('memory time'); -legend('cuda', 'openmp', 'mpi'); -xlabel('testcase Size(N)', 'FontSize', 16); -ylabel('run time(seconds)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',[1024:1024:9216],... - 'XTick', [1:9]); - - %% -figure; -io(1,:) = cuda(:, 4) / 1000; -io(2,:) = openmp(:, 4) / 1000; -io(3,:) = mpi(:, 5) / 1000; - -plot(io'); -title('IO time'); -legend('cuda', 'openmp', 'mpi'); -xlabel('testcase Size(N)', 'FontSize', 16); -ylabel('run time(seconds)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',[1024:1024:9216],... - 'XTick', [1:9]); - - %% Communication time - figure; - -comm(1,:) = mpi(:, 3) / 1000; - -plot(comm); -title('Communication time'); -legend('mpi'); -xlabel('testcase Size(N)', 'FontSize', 16); -ylabel('run time(seconds)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',[1024:1024:9216],... - 'XTick', [1:9]); - %% - -cuda_fast = load('report/HW4_102062111_fast_cuda.log'); -openmp_fast = load('report/HW4_102062111_fast_openmp.log'); -mpi_fast = load('report/HW4_102062111_fast_mpi.log'); - -total(1,:) = cuda_fast(:, 2) / 1000; -total(2,:) = openmp_fast(:, 2) / 1000; -total(3,:) = cuda(:, 2) / 1000; -total(4,:) = openmp(:, 2) / 1000; - -plot(total'); -title('computing time'); -legend('cuda\_optimize', 'openmp\_optimize', 'cuda', 'openmp'); -xlabel('testcase Size(N)', 'FontSize', 16); -ylabel('run time(seconds)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',[1024:1024:9216],... - 'XTick', [1:9]); -%% - -cuda_kernel = load('report/HW4_102062111_cuda_block_kernel.log'); -openmp_kernel = load('report/HW4_102062111_openmp_block_kernel.log'); -mpi_kernel = load('report/HW4_102062111_mpi_block_kernel.log'); - - -N = 1024 -cuda_gflops = []; -openmp_gflops = []; -mpi_gflops = []; -for i = 1:6, - blockSize = 2 ^ (i + 4); - blockNum = N / blockSize; - cuda_gflops(i) = 3 * N ^ 3 / 10^9 / (cuda_kernel(i,2) / 10^3); - openmp_gflops(i) = 3 * N ^ 3 / 10^9 / (openmp_kernel(i,2) / 10^3); - mpi_gflops(i) = 3 * N ^ 3 / 10^9 / (mpi_kernel(i,2) / 10^3); -end - -total = [cuda_gflops; openmp_gflops; mpi_gflops]'; - -label = cuda_block(:,1); -bar(total); -title('GFLOPS'); -legend('cuda', 'openmp', 'mpi'); -xlabel('Block Factor', 'FontSize', 16); -ylabel('GFLOPS', 'FontSize', 16); - -set(gca,... - 'XTickLabel',label,... - 'XTick', 1:6); - - %% - -cuda_kernel = load('report/HW4_102062111_cuda_block_memory.log'); -openmp_kernel = load('report/HW4_102062111_openmp_block_memory.log'); -mpi_kernel = load('report/HW4_102062111_mpi_block_memory.log'); - - -N = 1024 -cuda_gflops = []; -openmp_gflops = []; -mpi_gflops = []; -for i = 1:6, - blockSize = 2 ^ (i + 4); - blockNum = N / blockSize; - - % cuda_gflops(i) = blockNum * N ^ 2 / (cuda_kernel(i,2) / 10^3) / 10^9; - openmp_gflops(i) = blockNum * N ^ 2 / (openmp_kernel(i,2) / 10^3) / 10^9; - mpi_gflops(i) = blockNum * N ^ 2 / (mpi_kernel(i,2) / 10^3) / 10^9; -end - -total = [openmp_gflops; mpi_gflops]'; - -label = cuda_block(:,1); -plot(total); -title('Memory BandWidth'); -legend('openmp', 'mpi'); -xlabel('Block Factor', 'FontSize', 16); -ylabel('BandWidth(GB/s)', 'FontSize', 16); - -set(gca,... - 'XTickLabel',label,... - 'XTick', 1:6); \ No newline at end of file diff --git a/HW4_102062111/report/report/HW4_102062111_cuda.log b/HW4_102062111/report/report/HW4_102062111_cuda.log deleted file mode 100644 index 90a6544..0000000 --- a/HW4_102062111/report/report/HW4_102062111_cuda.log +++ /dev/null @@ -1,9 +0,0 @@ -925 419 4 491 -3537 2662 12 832 -9774 8251 32 1424 -20859 18471 47 2223 -39617 35878 83 3471 -67097 61885 119 4825 -101621 94897 150 6214 -147165 138546 271 7873 -217982 206733 493 10154 diff --git a/HW4_102062111/report/report/HW4_102062111_cuda_block.log b/HW4_102062111/report/report/HW4_102062111_cuda_block.log deleted file mode 100644 index b134c33..0000000 --- a/HW4_102062111/report/report/HW4_102062111_cuda_block.log +++ /dev/null @@ -1,6 +0,0 @@ -32 477915 -64 934534 -128 1809687 -256 3510857 -512 7077888 -1024 15728640 diff --git a/HW4_102062111/report/report/HW4_102062111_cuda_block_kernel.log b/HW4_102062111/report/report/HW4_102062111_cuda_block_kernel.log deleted file mode 100644 index e659597..0000000 --- a/HW4_102062111/report/report/HW4_102062111_cuda_block_kernel.log +++ /dev/null @@ -1,6 +0,0 @@ -32 403.853485 -64 344.607452 -128 345.932678 -256 418.571594 -512 583.395020 -1024 1011.617188 diff --git a/HW4_102062111/report/report/HW4_102062111_cuda_block_memory.log b/HW4_102062111/report/report/HW4_102062111_cuda_block_memory.log deleted file mode 100644 index 7845d5f..0000000 --- a/HW4_102062111/report/report/HW4_102062111_cuda_block_memory.log +++ /dev/null @@ -1,6 +0,0 @@ -32 6.627552 -64 6.714336 -128 3.686560 -256 3.466816 -512 3.414080 -1024 4.675904 diff --git a/HW4_102062111/report/report/HW4_102062111_fast_cuda.log b/HW4_102062111/report/report/HW4_102062111_fast_cuda.log deleted file mode 100644 index 82c8ae5..0000000 --- a/HW4_102062111/report/report/HW4_102062111_fast_cuda.log +++ /dev/null @@ -1,9 +0,0 @@ -774 264 4 497 -2495 1471 17 975 -5920 4315 37 1501 -11938 9376 68 2377 -21169 17412 76 3496 -34721 29237 155 5062 -52226 45121 152 6591 -74416 65612 212 8115 -104088 92888 253 10343 diff --git a/HW4_102062111/report/report/HW4_102062111_fast_mpi.log b/HW4_102062111/report/report/HW4_102062111_fast_mpi.log deleted file mode 100644 index 60506b7..0000000 --- a/HW4_102062111/report/report/HW4_102062111_fast_mpi.log +++ /dev/null @@ -1,9 +0,0 @@ -760 208 5 2 533 -2266 1029 41 15 1139 -4488 2829 135 47 1387 -8581 5843 309 108 2166 -14941 10518 590 203 3392 -24466 17297 1476 346 4812 -35101 26201 1637 539 6261 -51975 37420 4866 790 8070 -69149 52690 4203 1114 10393 diff --git a/HW4_102062111/report/report/HW4_102062111_fast_openmp.log b/HW4_102062111/report/report/HW4_102062111_fast_openmp.log deleted file mode 100644 index 67a6490..0000000 --- a/HW4_102062111/report/report/HW4_102062111_fast_openmp.log +++ /dev/null @@ -1,9 +0,0 @@ -836 209 7 485 -2125 1029 34 950 -4457 2828 87 1386 -8439 5846 204 2167 -14497 10500 348 3341 -22939 17300 646 4582 -34561 26198 961 6868 -47528 37418 1314 8123 -65165 52737 1518 10092 diff --git a/HW4_102062111/report/report/HW4_102062111_mpi.log b/HW4_102062111/report/report/HW4_102062111_mpi.log deleted file mode 100644 index 6b08b0a..0000000 --- a/HW4_102062111/report/report/HW4_102062111_mpi.log +++ /dev/null @@ -1,9 +0,0 @@ -843 287 5 2 537 -2517 1615 39 15 809 -6542 4810 141 47 1449 -13281 10425 330 108 2259 -24044 19624 738 205 3226 -40460 33554 1098 348 4907 -65052 51114 6752 541 6154 -86212 73891 2559 796 7731 -125866 109360 3246 1098 10135 diff --git a/HW4_102062111/report/report/HW4_102062111_mpi_block_kernel.log b/HW4_102062111/report/report/HW4_102062111_mpi_block_kernel.log deleted file mode 100644 index f8dc130..0000000 --- a/HW4_102062111/report/report/HW4_102062111_mpi_block_kernel.log +++ /dev/null @@ -1,6 +0,0 @@ -32 216.768158 -64 195.226990 -128 212.164673 -256 285.164307 -512 451.332886 -1024 758.123596 diff --git a/HW4_102062111/report/report/HW4_102062111_mpi_block_memory.log b/HW4_102062111/report/report/HW4_102062111_mpi_block_memory.log deleted file mode 100644 index dd3086a..0000000 --- a/HW4_102062111/report/report/HW4_102062111_mpi_block_memory.log +++ /dev/null @@ -1,6 +0,0 @@ -32 14.097184 -64 7.309344 -128 3.909856 -256 2.295200 -512 1.374656 -1024 1.021984 diff --git a/HW4_102062111/report/report/HW4_102062111_openmp.log b/HW4_102062111/report/report/HW4_102062111_openmp.log deleted file mode 100644 index 9d152b0..0000000 --- a/HW4_102062111/report/report/HW4_102062111_openmp.log +++ /dev/null @@ -1,9 +0,0 @@ -915 286 5 485 -2573 1615 29 809 -6410 4791 83 1383 -13052 10396 216 2216 -23487 19624 326 3234 -39068 33461 570 4635 -58843 51075 868 6368 -83893 73953 1351 7914 -121858 109257 2156 9619 diff --git a/HW4_102062111/report/report/HW4_102062111_openmp_block.log b/HW4_102062111/report/report/HW4_102062111_openmp_block.log deleted file mode 100644 index 4e43b47..0000000 --- a/HW4_102062111/report/report/HW4_102062111_openmp_block.log +++ /dev/null @@ -1,6 +0,0 @@ -32 932916 -64 1798516 -128 3440640 -256 6684672 -512 13762560 -1024 31457280 diff --git a/HW4_102062111/report/report/HW4_102062111_openmp_block_kernel.log b/HW4_102062111/report/report/HW4_102062111_openmp_block_kernel.log deleted file mode 100644 index b0ef16a..0000000 --- a/HW4_102062111/report/report/HW4_102062111_openmp_block_kernel.log +++ /dev/null @@ -1,6 +0,0 @@ -32 224.550034 -64 201.295090 -128 214.845184 -256 285.044647 -512 452.308838 -1024 758.169312 diff --git a/HW4_102062111/report/report/HW4_102062111_openmp_block_memory.log b/HW4_102062111/report/report/HW4_102062111_openmp_block_memory.log deleted file mode 100644 index a86e1ed..0000000 --- a/HW4_102062111/report/report/HW4_102062111_openmp_block_memory.log +++ /dev/null @@ -1,6 +0,0 @@ -32 24.517439 -64 14.753088 -128 9.391359 -256 5.909376 -512 4.830432 -1024 256.408905 diff --git a/HW4_102062111/report/report/blockfactor.rb b/HW4_102062111/report/report/blockfactor.rb deleted file mode 100755 index 69e1c3f..0000000 --- a/HW4_102062111/report/report/blockfactor.rb +++ /dev/null @@ -1,25 +0,0 @@ -#!/usr/bin/env ruby - -prefix = 'HW4_102062111' -apis = ['cuda', 'openmp'] - -apis.each do |api| - name = prefix + "_#{api}" - target = name + "_block.log" - f = File.open(target, 'w') - for i in 5..10 - block = 2 ** i - command = "nvprof --metrics inst_integer ./#{name} testcase/1024.txt ans.txt #{block} 2>&1 | grep -i inst_integer" - puts command - ret = `#{command}` - - total_time = 0 - ret.split("\n").each do |line| - time = line.split(" ")[-1].to_i - total_time = total_time + time - end - puts "#{name} #{block} #{total_time}" - f.write("#{block} #{total_time}\n") - end - f.close -end diff --git a/HW4_102062111/report/report/gflops.png b/HW4_102062111/report/report/gflops.png deleted file mode 100644 index 13d9452..0000000 Binary files a/HW4_102062111/report/report/gflops.png and /dev/null differ diff --git a/HW4_102062111/report/total.png b/HW4_102062111/report/total.png deleted file mode 100644 index e5089c8..0000000 Binary files a/HW4_102062111/report/total.png and /dev/null differ diff --git a/HW4_102062111/testcase/floyd.cpp b/HW4_102062111/testcase/floyd.cpp deleted file mode 100644 index 9d7c4f4..0000000 --- a/HW4_102062111/testcase/floyd.cpp +++ /dev/null @@ -1,52 +0,0 @@ -#include -#include -#include -#include -#include -#define INF 1e9 - -int d[10001][10001]; - -int main(int argc, const char *argv[]){ - - FILE *f = fopen("testcase.txt", "r"); - FILE *fout = fopen("output.txt", "w"); - - int N, M; - - fscanf(f, "%d %d", &N, &M); - for (int i = 1; i <=N; ++i) { - for (int j = 1; j <= N; ++j) - d[i][j] = INF; - d[i][i] = 0; - } - - for (int i = 0; i < M; i++) { - int x , y, w; - fscanf(f, "%d %d %d", &x, &y, &w); - d[x][y] = w; - } - for (int k = 1; k <= N; ++k) { - #pragma omp parallel num_threads(40) private(i) - { - int i; - #pragma omp for schedule(static) - for (i = 1; i <= N; ++i) { - //printf("thread %d, i = %d\n", omp_get_thread_num(), i ); - for (int j = 1; j <= N; ++j) - if (d[i][j] > d[i][k] + d[k][j]) - d[i][j] = d[i][k] + d[k][j]; - } - } - } - - for (int i = 1; i <= N; i++) { - for (int j = 1; j < N ; ++j) - fprintf(fout, "%d ", d[i][j]); - fprintf(fout, "%d\n", d[i][N]); - } - - fclose(fout); - fclose(f); - return 0; -} diff --git a/HW4_102062111/testcase/testcase.rb b/HW4_102062111/testcase/testcase.rb deleted file mode 100755 index 62d51aa..0000000 --- a/HW4_102062111/testcase/testcase.rb +++ /dev/null @@ -1,23 +0,0 @@ -#!/usr/bin/env ruby - -if ARGV.size < 2 - puts "Please enter n/m " - exit 1 -end - -f = File.open('testcase.txt', 'w') -N = ARGV[0].to_i -M = ARGV[1].to_i - -f.write("#{N} #{M}\n") -for i in 1..M - x = 2 - y = 2 - w = (rand() * 100).to_i + 1 - while x == y - x = (rand() * N).to_i + 1 - y = (rand() * N).to_i + 1 - end - f.write("#{x} #{y} #{w}\n") -end -f.close diff --git a/README.md b/README.md new file mode 100644 index 0000000..3bf1f93 --- /dev/null +++ b/README.md @@ -0,0 +1,10 @@ +# Introduction + +This repo is "NTHU Parallel Programming" course project. + +There are four topics: + +1. Use MPI to do "Odd & Even Sort" +2. Use pthread & OpenMP to do "NBody" +3. Use OpenMP & MPI & Hybrid to do "mandelbrot set" +4. Use CUDA to do "All Pairs Shortest Path"