Adds API for non-caching-store and load

CMakeLists.txt

@@ -399,6 +399,9 @@ set(SRC_FILES
   src/device/broadcast.h
   src/device/common.h
   src/device/common_kernel.h
+  src/device/non_caching_load.h
+  src/device/non_caching_store.h
+  src/device/non_caching_store_vec4.h
   src/device/op128.h
   src/device/primitives.h
   src/device/prims_ll128.h

src/device/non_caching_load.h

@@ -0,0 +1,77 @@
+#ifndef NON_CACHING_LOAD_H_
+#define NON_CACHING_LOAD_H_
+
+template<typename T>
+inline
+__attribute__((always_inline))
+__host__ __device__ T __non_caching_load(const T* p)
+{
+    #if !defined(__GFX11__) && !defined(GFX12)
+        #define LD "global_load_ubyte"
+        #define LD2 "global_load_ushort"
+        #define LD3 "global_load_dword"
+        #define LD4 "global_load_dwordx2"
+        #define LD5 "global_load_dwordx4"
+        #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+            #define BITS "sc0 sc1 nt"
+        #elif defined(__GFX9__) || defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || defined(__gfx1013__)
+            #define BITS "glc slc"
+        #else
+            #define BITS "glc slc dlc"
+        #endif
+        #define WAIT ((0 << 14) | (0x3f << 8) | (0x7) << 4)
+    #else
+        #define LD "global_load_u8"
+        #define LD2 "global_load_u16"
+        #define LD3 "global_load_b32"
+        #define LD4 "global_load_b64"
+        #define LD5 "global_load_b128"
+        #define BITS "glc slc dlc"
+        #define WAIT ((0 << 10) | (0x3f << 4) | 0x7)
+    #endif
+    #define LOAD LD " %0 %1 off " BITS
+    #define LOAD2 LD2 " %0 %1 off " BITS
+    #define LOAD3 LD3 " %0 %1 off " BITS
+    #define LOAD4 LD4 " %0 %1 off " BITS
+    #define LOAD5 LD5 " %0 %1 off " BITS
+
+    T r{};
+
+    switch (sizeof(T)) {
+    case 1:
+        asm volatile(LOAD : "={v0}"(r) : "v"(p));
+        break;
+    case 2:
+        asm volatile(LOAD2 : "={v0}"(r) : "v"(p));
+        break;
+    case 4:
+        asm volatile(LOAD3 : "={v0}"(r) : "v"(p));
+        break;
+    case 8:
+        asm volatile(LOAD4 : "={v[0:1]}"(r) : "v"(p));
+        break;
+    case 16:
+        asm volatile(LOAD5 : "=v"(r) : "v"(p));
+        break;
+    default: __builtin_trap();
+    }
+    __builtin_amdgcn_s_waitcnt(WAIT);
+
+    return r;
+
+    #undef LOAD5
+    #undef LOAD4
+    #undef LOAD3
+    #undef LOAD2
+    #undef LOAD
+    #undef WAIT
+    #undef BITS
+    #undef LD5
+    #undef LD4
+    #undef LD3
+    #undef LD2
+    #undef LD
+}
+
+#endif
+

src/device/non_caching_store.h

@@ -0,0 +1,62 @@
+#ifndef NON_CACHING_STORE_H_
+#define NON_CACHING_STORE_H_
+
+template<typename T>
+inline
+__attribute__((always_inline))
+__host__ __device__ T __non_caching_store(const T val, const T* p)
+{
+    #if !defined(__GFX11__) && !defined(GFX12)
+        #define ST "global_store_byte"
+        #define ST2 "global_store_short"
+        #define ST3 "global_store_dword"
+        #define ST4 "global_store_dwordx2"
+        #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+            #define BITS "sc0 sc1 nt"
+        #elif defined(__GFX9__) || defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || defined(__gfx1013__)
+            #define BITS "glc slc"
+        #else
+            #define BITS "glc slc dlc"
+        #endif
+    #else
+        #define ST "global_store_b8"
+        #define ST2 "global_store_b16"
+        #define ST3 "global_store_b32"
+        #define ST4 "global_store_b64"
+        #define BITS "glc slc dlc"
+    #endif
+    #define STORE ST " %0 %1 %2 " BITS
+    #define STORE2 ST2 " %0 %1 %2 " BITS
+    #define STORE3 ST3 " %0 %1 %2 " BITS
+    #define STORE4 ST4 " %0 %1 %2 " BITS
+
+    switch (sizeof(T)) {
+    case 1:
+        asm volatile(STORE :: "v"(0), "v"(uint32_t(val)) , "s"(p));
+        break;
+    case 2:
+        asm volatile(STORE2 :: "v"(0), "v"(val) , "s"(p));
+        break;
+    case 4:
+        asm volatile(STORE3 :: "v"(0), "v"(val) , "s"(p));
+        break;
+    case 8:
+        asm volatile(STORE4 :: "v"(0), "v"(val) , "s"(p));
+        break;
+    default: __builtin_trap();
+    }
+    asm volatile("s_endpgm");
+
+    #undef STORE4
+    #undef STORE3
+    #undef STORE2
+    #undef STORE
+    #undef BITS
+    #undef ST4
+    #undef ST3
+    #undef ST2
+    #undef ST
+}
+
+#endif
+

src/device/non_caching_store_vec4.h

@@ -0,0 +1,34 @@
+#ifndef NON_CACHING_STORE_VEC4_H_
+#define NON_CACHING_STORE_VEC4_H_
+
+template<typename T>
+inline
+__attribute__((always_inline))
+__host__ __device__ T __non_caching_store_vec4(const T val, const T* p)
+{
+    #if !defined(__GFX11__) && !defined(GFX12)
+        #define ST "global_store_dwordx4"
+        #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+            #define BITS "sc0 sc1 nt"
+        #elif defined(__GFX9__) || defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || defined(__gfx1013__)
+            #define BITS "glc slc"
+        #else
+            #define BITS "glc slc dlc"
+        #endif
+    #else
+        #define ST "global_store_b128"
+        #define BITS "glc slc dlc"
+    #endif
+
+    #define STORE ST " %0 %1 %2 " BITS
+
+    asm volatile(STORE :: "v"(0), "v"(val) , "s"(p));
+    asm volatile("s_endpgm");
+
+    #undef STORE
+    #undef BITS
+    #undef ST
+}
+
+#endif
+

tools/non-caching-load/Makefile

@@ -0,0 +1,22 @@
+# Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+
+# Set to where RCCL is installed
+RCCL_INSTALL=../../build/release
+
+HIP_PATH?= $(wildcard /opt/rocm)
+ifeq (,$(HIP_PATH))
+HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=non-caching-load
+CXXFLAGS = -std=c++11 -O3 -I../../src/include  -I../../src/device -I$(RCCL_INSTALL) -L$(RCCL_INSTALL) -lrccl
+
+all: $(EXE)
+
+$(EXE): $(EXE).cpp $(shell find -regex ".*\.\hpp")
+	$(HIPCC) $(CXXFLAGS) $< -o $@
+
+clean:
+	rm -f *.o $(EXE)
+

tools/non-caching-load/non-caching-load.cpp

@@ -0,0 +1,116 @@
+/*
+Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <sys/socket.h>
+#include <ifaddrs.h>
+#include <netdb.h>
+#include <unistd.h>
+#include <cstdio>
+#include <string>
+#include <chrono>
+#include <hip/hip_runtime.h>
+#include <rccl/rccl.h>
+#include <cstdlib>
+#include <fstream>
+#include <iostream> //cerr
+#include <cstring>
+#include "non-caching-load.hpp"
+#include "non_caching_load.h"
+
+typedef uint32_t uint32x4 __attribute__((ext_vector_type(4)));
+
+template<typename T>
+__global__ void nonCachingLoad(T* p, T* out){
+  if constexpr (std::is_same<T, uint32x4>::value)
+    p[0] = {22, 22, 22, 22};
+  else
+    p[0] = 22;
+	out[0] = __non_caching_load<T>(p);
+}
+
+template<typename T>
+__global__ void builtinTemporalLoad(T* p, T* out){
+  if constexpr (std::is_same<T, uint32x4>::value)
+    p[0] = {22, 22, 22, 22};
+  else
+    p[0] = 22;
+	out[0] = __builtin_nontemporal_load(p);
+}
+
+template<typename T>
+void caching_load() {
+  T* data;
+  T* out1;
+  T* out2;
+  size_t size = sizeof(data);
+
+  hipMalloc(&data, size);
+  hipMalloc(&out1, size);
+  hipMalloc(&out2, size);
+
+  hipLaunchKernelGGL(nonCachingLoad<T>, dim3(1), dim3(1), 0, 0, data, out1);
+  hipLaunchKernelGGL(builtinTemporalLoad<T>, dim3(1), dim3(1), 0, 0, data, out2);
+
+  hipDeviceSynchronize();
+
+  T* host_data = (T*)malloc(size);
+  T* h_o1 = (T*)malloc(size);
+  T* h_o2 = (T*)malloc(size);
+
+  hipMemcpy(host_data, data, size, hipMemcpyDeviceToHost);
+  hipMemcpy(h_o1, out1, size, hipMemcpyDeviceToHost);
+  hipMemcpy(h_o2, out2, size, hipMemcpyDeviceToHost);
+
+  if constexpr (std::is_same<T, uint32x4>::value)
+  {
+    if ( ((*h_o1)[0] == (*h_o2)[0]) && ((*h_o1)[1] == (*h_o2)[1]) && ((*h_o1)[2] == (*h_o2)[2])
+     && ((*h_o1)[3] == (*h_o2)[3]))
+      std::cout << "PASS" << std::endl;
+    else
+      std::cout << "FAIL" << std::endl; 
+  }
+  else {
+    if(*h_o1 == *h_o2)
+    {
+      std::cout << "PASS" << std::endl; 
+    }
+    else{
+      std::cout << "FAIL" << std::endl;
+    }
+  }
+
+  hipFree(data);
+  return;
+}
+
+int main(int argc, char **argv)
+{
+  caching_load<uint64_t>();
+  caching_load<uint32_t>();
+  caching_load<uint16_t>();
+  caching_load<uint8_t>();
+  using V2 = unsigned __attribute__((ext_vector_type(4)));
+  caching_load<V2>();
+
+  return 0;
+}
+

tools/non-caching-load/non-caching-load.hpp

@@ -0,0 +1,50 @@
+/*
+Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HELLORCCL_HPP
+#define HELLORCCL_HPP
+#include <iostream>
+
+#define HIP_CALL(cmd)                                                 \
+  do {                                                                \
+    hipError_t error = (cmd);                                         \
+    if (error != hipSuccess)                                          \
+    {                                                                   \
+      std::cerr << "Encountered HIP error (" << hipGetErrorString(error) << ") at line " \
+                << __LINE__ << " in file " << __FILE__ << "\n";         \
+      exit(-1);                                                         \
+    }                                                                   \
+  } while (0)
+
+#define NCCL_CALL(cmd) \
+  do { \
+    ncclResult_t error = (cmd);                 \
+    if (error != ncclSuccess)                   \
+    {                                           \
+      std::cerr << "Encountered NCCL error (" << ncclGetErrorString(error) << ") at line " \
+                << __LINE__ << " in file " << __FILE__ << "\n";         \
+      exit(-1);                                                         \
+    }                                                                   \
+  } while (0)
+
+#endif
+

tools/non-caching-store/Makefile

@@ -0,0 +1,22 @@
+# Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+
+# Set to where RCCL is installed
+RCCL_INSTALL=../../build/release
+
+HIP_PATH?= $(wildcard /opt/rocm)
+ifeq (,$(HIP_PATH))
+HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=non-caching-store
+CXXFLAGS = -std=c++11 -O3 -I../../src/include  -I../../src/device -I$(RCCL_INSTALL)/include -L$(RCCL_INSTALL) -lrccl
+
+all: $(EXE)
+
+$(EXE): $(EXE).cpp $(shell find -regex ".*\.\hpp")
+	$(HIPCC) $(CXXFLAGS) $< -o $@
+
+clean:
+	rm -f *.o $(EXE)
+