Replace aligned_ptr by vector_ptr

Author: stijnh

examples/vector_add/main.cu

@@ -13,13 +13,15 @@ void cuda_check(cudaError_t code) {
 }
 
 template<int N>
-__global__ void my_kernel(int length, const __half* input, double constant, float* output) {
+__global__ void my_kernel(
+    int length,
+    kf::vec_ptr<const __half, N> input,
+    double constant,
+    kf::vec_ptr<float, N> output) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
 
     if (i * N < length) {
-        auto a = kf::read_aligned<N>(input + i * N);
-        auto b = kf::fma(a, a, kf::cast<__half>(constant));
-        kf::write_aligned<N>(output + i * N, b);
+        output[i] = kf::fma(input[i], input[i], kf::cast<__half>(constant));
     }
 }
 
@@ -51,9 +53,9 @@ void run_kernel(int n) {
     int grid_size = (n + items_per_block - 1) / items_per_block;
     my_kernel<items_per_thread><<<grid_size, block_size>>>(
         n,
-        kf::aligned_ptr(input_dev),
+        kf::vector_ptr<const half, items_per_thread>(input_dev),
         constant,
-        kf::aligned_ptr(output_dev));
+        kf::vector_ptr<float, items_per_thread>(output_dev));
 
     // Copy results back
     cuda_check(cudaMemcpy(output_dev, output_result.data(), sizeof(float) * n, cudaMemcpyDefault));

examples/vector_add_tiling/main.cu

@@ -14,11 +14,7 @@ void cuda_check(cudaError_t code) {
 }
 
 template<int N, int B>
-__global__ void my_kernel(
-    int length,
-    kf::aligned_ptr<const __half> input,
-    double constant,
-    kf::aligned_ptr<float> output) {
+__global__ void my_kernel(int length, const __half* input, double constant, float* output) {
     auto tiling = kf::tiling<
         kf::tile_factor<N>,
         kf::block_size<B>,
@@ -27,9 +23,9 @@ __global__ void my_kernel(
     auto points = int(blockIdx.x * tiling.tile_size(0)) + tiling.local_points(0);
     auto mask = tiling.local_mask();
 
-    auto a = input.read(points, mask);
+    auto a = kf::read(input, points, mask);
     auto b = (a * a) * constant;
-    output.write(points, b, mask);
+    kf::write(output, points, b, mask);
 }
 
 template<int items_per_thread, int block_size = 256>
@@ -57,11 +53,8 @@ void run_kernel(int n) {
     // Launch kernel!
     int items_per_block = block_size * items_per_thread;
     int grid_size = (n + items_per_block - 1) / items_per_block;
-    my_kernel<items_per_thread, block_size><<<grid_size, block_size>>>(
-        n,
-        kf::aligned_ptr(input_dev),
-        constant,
-        kf::aligned_ptr(output_dev));
+    my_kernel<items_per_thread, block_size>
+        <<<grid_size, block_size>>>(n, input_dev, constant, output_dev);
 
     // Copy results back
     cuda_check(cudaMemcpy(output_dev, output_result.data(), sizeof(float) * n, cudaMemcpyDefault));