Add __restrict__

poyenc · poyenc · commit 6eb3fbc6f13d · 2025-09-09T04:46:33.000-05:00
diff --git a/include/ck/utility/amd_buffer_addressing_builtins.hpp b/include/ck/utility/amd_buffer_addressing_builtins.hpp
@@ -20,7 +20,7 @@ union BufferResource
 };
 
 template <typename T>
-__device__ int32x4_t make_wave_buffer_resource(T* p_wave, index_t element_space_size)
+__device__ int32x4_t make_wave_buffer_resource(T* __restrict__ p_wave, index_t element_space_size)
 {
     BufferResource<T> wave_buffer_resource;
 
@@ -35,7 +35,7 @@ __device__ int32x4_t make_wave_buffer_resource(T* p_wave, index_t element_space_
 }
 
 template <typename T>
-__device__ int32x4_t make_wave_buffer_resource_with_default_range(T* p_wave)
+__device__ int32x4_t make_wave_buffer_resource_with_default_range(T* __restrict__ p_wave)
 {
     BufferResource<T> wave_buffer_resource;
 
@@ -711,7 +711,7 @@ template <typename T,
           index_t N,
           AmdBufferCoherenceEnum coherence = AmdBufferCoherenceEnum::DefaultCoherence>
 __device__ void amd_buffer_store(const typename vector_type_maker<T, N>::type::type src_thread_data,
-                                 T* p_dst_wave,
+                                 T* __restrict__ p_dst_wave,
                                  const index_t dst_thread_element_offset,
                                  const bool dst_thread_element_valid,
                                  const index_t dst_element_space_size)
diff --git a/include/ck/utility/c_style_pointer_cast.hpp b/include/ck/utility/c_style_pointer_cast.hpp
@@ -12,7 +12,7 @@ namespace ck {
 template <typename PY,
           typename PX,
           typename enable_if<is_pointer_v<PY> && is_pointer_v<PX>, bool>::type = false>
-__host__ __device__ PY c_style_pointer_cast(PX p_x)
+__host__ __device__ PY c_style_pointer_cast(PX __restrict__ p_x)
 {
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Wold-style-cast"
diff --git a/include/ck_tile/core/arch/amd_buffer_addressing.hpp b/include/ck_tile/core/arch/amd_buffer_addressing.hpp
@@ -37,7 +37,7 @@ struct __attribute__((packed)) buffer_resource
     uint32_t config;
 };
 
-CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* ptr, uint32_t size = 0xffffffff)
+CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* __restrict__ ptr, uint32_t size = 0xffffffff)
 {
     buffer_resource res{ptr, size, CK_TILE_BUFFER_RESOURCE_3RD_DWORD};
     int32x4_t r = __builtin_bit_cast(int32x4_t, res);
diff --git a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp
@@ -28,7 +28,7 @@ struct __attribute__((packed)) buffer_resource
     uint32_t config;
 };
 
-CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* ptr, uint32_t size = 0xffffffff)
+CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* __restrict__  ptr, uint32_t size = 0xffffffff)
 {
     buffer_resource res{ptr, size, CK_TILE_BUFFER_RESOURCE_3RD_DWORD};
     int32x4_t r = __builtin_bit_cast(int32x4_t, res);
diff --git a/include/ck_tile/core/tensor/load_tile_transpose.hpp b/include/ck_tile/core/tensor/load_tile_transpose.hpp
@@ -403,7 +403,7 @@ CK_TILE_DEVICE auto
 load_tile_transpose(const tile_window_with_static_distribution<BottomTensorView_,
                                                                WindowLengths_,
                                                                TileDistribution_,
-                                                               NumCoord>& tile_window)
+                                                               NumCoord>& __restrict__ tile_window)
 {
     using OutTileDstrEncode = typename OutputTileDistributionTraits<
         typename TileDistribution_::DstrEncode,

Original file line number	Diff line number	Diff line change
`@@ -12,7 +12,7 @@ namespace ck {`
`12`	`12`	`template <typename PY,`
`13`	`13`	`typename PX,`
`14`	`14`	`typename enable_if<is_pointer_v<PY> && is_pointer_v<PX>, bool>::type = false>`
`15`		`-__host__ __device__ PY c_style_pointer_cast(PX p_x)`
	`15`	`+__host__ __device__ PY c_style_pointer_cast(PX __restrict__ p_x)`
`16`	`16`	`{`
`17`	`17`	`#pragma clang diagnostic push`
`18`	`18`	`#pragma clang diagnostic ignored "-Wold-style-cast"`
Original file line number	Diff line number	Diff line change
`@@ -37,7 +37,7 @@ struct __attribute__((packed)) buffer_resource`
`37`	`37`	`uint32_t config;`
`38`	`38`	`};`
`39`	`39`
`40`		`-CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* ptr, uint32_t size = 0xffffffff)`
	`40`	`+CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* __restrict__ ptr, uint32_t size = 0xffffffff)`
`41`	`41`	`{`
`42`	`42`	`buffer_resource res{ptr, size, CK_TILE_BUFFER_RESOURCE_3RD_DWORD};`
`43`	`43`	`int32x4_t r = __builtin_bit_cast(int32x4_t, res);`
Original file line number	Diff line number	Diff line change
`@@ -28,7 +28,7 @@ struct __attribute__((packed)) buffer_resource`
`28`	`28`	`uint32_t config;`
`29`	`29`	`};`
`30`	`30`
`31`		`-CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* ptr, uint32_t size = 0xffffffff)`
	`31`	`+CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* __restrict__ ptr, uint32_t size = 0xffffffff)`
`32`	`32`	`{`
`33`	`33`	`buffer_resource res{ptr, size, CK_TILE_BUFFER_RESOURCE_3RD_DWORD};`
`34`	`34`	`int32x4_t r = __builtin_bit_cast(int32x4_t, res);`
Original file line number	Diff line number	Diff line change
`@@ -403,7 +403,7 @@ CK_TILE_DEVICE auto`
`403`	`403`	`load_tile_transpose(const tile_window_with_static_distribution<BottomTensorView_,`
`404`	`404`	`WindowLengths_,`
`405`	`405`	`TileDistribution_,`
`406`		`- NumCoord>& tile_window)`
	`406`	`+ NumCoord>& __restrict__ tile_window)`
`407`	`407`	`{`
`408`	`408`	`using OutTileDstrEncode = typename OutputTileDistributionTraits<`
`409`	`409`	`typename TileDistribution_::DstrEncode,`