diff --git a/src/bincount-omp/CMakeLists.txt b/src/bincount-omp/CMakeLists.txt index ef7314456..b0031369b 100644 --- a/src/bincount-omp/CMakeLists.txt +++ b/src/bincount-omp/CMakeLists.txt @@ -6,3 +6,15 @@ add_hecbench_benchmark( SOURCES main.cpp CATEGORIES algorithms ) + +# Exact histogram validation requires the host reference and Intel GPU kernel +# to use the same floating-point division semantics. +if(CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM") + target_compile_options(bincount-omp PRIVATE -fp-model=precise) + target_compile_definitions(bincount-omp PRIVATE + __STRICT_ANSI__ + ) + target_link_options(bincount-omp PRIVATE + "SHELL:-Xopenmp-target-backend -cl-fp32-correctly-rounded-divide-sqrt" + ) +endif() diff --git a/src/bincount-omp/Makefile b/src/bincount-omp/Makefile index a8d7eb0b3..906512925 100644 --- a/src/bincount-omp/Makefile +++ b/src/bincount-omp/Makefile @@ -43,7 +43,7 @@ ifeq ($(OPTIMIZE),yes) endif ifeq ($(DEVICE),gpu) - CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ -DHAVE_OMPX_DEVICE_INFO + CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ LDFLAGS +=-Xopenmp-target-backend "-cl-fp32-correctly-rounded-divide-sqrt" else CFLAGS +=-qopenmp diff --git a/src/bincount-omp/main.cpp b/src/bincount-omp/main.cpp index 9eeb6c72a..c4e74bcc8 100644 --- a/src/bincount-omp/main.cpp +++ b/src/bincount-omp/main.cpp @@ -4,32 +4,22 @@ #include #include #include +#include #include #include "reference.h" -#define threadsPerBlock 256 - -// Capacity of the per-team histogram. OpenMP has no way to size a team-local -// array at run time, so this is fixed when the kernel is built; the usable -// limit is the smaller of this and what the device actually offers. -// It is kept below a device's full 64 KB local memory so that the compiler's -// own per-kernel local-memory overhead still fits under the hardware limit. -#define sharedMemoryCapacity (63 * 1024) - -// Largest per-team local memory the device provides. OpenMP has no portable -// query for this; ompx_get_device_info is an Intel extension that also needs -// the offload runtime, so the Makefile only defines HAVE_OMPX_DEVICE_INFO for -// an offload build. Everything else falls back to the built-in capacity. -static int getDeviceLocalMemSize() -{ -#ifdef HAVE_OMPX_DEVICE_INFO - size_t localMemSize = 0, sizeRet = 0; - if (ompx_get_device_info(omp_get_default_device(), ompx_devinfo_local_mem_size, - sizeof(localMemSize), &localMemSize, &sizeRet) == 0) - return (int) localMemSize; +#ifndef THREADS_PER_BLOCK +#define THREADS_PER_BLOCK 256 #endif - return sharedMemoryCapacity; -} + +// Bound the auto-tuner's temporary storage and search space independently of +// the compiler or target architecture. +static constexpr size_t PARTIAL_HISTOGRAM_MEMORY_BUDGET = + 64u * 1024u * 1024u; +static constexpr int MAX_PARTIAL_HISTOGRAMS = 1024; +static constexpr int MAX_LOCAL_HISTOGRAM_TEAMS = 1024; +static constexpr int MAX_LOCAL_HISTOGRAM_BINS = 4096; +static constexpr int CALIBRATION_ITERATIONS = 3; #pragma omp declare target template @@ -48,7 +38,9 @@ getBin(input_t v, input_t minvalue, input_t maxvalue, IndexType nbins) /* Calculate the frequency of the input values. - The GPU offloaded kernel atomically updates the global histogram tensor. + Runtime calibration selects direct atomics, global partial histograms, or + an OpenMP 4.5 team-scoped histogram. Selection depends on actual compiler + and device behavior rather than compiler-name heuristics. */ template void eval(IndexType input_size, int repeat) @@ -71,16 +63,11 @@ void eval(IndexType input_size, int repeat) input_t input_maxvalue = *max_iter; printf("Input min, max values: (%f %f)\n", (float)input_minvalue, (float)input_maxvalue); - #pragma omp target enter data map(to: input[0:input_size]) - - const int maxSharedMemory = - std::min(getDeviceLocalMemSize(), (int)sharedMemoryCapacity); - printf("Maximum shared local memory size per block in bytes: %d\n", maxSharedMemory); - + #pragma omp target data map(to: input[0:input_size]) + { for (IndexType nbins = 768; nbins <= 768 * 32; nbins = nbins * 2) { printf("\nNumber of bins: %d\n", nbins); - IndexType sharedMem = nbins * sizeof(output_t); IndexType output_size = nbins; size_t output_size_bytes = sizeof(output_t) * output_size; @@ -92,112 +79,274 @@ void eval(IndexType input_size, int repeat) output_r, input, nbins, input_minvalue, input_maxvalue, input_size, output_size, repeat); - #pragma omp target enter data map(alloc: output[0:output_size]) - input_t minvalue = input_minvalue; input_t maxvalue = input_maxvalue; - // determine memory type to use in the kernel - printf("bincount using global atomics\n"); - - #pragma omp target teams distribute parallel for - for (IndexType i = 0; i < output_size; i++) output[i] = 0; - - auto start = std::chrono::steady_clock::now(); - for (int n = 0; n < repeat; n++) { + const IndexType inputTeams = + (input_size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + const size_t histogramBytes = (size_t)nbins * sizeof(output_t); + const IndexType memoryLimitedHistograms = + (IndexType)(PARTIAL_HISTOGRAM_MEMORY_BUDGET / histogramBytes); + const IndexType maxPartialHistograms = + std::min( + inputTeams, + std::min(MAX_PARTIAL_HISTOGRAMS, + memoryLimitedHistograms)); + const size_t partialSize = + (size_t)maxPartialHistograms * (size_t)nbins; + output_t *partial = + (output_t*) malloc(partialSize * sizeof(output_t)); + + std::vector candidates; + candidates.push_back(0); + if (maxPartialHistograms > 1) { + for (IndexType count = 32; count < maxPartialHistograms; count *= 4) + candidates.push_back(count); + if (candidates.back() != maxPartialHistograms) + candidates.push_back(maxPartialHistograms); + } + if (nbins <= MAX_LOCAL_HISTOGRAM_BINS) { + const IndexType maxLocalTeams = + std::min(inputTeams, MAX_LOCAL_HISTOGRAM_TEAMS); + for (IndexType count = 128; count < maxLocalTeams; count *= 4) + candidates.push_back(-count); + if (maxLocalTeams > 0 && + (candidates.empty() || candidates.back() != -maxLocalTeams)) + candidates.push_back(-maxLocalTeams); + } + #pragma omp target data map(from: output[0:output_size]) \ + map(alloc: partial[0:partialSize]) + { #pragma omp target teams distribute parallel for - for (IndexType linearIndex = 0; linearIndex < input_size; linearIndex++) { - const input_t v = input[linearIndex]; - if (v >= minvalue && v <= maxvalue) { - const IndexType bin = getBin( - v, minvalue, maxvalue, nbins); - #pragma omp atomic update - output[bin] += 1; + for (IndexType i = 0; i < output_size; i++) output[i] = 0; + #pragma omp target teams distribute parallel for + for (size_t i = 0; i < partialSize; i++) partial[i] = 0; + + long long bestTime = 0; + IndexType bestConfiguration = 0; + for (const IndexType activeConfiguration : candidates) { + // The first trial warms this exact target region. Only the second trial + // participates in selection, excluding JIT and runtime initialization. + for (int trial = 0; trial < 2; trial++) { + #pragma omp target teams distribute parallel for + for (IndexType i = 0; i < output_size; i++) output[i] = 0; + + const auto calibrationStart = std::chrono::steady_clock::now(); + if (activeConfiguration == 0) { + for (int iteration = 0; + iteration < CALIBRATION_ITERATIONS; iteration++) { + #pragma omp target teams distribute parallel for + for (IndexType linearIndex = 0; + linearIndex < input_size; linearIndex++) { + const input_t v = input[linearIndex]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + #pragma omp atomic update + output[bin] += 1; + } + } + } + } else if (activeConfiguration < 0) { + const IndexType localTeams = -activeConfiguration; + for (int iteration = 0; + iteration < CALIBRATION_ITERATIONS; iteration++) { + #pragma omp target teams num_teams(localTeams) \ + thread_limit(THREADS_PER_BLOCK) + { + output_t histogram[MAX_LOCAL_HISTOGRAM_BINS]; + + #pragma omp parallel shared(histogram) + { + const int thread = omp_get_thread_num(); + const int threads = omp_get_num_threads(); + const int team = omp_get_team_num(); + const int teams = omp_get_num_teams(); + + for (IndexType bin = thread; bin < nbins; bin += threads) + histogram[bin] = 0; + + #pragma omp barrier + + for (IndexType i = (IndexType)team * threads + thread; + i < input_size; i += (IndexType)teams * threads) { + const input_t v = input[i]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + #pragma omp atomic update + histogram[bin] += 1; + } + } + + #pragma omp barrier + + for (IndexType bin = thread; bin < nbins; bin += threads) { + #pragma omp atomic update + output[bin] += histogram[bin]; + } + } + } + } + } else { + const IndexType activePartialHistograms = activeConfiguration; + for (int iteration = 0; + iteration < CALIBRATION_ITERATIONS; iteration++) { + #pragma omp target teams distribute parallel for \ + num_teams(activePartialHistograms) \ + thread_limit(THREADS_PER_BLOCK) + for (IndexType linearIndex = 0; + linearIndex < input_size; linearIndex++) { + const input_t v = input[linearIndex]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + const IndexType team = omp_get_team_num(); + #pragma omp atomic update + partial[(size_t)team * nbins + bin] += 1; + } + } + + // Threads own output bins, so the reduction needs no atomics. + #pragma omp target teams distribute parallel for + for (IndexType bin = 0; bin < nbins; bin++) { + output_t sum = 0; + for (IndexType histogram = 0; + histogram < activePartialHistograms; histogram++) { + const size_t offset = (size_t)histogram * nbins + bin; + sum += partial[offset]; + partial[offset] = 0; + } + output[bin] += sum; + } + } + } + const auto calibrationEnd = std::chrono::steady_clock::now(); + const long long candidateTime = + std::chrono::duration_cast( + calibrationEnd - calibrationStart).count(); + if (trial == 1 && (bestTime == 0 || candidateTime < bestTime)) { + bestTime = candidateTime; + bestConfiguration = activeConfiguration; } } } - auto end = std::chrono::steady_clock::now(); - auto time = std::chrono::duration_cast(end - start).count(); - printf("Average execution time of bincount kernel: %f (us)\n", - (time * 1e-3f) / repeat); - - #pragma omp target update from(output[0:output_size]) - - int status = memcmp(output, output_r, output_size_bytes); - printf("%s\n", status ? "FAIL" : "PASS"); - - if (sharedMem <= maxSharedMemory) { - printf("\n"); - printf("bincount using global and local atomics\n"); - - // number of teams mirrors the CUDA grid dimension - const IndexType numTeams = - (input_size + threadsPerBlock - 1) / threadsPerBlock; - #pragma omp target teams distribute parallel for - for (IndexType i = 0; i < output_size; i++) output[i] = 0; + #pragma omp target teams distribute parallel for + for (IndexType i = 0; i < output_size; i++) output[i] = 0; - start = std::chrono::steady_clock::now(); - for (int n = 0; n < repeat; n++) { - #pragma omp target teams num_teams(numTeams) thread_limit(threadsPerBlock) + if (bestConfiguration == 0) + printf("bincount using global atomics (auto-selected)\n"); + else if (bestConfiguration < 0) + printf("bincount using team-local histogram with %d teams " + "(auto-selected)\n", -bestConfiguration); + else + printf("bincount using %d global partial histograms (auto-selected)\n", + bestConfiguration); + + const auto start = std::chrono::steady_clock::now(); + if (bestConfiguration == 0) { + for (int iteration = 0; iteration < repeat; iteration++) { + #pragma omp target teams distribute parallel for + for (IndexType linearIndex = 0; + linearIndex < input_size; linearIndex++) { + const input_t v = input[linearIndex]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + #pragma omp atomic update + output[bin] += 1; + } + } + } + } else if (bestConfiguration < 0) { + const IndexType localTeams = -bestConfiguration; + for (int iteration = 0; iteration < repeat; iteration++) { + #pragma omp target teams num_teams(localTeams) \ + thread_limit(THREADS_PER_BLOCK) { - // Per-team histogram. Declaring it in the teams region makes it - // shared by the team's threads and lets the compiler place it in - // team-local memory (LDS/SLM). - output_t smem[sharedMemoryCapacity / sizeof(output_t)]; + output_t histogram[MAX_LOCAL_HISTOGRAM_BINS]; - #pragma omp parallel + #pragma omp parallel shared(histogram) { - const int nthreads = omp_get_num_threads(); - const int tid = omp_get_thread_num(); + const int thread = omp_get_thread_num(); + const int threads = omp_get_num_threads(); const int team = omp_get_team_num(); - const int nteams = omp_get_num_teams(); + const int teams = omp_get_num_teams(); - // zero the shared histogram - for (IndexType i = tid; i < nbins; i += nthreads) smem[i] = 0; + for (IndexType bin = thread; bin < nbins; bin += threads) + histogram[bin] = 0; #pragma omp barrier - // atomically accumulate into the shared histogram - for (IndexType linearIndex = (IndexType)team * nthreads + tid; - linearIndex < input_size; - linearIndex += (IndexType)nteams * nthreads) { - const input_t v = input[linearIndex]; + for (IndexType i = (IndexType)team * threads + thread; + i < input_size; i += (IndexType)teams * threads) { + const input_t v = input[i]; if (v >= minvalue && v <= maxvalue) { const IndexType bin = getBin( v, minvalue, maxvalue, nbins); #pragma omp atomic update - smem[bin] += 1; + histogram[bin] += 1; } } #pragma omp barrier - // flush the shared histogram to the global output - for (IndexType i = tid; i < nbins; i += nthreads) { + for (IndexType bin = thread; bin < nbins; bin += threads) { #pragma omp atomic update - output[i] += smem[i]; + output[bin] += histogram[bin]; } } } } - end = std::chrono::steady_clock::now(); - time = std::chrono::duration_cast(end - start).count(); - printf("Average execution time of bincount kernel: %f (us)\n", - (time * 1e-3f) / repeat); - - #pragma omp target update from(output[0:output_size]) + } else { + const IndexType bestPartialHistograms = bestConfiguration; + for (int iteration = 0; iteration < repeat; iteration++) { + #pragma omp target teams distribute parallel for \ + num_teams(bestPartialHistograms) \ + thread_limit(THREADS_PER_BLOCK) + for (IndexType linearIndex = 0; + linearIndex < input_size; linearIndex++) { + const input_t v = input[linearIndex]; + if (v >= minvalue && v <= maxvalue) { + const IndexType bin = getBin( + v, minvalue, maxvalue, nbins); + const IndexType team = omp_get_team_num(); + #pragma omp atomic update + partial[(size_t)team * nbins + bin] += 1; + } + } - int status = memcmp(output, output_r, output_size_bytes); - printf("%s\n", status ? "FAIL" : "PASS"); + #pragma omp target teams distribute parallel for + for (IndexType bin = 0; bin < nbins; bin++) { + output_t sum = 0; + for (IndexType histogram = 0; + histogram < bestPartialHistograms; histogram++) { + const size_t offset = (size_t)histogram * nbins + bin; + sum += partial[offset]; + partial[offset] = 0; + } + output[bin] += sum; + } + } } + const auto end = std::chrono::steady_clock::now(); + const long long time = + std::chrono::duration_cast( + end - start).count(); + printf("Average execution time of bincount kernel: %f (us)\n", + (time * 1e-3f) / repeat); + } + + int status = memcmp(output, output_r, output_size_bytes); + printf("%s\n", status ? "FAIL" : "PASS"); + + free(partial); - #pragma omp target exit data map(delete: output[0:output_size]) free(output); free(output_r); } - - #pragma omp target exit data map(delete: input[0:input_size]) + } free(input); } diff --git a/src/dwconv-cuda/main.cu b/src/dwconv-cuda/main.cu index ce681ee60..fdc67555a 100644 --- a/src/dwconv-cuda/main.cu +++ b/src/dwconv-cuda/main.cu @@ -227,7 +227,7 @@ void dwconv2d_forward (const int m, int errors = 0; for (int i = 0; i < output_size; i++) { const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); - if (fabs(h_output[i] - h_reference[i]) > tolerance) + if (!(fabs(h_output[i] - h_reference[i]) <= tolerance)) errors++; } printf("%s\n", errors == 0 ? "PASS" : "FAIL"); diff --git a/src/dwconv-hip/main.cu b/src/dwconv-hip/main.cu index c5a35e243..ae4e65896 100644 --- a/src/dwconv-hip/main.cu +++ b/src/dwconv-hip/main.cu @@ -227,7 +227,7 @@ void dwconv2d_forward (const int m, int errors = 0; for (int i = 0; i < output_size; i++) { const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); - if (fabs(h_output[i] - h_reference[i]) > tolerance) + if (!(fabs(h_output[i] - h_reference[i]) <= tolerance)) errors++; } printf("%s\n", errors == 0 ? "PASS" : "FAIL"); diff --git a/src/dwconv-omp/main.cpp b/src/dwconv-omp/main.cpp index 01c4d4aa8..a6b6c1a18 100644 --- a/src/dwconv-omp/main.cpp +++ b/src/dwconv-omp/main.cpp @@ -199,7 +199,7 @@ void dwconv2d_forward (const int m, int errors = 0; for (int i = 0; i < output_size; i++) { const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); - if (fabs(h_output[i] - h_reference[i]) > tolerance) + if (!(fabs(h_output[i] - h_reference[i]) <= tolerance)) errors++; } printf("%s\n", errors == 0 ? "PASS" : "FAIL"); diff --git a/src/dwconv-sycl/main.cpp b/src/dwconv-sycl/main.cpp index d633a950b..52ed8dafe 100644 --- a/src/dwconv-sycl/main.cpp +++ b/src/dwconv-sycl/main.cpp @@ -246,7 +246,7 @@ void dwconv2d_forward (sycl::queue &q, int errors = 0; for (int i = 0; i < output_size; i++) { const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]); - if (fabs(h_output[i] - h_reference[i]) > tolerance) + if (!(fabs(h_output[i] - h_reference[i]) <= tolerance)) errors++; } printf("%s\n", errors == 0 ? "PASS" : "FAIL"); diff --git a/src/pointerchase-omp/CMakeLists.txt b/src/pointerchase-omp/CMakeLists.txt deleted file mode 100644 index 03034a3bc..000000000 --- a/src/pointerchase-omp/CMakeLists.txt +++ /dev/null @@ -1,8 +0,0 @@ -# pointerchase-omp/CMakeLists.txt - -add_hecbench_benchmark( - NAME pointerchase - MODEL omp - SOURCES main.cpp - CATEGORIES algorithms -) diff --git a/src/pointerchase-omp/Makefile b/src/pointerchase-omp/Makefile deleted file mode 100644 index b5622e48a..000000000 --- a/src/pointerchase-omp/Makefile +++ /dev/null @@ -1,62 +0,0 @@ -#=============================================================================== -# User Options -#=============================================================================== - -# Compiler can be set below, or via environment variable -CC = icpx -OPTIMIZE = yes -DEBUG = no -DEVICE = gpu -LAUNCHER ?= - -#=============================================================================== -# Program name & source code list -#=============================================================================== - -program = main - -source = main.cpp - -obj = $(source:.cpp=.o) - -#=============================================================================== -# Sets Flags -#=============================================================================== - -# Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall - -# Linker Flags -LDFLAGS = - -# Debug Flags -ifeq ($(DEBUG),yes) - CFLAGS += -g - LDFLAGS += -g -endif - -# Optimization Flags -ifeq ($(OPTIMIZE),yes) - CFLAGS += -O3 -endif - -ifeq ($(DEVICE),gpu) - CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ -else - CFLAGS +=-qopenmp -endif -#=============================================================================== -# Targets to Build -#=============================================================================== - -$(program): $(obj) Makefile - $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) - -%.o: %.cpp Makefile - $(CC) $(CFLAGS) -c $< -o $@ - -clean: - rm -rf $(program) $(obj) - -run: $(program) - $(LAUNCHER) ./$(program) diff --git a/src/pointerchase-omp/Makefile.aomp b/src/pointerchase-omp/Makefile.aomp deleted file mode 100644 index 7ee3df7ee..000000000 --- a/src/pointerchase-omp/Makefile.aomp +++ /dev/null @@ -1,66 +0,0 @@ -#=============================================================================== -# User Options -#=============================================================================== - -# Compiler can be set below, or via environment variable -CC = clang++ -OPTIMIZE = yes -DEBUG = no -DEVICE = gpu -ARCH = gfx906 -LAUNCHER ?= - -#=============================================================================== -# Program name & source code list -#=============================================================================== - -program = main - -source = main.cpp - -obj = $(source:.cpp=.o) - -#=============================================================================== -# Sets Flags -#=============================================================================== - -# Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall - -# Linker Flags -LDFLAGS = - -# Debug Flags -ifeq ($(DEBUG),yes) - CFLAGS += -g - LDFLAGS += -g -endif - -# Optimization Flags -ifeq ($(OPTIMIZE),yes) - CFLAGS += -O3 -endif - -ifeq ($(DEVICE),gpu) - CFLAGS += -target x86_64-pc-linux-gnu \ - -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ - -Xopenmp-target=amdgcn-amd-amdhsa \ - -march=$(ARCH) -else - CFLAGS +=-fopenmp -endif -#=============================================================================== -# Targets to Build -#=============================================================================== - -$(program): $(obj) Makefile.aomp - $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) - -%.o: %.cpp Makefile.aomp - $(CC) $(CFLAGS) -c $< -o $@ - -clean: - rm -rf $(program) $(obj) - -run: $(program) - $(LAUNCHER) ./$(program) diff --git a/src/pointerchase-omp/Makefile.nvc b/src/pointerchase-omp/Makefile.nvc deleted file mode 100644 index f113f48f8..000000000 --- a/src/pointerchase-omp/Makefile.nvc +++ /dev/null @@ -1,63 +0,0 @@ -#=============================================================================== -# User Options -#=============================================================================== - -# Compiler can be set below, or via environment variable -CC = nvc++ -OPTIMIZE = yes -DEBUG = no -DEVICE = gpu -SM ?= cc70 -LAUNCHER ?= - -#=============================================================================== -# Program name & source code list -#=============================================================================== - -program = main - -source = main.cpp - -obj = $(source:.cpp=.o) - -#=============================================================================== -# Sets Flags -#=============================================================================== - -# Standard Flags -CFLAGS := $(EXTRA_CFLAGS) -std=c++17 -Wall - -# Linker Flags -LDFLAGS = - -# Debug Flags -ifeq ($(DEBUG),yes) - CFLAGS += -g - LDFLAGS += -g -endif - -# Optimization Flags -ifeq ($(OPTIMIZE),yes) - CFLAGS += -O3 -endif - -ifeq ($(DEVICE),gpu) - CFLAGS +=-Minfo -mp=gpu -gpu=$(SM) -else - CFLAGS +=-mp -endif -#=============================================================================== -# Targets to Build -#=============================================================================== - -$(program): $(obj) Makefile.nvc - $(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS) - -%.o: %.cpp Makefile.nvc - $(CC) $(CFLAGS) -c $< -o $@ - -clean: - rm -rf $(program) $(obj) - -run: $(program) - $(LAUNCHER) ./$(program) diff --git a/src/pointerchase-omp/main.cpp b/src/pointerchase-omp/main.cpp deleted file mode 100644 index 8b764af12..000000000 --- a/src/pointerchase-omp/main.cpp +++ /dev/null @@ -1,146 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include - -const uint64_t latencyMemAccessCnt = 1000000; /* 1M total read accesses to gauge latency */ -const uint32_t _2MiB = 2 * 1024 * 1024; -const uint32_t strideLen = 16; /* cacheLine size 128 Bytes, 16 words */ - -/* Upper bound on the number of teams used */ -const uint32_t maxTeamCount = 256; - -struct LatencyNode { - struct LatencyNode *next; -}; - -void initBuffer(void* buffer, uint64_t buffer_size, bool measureDeviceToDeviceLatency) { - uint64_t n_ptrs = buffer_size / sizeof(struct LatencyNode); - - if (measureDeviceToDeviceLatency) { - // For device-to-device latency, create and initialize pattern on device - const int dev = omp_get_default_device(); - const int host = omp_get_initial_device(); - for (uint64_t i = 0; i < n_ptrs; i++) { - struct LatencyNode node; - uint64_t nextOffset = ((i + strideLen) % n_ptrs) * sizeof(struct LatencyNode); - // Set up pattern with device addresses - node.next = (struct LatencyNode*)((uint8_t*)buffer + nextOffset); - int status = omp_target_memcpy(buffer, &node, sizeof(struct LatencyNode), - i * sizeof(struct LatencyNode), 0, - dev, host); - if (status != 0) { - fprintf(stderr, "omp_target_memcpy failed with status %d\n", status); - exit(EXIT_FAILURE); - } - } - } else { - // For host-device latency, initialize pattern with host addresses - struct LatencyNode* hostMem = (struct LatencyNode*)buffer; - for (uint64_t i = 0; i < n_ptrs; i++) { - hostMem[i].next = &hostMem[(i + strideLen) % n_ptrs]; - } - } -} - - -double latencyPtrChaseKernel(void* data, uint64_t memAccessCnt, - uint32_t smCount) -{ - double latencySum = 0.0f; - uint32_t measuredTeamCount = 0; - struct LatencyNode *nodes = static_cast(data); - - // For smCount teams, each team has memAccessCnt pointer chases - for (uint32_t targetBlock = 0; targetBlock < smCount; ++targetBlock) { - int executed = 0; - auto start = std::chrono::steady_clock::now(); - - // The body of a teams region is executed by one thread per team - #pragma omp target teams num_teams(smCount) thread_limit(1) \ - is_device_ptr(nodes) map(tofrom: executed) - { - if ((uint32_t)omp_get_team_num() == targetBlock) { - executed = 1; - struct LatencyNode *p = nodes; - for (uint32_t i = 0; i < memAccessCnt; ++i) { - p = p->next; - } - - // Avoid compiler optimization: the store is never reached, but the - // compiler cannot prove it and therefore has to keep the pointer chase - // above. An assert() would not do, as it is compiled out when NDEBUG is - // defined. - if (p == nullptr) { - nodes[0].next = nullptr; - } - } - } - - auto end = std::chrono::steady_clock::now(); - if (executed) { - auto latency = - std::chrono::duration_cast(end - start) - .count(); - latencySum += latency; - ++measuredTeamCount; - } - } - if (measuredTeamCount == 0) { - fprintf(stderr, "No OpenMP target team executed the pointer chase\n"); - exit(EXIT_FAILURE); - } - return latencySum / - (memAccessCnt * measuredTeamCount); // finalLatencyPerAccessNs -} - -class MemPtrChaseOperation { - public: - MemPtrChaseOperation() { - // OpenMP has no portable query for the number of compute units - int teams = 0; - #pragma omp target teams num_teams(maxTeamCount) thread_limit(1) \ - map(tofrom: teams) - { - if (omp_get_team_num() == 0) teams = omp_get_num_teams(); - } - if (teams < 1) teams = 1; - smCount = (uint32_t)teams; - if (smCount > maxTeamCount) smCount = maxTeamCount; - } - ~MemPtrChaseOperation() = default; - double doPtrChase(void* peerBuffer) { - double lat = - latencyPtrChaseKernel(peerBuffer, latencyMemAccessCnt, smCount); - return lat; - } - private: - uint32_t smCount; -}; - -int main() { - const uint64_t buffer_size = _2MiB; - const bool measureDeviceToDeviceLatency = true; - - void *buffer = omp_target_alloc(buffer_size, omp_get_default_device()); - if (buffer == nullptr) { - fprintf(stderr, "Failed to allocate %lu bytes on the device\n", - (unsigned long) buffer_size); - return 1; - } - - // initialize the buffer - initBuffer(buffer, buffer_size, measureDeviceToDeviceLatency); - - // compute the latency of pointer chasing on the default device - MemPtrChaseOperation mpc; - - double lat = mpc.doPtrChase(buffer); - - printf("Latency per access on device: %lf (ns)\n", lat); - omp_target_free(buffer, omp_get_default_device()); - return 0; -}