<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class="">
<div class="">Hi everyone! </div>
<div class=""><br class="">
</div>
<div class="">I am following Jonas Hahnfeld’s blog post (<a href="https://www.hahnjo.de/blog/2018/10/08/clang-7.0-openmp-offloading-nvidia.html" class="">https://www.hahnjo.de/blog/2018/10/08/clang-7.0-openmp-offloading-nvidia.html</a>) to build clang to test
some simple target offloading workloads. </div>
<div class=""><br class="">
</div>
<div class="">The complete code can be found at the end. </div>
<div class=""><br class="">
</div>
<div class="">I tried clang 7.0.1 and llvm trunk (9.0.0). omp_get_num_devices() returned 0 for both compilers although the server I am running this on has Quadro P4000 on it. As an attempt to fix the issue, I added compute compatibility 61 and set it to default.
This did not help. </div>
<div class=""><br class="">
</div>
<div class="">Additionally, the binary built by trunk cannot even offload to CPU targets, reporting error </div>
<div class=""><br class="">
</div>
<div class="">"Libomptarget fatal error 1: default offloading policy must switched to mandatory or disabled”</div>
<div class=""><br class="">
</div>
<div class="">May I get some help from the community on what to look into to get this microbenchmark to run on the GPU? </div>
<div class=""><br class="">
</div>
<div class="">Thanks! </div>
<div class=""><br class="">
</div>
<div class="">######################################################################################################</div>
<div class="">
<div class="">#include <stdio.h></div>
<div class="">#include <stdlib.h></div>
<div class="">#include <time.h></div>
<div class="">#include <omp.h></div>
<div class=""><br class="">
</div>
<div class="">#define VALUE_TYPE double</div>
<div class="">#define ARR_SIZE 128 * 1024 * 1024</div>
<div class=""><br class="">
</div>
<div class="">struct timespec timespec_diff(struct timespec t1, struct timespec t2) {</div>
<div class=""> struct timespec diff;</div>
<div class=""><br class="">
</div>
<div class=""> diff.tv_nsec = t1.tv_nsec - t2.tv_nsec;</div>
<div class=""> diff.tv_sec = t1.tv_sec - t2.tv_sec;</div>
<div class=""> if (diff.tv_nsec < 0) {</div>
<div class=""> diff.tv_nsec += 1000000000;</div>
<div class=""> diff.tv_sec -= 1;</div>
<div class=""> }</div>
<div class=""><br class="">
</div>
<div class=""> return diff;</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">double timespec_to_sec(struct timespec t) {</div>
<div class=""> double sec = (double)t.tv_sec;</div>
<div class=""> double nsec = (double)t.tv_nsec;</div>
<div class=""> return sec + nsec * 1E-9;</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">void gen_rand_number(VALUE_TYPE *arr, int size) {</div>
<div class=""> for (int i = 0; i < size; i++) {</div>
<div class=""> arr[i] = (VALUE_TYPE)rand();</div>
<div class=""> }</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">void cpu_add(VALUE_TYPE *out, VALUE_TYPE *in_1, VALUE_TYPE *in_2, int size) {</div>
<div class=""> for (int i = 0; i < size; i++) {</div>
<div class=""> out[i] = in_1[i] + in_2[i];</div>
<div class=""> }</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class=""><br class="">
</div>
<div class="">void omp_cpu_add(VALUE_TYPE *out, VALUE_TYPE *in_1, VALUE_TYPE *in_2, int size) {</div>
<div class="">#pragma omp parallel for</div>
<div class=""> for (int i = 0; i < size; i++) {</div>
<div class=""> out[i] = in_1[i] + in_2[i];</div>
<div class=""> }</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">void omp_gpu_add(VALUE_TYPE *out, VALUE_TYPE *in_1, VALUE_TYPE *in_2, int size) {</div>
<div class="">#pragma omp target map(to:in_1, in_2) map(from:out)</div>
<div class="">#pragma omp teams</div>
<div class="">#pragma omp distribute parallel for</div>
<div class=""> for (int i = 0; i < size; i++) {</div>
<div class=""> out[i] = in_1[i] + in_2[i];</div>
<div class=""> }</div>
<div class="">}</div>
<div class=""><br class="">
</div>
<div class="">int main() {</div>
<div class=""> VALUE_TYPE *a = malloc(sizeof(VALUE_TYPE) * ARR_SIZE);</div>
<div class=""> VALUE_TYPE *b = malloc(sizeof(VALUE_TYPE) * ARR_SIZE);</div>
<div class=""> VALUE_TYPE *c_1 = malloc(sizeof(VALUE_TYPE) * ARR_SIZE);</div>
<div class=""> VALUE_TYPE *c_2 = malloc(sizeof(VALUE_TYPE) * ARR_SIZE);</div>
<div class=""> VALUE_TYPE *c_3 = malloc(sizeof(VALUE_TYPE) * ARR_SIZE);</div>
<div class=""><br class="">
</div>
<div class=""> struct timespec start, end, diff;</div>
<div class=""><br class="">
</div>
<div class=""> gen_rand_number(a, ARR_SIZE);</div>
<div class=""> gen_rand_number(b, ARR_SIZE);</div>
<div class=""><br class="">
</div>
<div class=""> printf("number of omp devices %d\n", omp_get_num_devices());</div>
<div class=""> printf("number of omp teams %d\n", omp_get_num_teams());</div>
<div class=""><br class="">
</div>
<div class=""> clock_gettime(CLOCK_MONOTONIC, &start);</div>
<div class=""> cpu_add(c_1, a, b, ARR_SIZE);</div>
<div class=""> clock_gettime(CLOCK_MONOTONIC, &end);</div>
<div class=""> printf("CPU took %f seconds. \n", timespec_to_sec(timespec_diff(end, start)));</div>
<div class=""><br class="">
</div>
<div class=""> clock_gettime(CLOCK_MONOTONIC, &start);</div>
<div class=""> omp_cpu_add(c_2, a, b, ARR_SIZE);</div>
<div class=""> clock_gettime(CLOCK_MONOTONIC, &end);</div>
<div class=""> printf("CPU omp took %f seconds. \n", timespec_to_sec(timespec_diff(end, start)));</div>
<div class=""><br class="">
</div>
<div class=""> clock_gettime(CLOCK_MONOTONIC, &start);</div>
<div class=""> for (int i = 0; i < 100; i++) {</div>
<div class=""> omp_gpu_add(c_3, a, b, ARR_SIZE);</div>
<div class=""> }</div>
<div class=""> clock_gettime(CLOCK_MONOTONIC, &end);</div>
<div class=""> printf("GPU omp offload took %f seconds. \n", timespec_to_sec(timespec_diff(end, start)) / 100.0);</div>
<div class="">}</div>
</div>
<div class=""><br class="">
</div>
<div class="">
<div class="">######################################################################################################</div>
<div class=""></div>
</div>
<br class="">
<div class="">
<div dir="auto" style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class="">
<div style="caret-color: rgb(0, 0, 0); color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;">
Sincerely,</div>
<div style="caret-color: rgb(0, 0, 0); color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;">
Qiongsi</div>
<div style="caret-color: rgb(0, 0, 0); color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class="">
<br class="">
</div>
<br class="Apple-interchange-newline">
</div>
<br class="Apple-interchange-newline">
</div>
<br class="">
</body>
</html>