diff --git a/ds/25-1/1e/Makefile b/ds/25-1/1e/Makefile
new file mode 100644
index 0000000..88785e1
--- /dev/null
+++ b/ds/25-1/1e/Makefile
@@ -0,0 +1,29 @@
+CC = nvcc -arch=sm_75
+DEBUG ?= false
+DIRS = dist build
+
+ifeq ($(DEBUG), false)
+    CC += -O3
+else
+    CC += -g -G
+endif
+
+.PHONY: all run
+
+all: $(DIRS) dist/app
+
+dist/app: build/main.o build/op.o
+	$(CC) $^ -o $@ -lcuda
+
+build/op.o: op.ptx
+	$(CC) $^ -dc -o $@
+
+build/main.o: main.cu
+	$(CC) $^ -ptx -o build/main.ptx
+	$(CC) $^ -rdc=true -dc -o $@
+
+$(DIRS):
+	mkdir -p $@
+
+clean:
+	rm -rf $(DIRS)
\ No newline at end of file
diff --git a/ds/25-1/1e/README.md b/ds/25-1/1e/README.md
index 3bc2693..77d21e3 100644
--- a/ds/25-1/1e/README.md
+++ b/ds/25-1/1e/README.md
@@ -1,3 +1,23 @@
 [euclidean](https://web.archive.org/web/20230212044931/http://www-math.ucdenver.edu/~wcherowi/courses/m5410/exeucalg.html)
 [ecdsa1](https://sefiks.com/2018/02/16/elegant-signatures-with-elliptic-curve-cryptography/)
 [ecdsa2](https://learnmeabitcoin.com/technical/cryptography/elliptic-curve/ecdsa/)
+[ptx](https://philipfabianek.com/posts/cuda-ptx-introduction)
+
+высокий приоритет
+
+6, 7	State Spaces / Properties of State Spaces	Ключевое отличие от CPU! В CPU память в основном плоская (RAM, кэш). В GPU есть много типов памяти: глобальная (.global), общая для блока потоков (.shared), константная (.const), локальная (.local) и т.д. Это фундамент для написания производительного кода.
+19	Cost Estimates for Accessing State-Spaces	Прямое продолжение предыдущего пункта. Объясняет, какая память быстрая, а какая медленная. Критично для оптимизации.
+4	Operator Precedence	Синтаксис PTX похож на ассемблер, но с выражениями. Знать приоритет операторов необходимо.
+8	Fundamental Type Specifiers	Типы данных в PTX (.b8, .s16, .f32, .b64 и т.д.). Аналог byte, word, dword в x86, но с учетом специфики GPU.
+3	Predefined Identifiers	Предопределенные константы, такие как %tid, %ctaid, %ntid. Это основа модели выполнения CUDA! Вместо одного потока (RIP/EIP) у вас есть идентификаторы потока, блока и сетки.
+20	Operation Types	Классификация инструкций PTX. Поможет быстро ориентироваться в мануале.
+1	PTX Directives	Директивы ассемблера (.version, .target, .global). Аналог секций и директив в NASM (SECTION .text, global _start)
+
+средний приоритет
+
+21	Scopes	Области видимости для атомарных операций и барьеров (.cta, .cluster, .gpu, .sys). Важно для синхронизации.
+14, 40, 56	Различные таблицы про Swizzling и Layout	Касаются продвинутых техник работы с памятью и матрицами для оптимизации доступа. Актуально для low-level оптимизаций, похоже на работу с выравниванием и SIMD в x86.
+29	Summary of Floating-Point Instructions	Обзор инструкций для чисел с плавающей точкой. На GPU они крайне важны.
+30-32	Cache Operators / Eviction Priority Hints	Управление кэшем. Продвинутая тема для тонкой настройки, аналогичная prefetch-инструкциям в x86.
+53, 55, 56	Таблицы про MMA (Matrix Multiply-Accumulate)	Инструкции для тензорных ядер (аналог FMA в x86, но для матриц). Сердце производительности в AI/HPC.
+22-25	Comparison Operators	Особенности сравнений для целых и вещественных чисел (учет NaN).
diff --git a/ds/25-1/1e/main.cu b/ds/25-1/1e/main.cu
new file mode 100644
index 0000000..c39ed0d
--- /dev/null
+++ b/ds/25-1/1e/main.cu
@@ -0,0 +1,35 @@
+#include <stdio.h>
+
+extern "C" __device__ void add_u32(
+    ulonglong4 *out_c,
+    ulonglong4 in_a,
+    ulonglong4 in_b
+);
+
+__constant__ char ok[] = "ok";
+__constant__ char not_ok[] = "not ok";
+
+__global__ void kernel(char *buf) {
+    ulonglong4 a = {0, 1, 2, 3};
+    ulonglong4 b = {1, 1, 1, 1};
+    ulonglong4 c = {1, 2, 3, 4};
+
+    add_u32(&c, a, b);
+
+    memcpy(buf, ok, sizeof(ok));
+}
+
+int main() {
+	char h_buf[32];
+    char *d_buf;   
+    cudaMalloc(&d_buf, 32);
+
+	kernel<<<1, 1>>>(d_buf);
+
+    cudaDeviceSynchronize();
+    cudaMemcpy(h_buf, d_buf, 32, cudaMemcpyDeviceToHost);
+    
+	printf("%s\n", h_buf);
+    cudaFree(d_buf);
+	return 0;
+}
\ No newline at end of file
diff --git a/ds/25-1/1e/op.ptx b/ds/25-1/1e/op.ptx
new file mode 100644
index 0000000..e6077d5
--- /dev/null
+++ b/ds/25-1/1e/op.ptx
@@ -0,0 +1,29 @@
+.version 8.4
+.target sm_75
+.address_size 64
+
+.visible .func add_u32(
+    .param .b64 out_c,
+    .param .align 16 .b8 in_a[32],
+    .param .align 16 .b8 in_b[32]
+) {
+    .reg .u64 %ra<4>, %rb<4>;
+    .reg .b64 %rdc;
+    
+    ld.param.b64 %rdc, [out_c];
+    
+    ld.param.v2.u64 {%ra0, %ra1}, [in_a];
+    ld.param.v2.u64 {%ra2, %ra3}, [in_a + 16];
+    ld.param.v2.u64 {%rb0, %rb1}, [in_b];
+    ld.param.v2.u64 {%rb2, %rb3}, [in_b + 16];
+
+    add.cc.u64 %ra0, %ra0, %rb0;
+    addc.cc.u64 %ra1, %ra1, %rb1;
+    addc.cc.u64 %ra2, %ra2, %rb2;
+    addc.u64 %ra3, %ra3, %rb3;
+
+    st.v2.u64 [%rdc], {%ra0, %ra1};
+    st.v2.u64 [%rdc + 16], {%ra2, %ra3};
+
+    ret;
+}
\ No newline at end of file